HipKittens: เคอร์เนล AMD ที่รวดเร็วและทรงพลัง

 (hazyresearch.stanford.edu)
3 คะแนน โดย GN⁺ 2025-11-16 | 1 ความคิดเห็น | แชร์ทาง WhatsApp
  • HipKittens เป็นโปรเจกต์ที่มอบ เคอร์เนลประสิทธิภาพสูงและพริติมิทีฟการเขียนโปรแกรมบน C++ สำหรับ AMD GPU เพื่อเพิ่มประสิทธิภาพการประมวลผล AI
  • ในระบบนิเวศ AMD เดิม AITER, PyTorch, Triton, TileLang, Composable Kernel เป็นต้น ได้เผยให้เห็นข้อจำกัดจาก ประสิทธิภาพที่ไม่เสถียรและการรองรับที่ยังไม่สมบูรณ์
  • HipKittens ยึด tile abstraction เป็นแกนกลาง โดย คงอินเทอร์เฟซร่วมระหว่าง NVIDIA และ AMD ไว้ พร้อมแยกส่วนการติดตั้งใช้งานตามฮาร์ดแวร์
  • เคอร์เนลที่เขียนด้วยโค้ดไม่เกินราว 500 บรรทัด ทำประสิทธิภาพได้เร็วกว่าคอร์เนลแอสเซมบลีแบบเขียนมือของ AMD เดิม
  • นำเสนอฐานที่ใช้งานได้จริงสำหรับการขยายงานประมวลผล AI ไปสู่ สภาพแวดล้อมแบบ multi-silicon พร้อมชี้ให้เห็นความเป็นไปได้ในการเปลี่ยนผ่านสู่ ระบบนิเวศฮาร์ดแวร์แบบเปิด

สถานะปัจจุบันและปัญหาของระบบนิเวศ AMD GPU

  • ที่ผ่านมา งานประมวลผล AI เติบโตโดยมี ผู้จำหน่ายฮาร์ดแวร์รายเดียวเป็นศูนย์กลาง แต่ AMD GPU ในเจเนอเรชันล่าสุดมอบ สมรรถนะการประมวลผลและแบนด์วิดท์หน่วยความจำระดับแนวหน้า
    • ตัวอย่าง: AMD MI355X OAM มี BF16 2.5 PFLOPs, MXFP8 5.0 PFLOPs, หน่วยความจำ 288GB, แบนด์วิดท์ 8.0TB/s
  • อย่างไรก็ตาม เนื่องจาก ขาดซอฟต์แวร์สแตกที่สุกงอมเพียงพอ ประสิทธิภาพนี้จึงยังถูกนำไปใช้ในเวิร์กโฟลว์ AI จริงได้ไม่เต็มที่
  • องค์ประกอบซอฟต์แวร์หลักของ AMD ประกอบด้วย AITER, PyTorch, Triton, Mojo, TileLang, Composable Kernel (CK) เป็นต้น
    • เคอร์เนล SDPA Llama GQA backward ของ AITER และ PyTorch ทำได้เพียง 30% และ 24% ของประสิทธิภาพระดับ SoTA ตามลำดับ
    • Mojo ติดข้อจำกัดอยู่ที่ ประสิทธิภาพระดับ 50% เนื่องจากปัญหา bank conflict
    • TileLang รองรับถึงเพียง CDNA3 และมี ฟีเจอร์สำคัญไม่ครบพร้อมการพึ่งพา CK ทำให้ความซับซ้อนเพิ่มขึ้น
    • Triton มีความยากลำบากในการติดตามอายุการใช้งานรีจิสเตอร์และการปรับแต่งการเข้าถึงหน่วยความจำ
  • ผลลัพธ์คือ เคอร์เนล AMD ที่ทำประสิทธิภาพสูงสุดยังต้องให้ผู้เชี่ยวชาญเขียนแอสเซมบลีด้วยตนเอง ซึ่งเป็นข้อจำกัดใหญ่ด้านการขยายและการบำรุงรักษา

เปรียบเทียบกับระบบนิเวศที่มี CUDA เป็นศูนย์กลาง

  • ในชุมชน AI มีการประเมินว่ามี กำแพงการเข้าถึงของระบบนิเวศ CUDA (CUDA moat) อยู่
  • ในอดีต การพัฒนาเคอร์เนลของ NVIDIA ก็ต้องอาศัย ความพยายามหลายปีบนพื้นฐาน CUDA/CUTLASS ระดับล่าง
  • ช่วงหลังมานี้ การพัฒนาเคอร์เนล NVIDIA ง่ายขึ้นด้วยความก้าวหน้าของ DSL (ภาษาจำเพาะโดเมน) และ เครื่องมือช่วยด้วย AI
    • ตัวอย่าง: ThunderKittens (TK), CuTe DSL, TinyGrad “tinykittens”, TileLang, Gluon เป็นต้น
  • จากกระแสดังกล่าว ฝั่ง AMD จึงเริ่มสำรวจ ความจำเป็นของพริติมิทีฟการเขียนโปรแกรมแบบใหม่

หลักการออกแบบของ HipKittens

  • HipKittens เป็น เวอร์ชันที่พัฒนาสำหรับ AMD ต่อจาก ThunderKittens (NVIDIA) และ ThunderMittens (Apple Silicon)
  • แนวคิดหลัก: tile abstraction
    1. ความเป็นสากลของ tile abstraction – โมเดลการคำนวณแบบอิงไทล์ที่ใช้ได้ผลบน NVIDIA สามารถนำมาใช้กับ AMD ได้อย่างเป็นธรรมชาติ
    2. การติดตั้งใช้งานแบ็กเอนด์ที่เฉพาะกับสถาปัตยกรรม – รูปแบบการเข้าถึงหน่วยความจำและการจัดตารางรีจิสเตอร์ถูกออกแบบต่างกันตามฮาร์ดแวร์
    3. ความยืดหยุ่นของกลยุทธ์การจัดตาราง – บน CDNA3·CDNA4 ของ AMD การจัดตารางแบบอิง wave ไม่มีประสิทธิภาพ แต่การจัดตารางในระดับไทล์ยังคงรักษาทั้งความเรียบง่ายและประสิทธิภาพไว้ได้
  • ด้วยการแยก อินเทอร์เฟซ (ไทล์และการดำเนินการ) ออกจาก การติดตั้งใช้งาน (การแมปเข้ากับฮาร์ดแวร์) จึงนำเสนอ โมเดลที่ใช้ร่วมกันได้กับสถาปัตยกรรม GPU หลากหลายแบบ

ผลลัพธ์ด้านประสิทธิภาพ

  • เคอร์เนล Attention Forward: เขียนด้วยโค้ดเพียงประมาณ 500 บรรทัด และทำ ประสิทธิภาพได้ดีกว่าเคอร์เนลแอสเซมบลีของ AITER
    • เหนือกว่าทั้งในการตั้งค่า Causal/Non-Causal และใน head dimension (D) กับ sequence length (N) ที่หลากหลาย
  • เคอร์เนล GEMM: ประกอบด้วย ลูปหลักไม่ถึง 100 บรรทัด และทำ ประสิทธิภาพสูงสุดทั้งในการคำนวณ BF16 และ FP8
  • ยังยืนยันการปรับปรุงเมื่อเทียบกับประสิทธิภาพสูงสุดเดิมใน เคอร์เนล Attention Backward, Rotary, Fused Dropout-Residual-LayerNorm
  • ทุกเคอร์เนลยังคง อ่านเข้าใจง่ายและแก้ไขได้สะดวก ขณะเดียวกันก็ได้ ประสิทธิภาพระดับแอสเซมบลีแบบเขียนมือ

การขยายสู่ AI แบบ multi-silicon

  • หากต้องการปลดปล่อยศักยภาพของ AI จำเป็นต้องใช้ ฮาร์ดแวร์ที่หลากหลายและเปิดกว้าง
  • HipKittens มีเป้าหมายทำให้ AMD GPU กลายเป็น แพลตฟอร์มที่นักพัฒนา AI เข้าถึงได้จริงในทางปฏิบัติ
  • เมื่อนำไปรวมกับ ซอฟต์แวร์สแตกโอเพนซอร์ส ของ AMD ก็ชี้ให้เห็นความเป็นไปได้ของการเปลี่ยนผ่านไปสู่ ระบบนิเวศ AI แบบ multi-silicon
  • บทความถัดไปจะกล่าวถึง รายละเอียดเชิงเทคนิคของโครงสร้าง HipKittens (เกริ่น part two)

บทความที่เกี่ยวข้อง

1 ความคิดเห็น

 
GN⁺ 2025-11-16
ความคิดเห็นบน Hacker News

  • เราทำสัญญากับ AMD และกำลังเทรน Llama 405B บน MI350X สำหรับ MLPerf
    สถานการณ์ของ AMD ดีขึ้นอย่างชัดเจน ถ้าคุณมี AMD GPU แนะนำให้กด Linux+ROCm ที่ pytorch.org แล้วลองติดตั้ง PyTorch ดู เมื่อ 3 ปีก่อนยังสิ้นหวังมาก แต่ตอนนี้ฟีเจอร์หลัก ๆ ส่วนใหญ่ใช้งานได้ดีแล้ว ผมรัน nanochat บน MI300X แล้วใช้ได้ทันที MI350X ก็เสถียรเหมือนกัน
    แน่นอนว่ายังตามหลัง NVIDIA อยู่ และยังต้องลงทุนกับ ecosystem ซอฟต์แวร์ คอมไพเลอร์ และไดรเวอร์ อีกมาก แต่ถ้าเทียบกับสภาพอันสิ้นหวังเมื่อ 2 ปีก่อน ตอนนี้เริ่มมองเห็นความหวังแล้ว
    ถ้าอยากเห็นจุดที่ LLVM backend ของ AMD ยังขาดอยู่ ลองเทียบโค้ดของ HipKittens กับ CUDA Kittens ดูได้
    สำหรับงานเทรน NVIDIA กับ Google คืออันดับ 1, AMD คืออันดับ 2, ที่เหลือแทบไม่มีอะไรเลย Intel กับ Tenstorrent ยังอีกไกล, Huawei รันตัวอย่างแล้วตายด้วย segfault, Groq เลิกขายชิปไปแล้ว, Cerebras ก็หาของไม่ได้เลย ส่วน Trainium ใช้เวลา 5 วันกว่าจะได้อินสแตนซ์เดียวจนหมดความสนใจ

    • สงสัยว่า tinygrad ยังห่างแค่ไหนกว่าจะสามารถแสดงหรือสำรวจฟีเจอร์แบบนี้ได้ เช่น การเพิ่มประสิทธิภาพหน่วยความจำหรือการปรับเฉพาะ warp และอยากรู้ด้วยว่าการเพิ่มฟีเจอร์พวกนี้เข้า tinygrad ซับซ้อนแค่ไหน
    • สงสัยว่าถ้าจะรัน ROCm + PyTorch บนฮาร์ดแวร์ฝั่งผู้บริโภค (ไม่ใช่ MI) ยังจำเป็นต้องใช้ ไดรเวอร์เคอร์เนลแบบ proprietary หรือเปล่า
    • คำว่า “Cerebras หาจากที่ไหนก็ไม่ได้” ฟังดูเหมือนเป็น สัญญาณแห่งชัยชนะ มากกว่า
    • ผมเป็น CEO ของ AMD NeoCloud มาตลอด 2 ปีที่ผ่านมา เลยดีใจมากที่ได้เห็น AMD พลิกฟื้น แบบนี้ด้วยตาตัวเอง
      การตั้งค่าเริ่มต้นยังค่อนข้างหยาบอยู่บ้าง แต่ดีขึ้นกว่าสมัยก่อนมาก ตัวอย่างเช่น แค่เดือนก่อน nanochat ยังทำงานไม่ค่อยได้ แต่ตอนนี้ใช้ได้แล้ว เรื่องสำคัญคือ ตอนนี้ผู้คนเริ่มสนใจ ecosystem ของ AMD แล้ว
      AI ต้องการ ความหลากหลายของฮาร์ดแวร์ การปล่อยให้บริษัทเดียวผูกขาดทั้งฮาร์ดแวร์และซอฟต์แวร์ AI อาจดีต่อผู้ถือหุ้น แต่เป็นผลเสียต่อความก้าวหน้าทางเทคโนโลยี

  • ผมไม่เข้าใจ มูลค่ากิจการ ของ NVIDIA ตอนนี้มีอัลกอริทึมไม่กี่แบบอย่าง Transformer ที่ชนะไปแล้ว และข้อมูลก็สำคัญขึ้นมาก ไม่เหมือนในอดีตที่ต้องรองรับโค้ด HPC หลากหลายแบบ ดังนั้นตอนนี้คู่แข่งแค่ต้องปรับให้เหมาะกับชุดอัลกอริทึมแคบ ๆ เท่านั้น
    ถ้ารัน vLLM กับ Transformer ได้อย่างมีประสิทธิภาพ ก็จะมีตลาดมหาศาลเปิดขึ้นมา ถ้าอย่างนั้น AMD หรือ Huawei ก็น่าจะตามทันได้ไม่ยาก เลยสงสัยว่า คูเมืองป้องกันธุรกิจ (moat) ที่แท้จริงของ NVIDIA คืออะไร หรือว่าแค่ InfiniBand ก็พอแล้ว?

    • ใช่แล้ว คูเมืองของ NVIDIA กำลังอ่อนลงเรื่อย ๆ บริษัทยักษ์ใหญ่อย่าง MS, Google, Amazon, Apple ต่างก็พัฒนาชิปของตัวเองเพื่อหลีกเลี่ยงการพึ่งพา NVidia
      ใน GPU สำหรับดาต้าเซ็นเตอร์ NVIDIA ยังแข็งแกร่งอยู่ แต่ Google ก็ใช้ TPU ในสเกลใหญ่ และ OpenAI ก็สั่งฮาร์ดแวร์จาก AMD ด้วย
      ecosystem ของ CUDA ยังสำคัญอยู่ แต่ทุกคนก็กำลังเคลื่อนไหวเพื่อออกจากตรงนั้น AMD, Intel, Qualcomm และรายอื่น ๆ ก็ลงมาแข่งกันแล้ว HipKittens ดูเหมือนเป็นก้าวสำคัญไปสู่ ecosystem ซอฟต์แวร์แบบเป็นกลาง
    • วิธีที่ง่ายที่สุดในการรองรับ “อัลกอริทึมไม่กี่แบบ” ก็คือสร้าง ฮาร์ดแวร์คำนวณเอนกประสงค์ เพราะฮาร์ดแวร์มีรอบพัฒนาที่ยาว วิธีนี้เลยปลอดภัยกว่า
    • InfiniBand กำลังถูกแทนที่ด้วย UEC สำหรับ inference ไม่จำเป็นต้องใช้ InfiniBand เพราะงั้นตลาด inference ไม่มีคูเมืองป้องกัน AMD หรือ Google TPU เป็นตัวเลือกที่ฉลาดกว่า
    • อาวุธที่แท้จริงของ NVIDIA คือ ecosystem ไลบรารี CUDA อันมหาศาล ซึ่งมีโค้ดรองรับแทบทุกสาขา
    • Transformer ไม่ใช่เทคนิคเดียวจบ วิธี implement มีได้หลากหลายมาก นั่นจึงเป็นเหตุผลว่าทำไม vLLM หรือ TRL ถึงไม่ง่ายอย่างที่คิด

  • AMD น่าจะลงทุนสนับสนุนโครงการแบบนี้ได้ แต่ผมคิดว่าคง พลาดโอกาสอีกแล้ว ตามสไตล์เดิมเรื่อง GPU และ AI

    • AMD ลงทุนแค่ซอฟต์แวร์ขั้นต่ำที่พอให้ส่งสินค้าได้เท่านั้น แม้แต่ ระบบทดสอบประสิทธิภาพก็ยังอ่อนมาก และบั๊ก regression ก็ถูกส่งต่อถึงลูกค้าแบบนั้นเลย
      HipKittens เป็นการพัฒนาที่ดีขึ้น แต่ภายใน AMD ยังขาดความสามารถในการติดตามประสิทธิภาพของเคอร์เนล ฝั่ง DevOps ก็เอาต์ซอร์สให้ TCS อินเดีย ซึ่งคนฝั่งนั้นก็ไม่เข้าใจสถานการณ์ดีพอ
      ทีมที่มีผู้นำดี ๆ จะตั้ง ทีม IT เงา ของตัวเอง ROCm ไม่เคยมีทีมแบบนั้น สุดท้ายต้องรอให้ลูกค้าคลาวด์รายใหญ่โวยวายก่อนถึงเริ่มดีขึ้น
      ต่อให้รับคนเก่งเข้ามา ก็ยังเสนอ เงินเดือนต่ำกว่าตลาด เพราะไปเทียบกับระดับ Qualcomm หรือ Walmart
      ตลอด 4 ปีที่ผ่านมา ไม่เคยมีปีไหนที่โบนัสจ่ายครบเต็มจำนวน
    • คำพูดที่ว่า “AMD ไม่เคยพลาดโอกาสที่จะพลาดโอกาส” นี่ตรงมาก ฮาร์ดแวร์ Instinct จริง ๆ แล้ว แข่งขันกับ NVidia ได้แทบเต็มตัว แต่การรองรับด้านซอฟต์แวร์แย่มาก
      ในอดีต AMD เคยหยุดซัพพอร์ตแบบ Radeon VII หรือชอบรื้อ ecosystem ใหม่บ่อย ๆ จนไม่เคยสุกงอม ด้วยความที่เข้ากันได้กับ CUDA ไม่ดีและไม่ยอมลงทุน องค์กรส่วนใหญ่เลยเลือก NVIDIA ไปเลย
      ถึงอย่างนั้นช่วงหลัง AMD ก็ลงทุนกับ ROCm และ ecosystem ของ Instinct อย่างต่อเนื่อง เลยค่อย ๆ ดีขึ้น แต่ใน ด้านเครือข่าย NVIDIA ก็ยังนำห่างอยู่มาก
    • ดูจากตารางเทียบประสิทธิภาพแล้ว AMD อาจไปถึงระดับ NVIDIA ได้ตั้งนานแล้ว แต่กลับล้มเหลวเพราะ ไม่มีซอฟต์แวร์รองรับ ทั้งที่การออกแบบชิปเป็นงานที่ยากกว่าเสียอีก ปัญหาคือพวกเขากลับไม่เข้าใจซอฟต์แวร์
    • ก่อนหน้านี้เคยมีคนส่งเคอร์เนลที่ปรับแต่งสำหรับ ROCm เข้าไป แต่ AMD กลับ ปิด PR แล้วไม่ merge เป็นพฤติกรรมที่เข้าใจไม่ได้จริง ๆ
    • ตอนนี้มีเงินสนับสนุนแล้ว และโครงการก็ ทำงานได้ตามปกติ

  • สงสัยว่ามีโครงการไหนสร้างอยู่บน ThunderKittens บ้างไหม
    ถ้านี่คือเวอร์ชันที่พอร์ตมาเป็น HIP จริง ๆ ก็ดูมี คุณค่าเชิงใช้งานจริง มากกว่าการพอร์ต CUDA แบบตรง ๆ เยอะ

  • คำว่า “raw assembly vs cooked assembly” น่าสนใจดี
    สมัยก่อนการ เขียนโค้ดแอสเซมบลี สำหรับ CPU ด้วยมือเป็นเรื่องปกติ บน GPU ก็ไม่จำเป็นต้องกลัวเกินไป สุดท้ายแล้วคอมไพเลอร์ก็เป็นคนสร้างโค้ดนั้นออกมาอยู่ดี

  • ถ้ามองในเชิงคณิตศาสตร์ inference ก็คือการคำนวณ linear algebra/BLAS พื้นฐาน
    เลยสงสัยว่าอาจทำ API inference แบบกระชับ ที่ครอบคลุม use case ได้ 80% โดยคำนึงถึง dtype และ sparsity ได้หรือไม่ มันคงไม่จำเป็นต้องซับซ้อนแบบ CUDA

  • composable-kernel คือซอฟต์แวร์ที่ทำให้เครื่อง Gentoo ของผมเจอ OOM (หน่วยความจำไม่พอ) บ่อยที่สุด ตอนที่ Clang คอมไพล์ CK มันกินเกิน 2.5GB ต่อเธรด

    • ผมเคยตรวจแพ็กเกจ CK สำหรับ Debian แล้ว ตอน build ด้วย -j32 แม้แต่เวิร์กสเตชัน RAM 64GB ก็ยัง OOM แต่พอลองรันด้วย -j1 ก็สำเร็จหลังผ่านไป 190 ชั่วโมง
      บนเซิร์ฟเวอร์ build อย่างเป็นทางการน่าจะต้องลดจำนวนสถาปัตยกรรมลง ได้ยินมาว่าแพ็กเกจที่เป็น dependency ส่วนใหญ่ต้องใช้แค่ header เท่านั้น ตอนนี้กำลังมีงานปรับปรุงเพื่อลดเวลา build อยู่
    • แค่การ instantiate template สำหรับเคอร์เนลเดียวยังใช้เวลาเกิน 10 นาที นี่มัน น่าทึ่งมาก
      device_grouped_conv2d_fwd_xdl_ngchw_gkcyx_ngkhw_f16_instance ไม่รู้เลยว่าพวกเขาไปทำอะไรกับ clang บ้าง

  • สงสัยว่าความก้าวหน้าแบบนี้จะช่วยให้ชิป AMD ฝั่งผู้บริโภค รัน LLM ได้ดีขึ้น หรือเปล่า
    เช่น ผมกำลังมองโน้ตบุ๊ก HP OMEN MAX 16 ที่มี AMD CPU/iGPU กับ RTX 5080 อยู่ด้วยกัน ฝั่ง AMD จะพอสู้กับ RTX ได้ไหม?

    • หลายคนอาจคิดว่า dGPU ต้องเร็วกว่าตลอด แต่ข้อจำกัดด้าน ความจุหน่วยความจำ มักเป็นตัวถ่วง
      บล็อกโพสต์นี้ หรือผลลัพธ์บน Mac ระดับสูงก็น่าสนใจ
    • ผมกำลังรัน Qwen3 Coder 30B บน RTX7900XTX ด้วย Ollama อยู่ มันทำงานได้ค่อนข้างดี ดูเหมือนว่าบางโหลดจะไหลไปที่หน่วยความจำระบบกับ Ryzen 7 CPU ด้วย
      มันช้ากว่า Sonnet 4 API นิดหน่อย แต่ถ้าได้ ประสิทธิภาพระดับนี้บนเครื่องตัวเอง ก็ถือว่าน่าพอใจมาก
      ชุด Opencode + Ollama + Qwen3 Coder เป็นทางเลือกที่ยอดเยี่ยมของ ClaudeCode + Sonnet4
      แต่ถ้าคุณอยู่ในสถานการณ์ที่ต้องให้ AI เขียนโค้ดแทนทั้งหมด ก็อาจรู้สึกต่างออกไป อย่างไรก็ตาม ถ้าใช้เป็นผู้ช่วยส่วนตัวถือว่าสมบูรณ์แบบ

  • ไม่เข้าใจว่าทำไม AMD ถึง เมิน B300 ไปแบบสมบูรณ์