วิธีทำให้ AMD GPU วิ่งด้วยความเร็วระดับ ‘brrr’

 (hazyresearch.stanford.edu)
1 คะแนน โดย GN⁺ 2025-11-16 | 1 ความคิดเห็น | แชร์ทาง WhatsApp
  • HipKittens คือ ชุดโปรแกรมมิงพริมิตีฟ ที่ออกแบบมาเพื่อดึงศักยภาพของ AMD GPU ออกมาให้ได้มากที่สุด โดยปรับแต่งการเข้าถึงหน่วยความจำ การจัดตารางทำงาน และการนำแคชกลับมาใช้ซ้ำให้เหมาะสม
  • AMD MI355X GPU มีโครงสร้าง 256 compute units และ 8 chiplets (XCD) พร้อม ไฟล์รีจิสเตอร์ขนาดใหญ่ และ ชุดคำสั่ง matrix core แบบละเอียด
  • ต่างจาก NVIDIA ตรงที่ AMD ไม่มี การจัดสรรรีจิสเตอร์ใหม่, คำสั่งเมทริกซ์แบบอะซิงโครนัส และ mbarrier ทำให้ 8-wave ping-pong และ 4-wave interleave มีประสิทธิภาพมากกว่า wave specialization
  • HipKittens ใช้ การจัดตารางแบบรับรู้ชิปเล็ต (grid) เพื่อปรับปรุง locality ของแคช L2 และ LLC และเพิ่ม แบนด์วิดท์สูงสุดและ TFLOPS ในงาน GEMM และ Attention
  • แนวทางนี้ช่วยชดเชย ความไม่สมบูรณ์ของซอฟต์แวร์ใน ecosystem ของ AMD GPU และเป็นฐานสำหรับเพิ่ม ความสามารถในการขยายตัวของ AI computing บนฮาร์ดแวร์ที่หลากหลาย

สถาปัตยกรรมและลักษณะประสิทธิภาพของ AMD CDNA GPU

  • AMD MI355X GPU มี 256 compute units (CU) โดยแต่ละ CU ประกอบด้วย SIMD 4 ชุด
    • หนึ่ง SIMD จะรัน wave ที่มี 64 เธรด ซึ่งต่างจาก warp ของ NVIDIA ที่มี 32 เธรด
  • MI355X มี SRAM 165KB หรือราว 70% ของ B200 และไม่มีฟีเจอร์ คำสั่งคูณเมทริกซ์แบบอะซิงโครนัส, การจัดสรรรีจิสเตอร์ใหม่, tensor memory accelerator, mbarrier
  • ในทางกลับกัน มันมี ไฟล์รีจิสเตอร์ใหญ่กว่า 2 เท่า และ จำนวนโปรเซสเซอร์มากกว่า 60% (256 CU เทียบกับ 160 SM)
    • รองรับ ชุดคำสั่ง matrix core ขนาดเล็กและละเอียด และมีฟังก์ชัน โหลดจาก global ไป shared memory โดยตรง (คล้าย TMA)
  • AMD ใช้ สถาปัตยกรรมแบบชิปเล็ต ที่ประกอบด้วย 8 chiplets (XCD) โดยแต่ละ XCD มีแคช L2 ของตัวเอง และมี แคช LLC อยู่ชั้นบน
  • ตามตาราง MI355X มีสมรรถนะประมวลผล BF16 2.5 PFLOPs, MXFP8 5.0 PFLOPs, MXFP6 10.1 PFLOPs พร้อม หน่วยความจำ 288GB และแบนด์วิดท์ 8TB/s

ความท้าทายในการออกแบบเคอร์เนลสำหรับ AMD

  • การปรับแต่งการเข้าถึงหน่วยความจำ: เนื่องจากข้อจำกัดของคอมไพเลอร์ HIPCC และพฤติกรรม I/O ที่ไม่เปิดเผย จึงต้องให้ความสำคัญกับการออกแบบ การจัดวางข้อมูลและแพตเทิร์น swizzle
  • การจัดตารางภายในโปรเซสเซอร์: AMD ต้องอาศัย ไฟล์รีจิสเตอร์และชุดคำสั่งเมทริกซ์ขนาดเล็ก แทน shared memory
  • การจัดตารางระหว่างโปรเซสเซอร์: โครงสร้างแบบชิปเล็ตทำให้จำเป็นต้องกระจายงานโดยคำนึงถึง ผลของ NUMA ในระดับแคช

แพตเทิร์นการเข้าถึงหน่วยความจำของ HipKittens

  • HipKittens(HK) ใช้ tile เป็นหน่วยข้อมูลพื้นฐาน และมี ฟังก์ชันโอเปอเรชันคล้าย PyTorch
    • tile ถูกนิยามด้วย ชนิดข้อมูล ขนาด และเลย์เอาต์ และรองรับอินพุตที่หลากหลายด้วย C++ template metaprogramming
  • การจัดตารางรีจิสเตอร์: HIPCC ไม่สามารถใช้รีจิสเตอร์บางตัวเป็นอินพุตของ MFMA ได้ ดังนั้น HK จึงมี ฟังก์ชันตรึงรีจิสเตอร์แบบชัดเจน
    • นักพัฒนาสามารถกำหนดรีจิสเตอร์ด้วยตนเองเพื่อเขียน เคอร์เนลที่ให้ประสิทธิภาพสูงสุด
  • เลย์เอาต์ของรีจิสเตอร์: AMD มีเลย์เอาต์ที่ต่างกันตามชนิดข้อมูลและรูปแบบเมทริกซ์ จึง ไม่สามารถใช้แพตเทิร์น swizzle แบบเดียวได้
    • ตัวอย่างเช่น tile bf16 ขนาด 16×16 และ tile bf16 ขนาด 16×32 ต้องใช้แพตเทิร์น swizzle คนละแบบ
  • โครงสร้างเฟสของคำสั่ง: คำสั่ง shared memory ของ AMD มี กลุ่มเฟสที่ไม่ต่อเนื่องกัน และ มีเอกสารภายในไม่เพียงพอ
    • HK จึงมี solver ที่ย้อนวิศวกรรมขึ้นมา เพื่อจัดการเรื่องนี้
  • การสร้างแอดเดรส: AMD รองรับ การโหลดแบบอะซิงโครนัสจาก HBM ไป shared memory และทำ optimization ด้วย HBM address swizzle

การจัดตารางภายในโปรเซสเซอร์: แพตเทิร์นของ wave

  • Wave specialization มีประสิทธิภาพบน NVIDIA แต่สำหรับ AMD กลับทำให้ประสิทธิภาพลดลงเพราะ ไม่มีการจัดสรรรีจิสเตอร์ใหม่
    • wave ฝั่ง producer ยึดรีจิสเตอร์ที่ไม่จำเป็นไว้ ขณะที่ wave ฝั่ง consumer มีรีจิสเตอร์ไม่พอจนเกิด spill
  • จากผลทดลองของ HK, wave specialization บน AMD ทำให้เกิด ความเข้มข้นเชิงคณิตศาสตร์ลดลงและคอขวดด้านหน่วยความจำ
    • ตัวอย่าง: ใน GEMM การจัดแบบ HK 0/8 ได้ 1605 TFLOPs ส่วน CUTLASS ได้ 1570 TFLOPs
  • แพตเทิร์นการจัดตารางทางเลือก
    • 8-wave ping-pong: สอง wave สลับกันรัน คลัสเตอร์หน่วยความจำ/การคำนวณ
    • 4-wave interleave: หนึ่ง wave สลับการทำงานของ หน่วยความจำและการคำนวณ อย่างละเอียด
    • 8-wave เขียนโค้ดได้กระชับกว่า ส่วน 4-wave ละเอียดกว่าแต่โค้ดยาวขึ้น
    • ใน GEMM และ Attention Forward, 8-wave ทำประสิทธิภาพได้ในระดับ SoTA

การจัดตารางระหว่างโปรเซสเซอร์: แนวทางแบบรับรู้ชิปเล็ต

  • AMD MI355X มี ชิปเล็ต XCD 8 ตัว และแต่ละตัวมี แคช L2 แยกอิสระ
    • thread block ถูกจัดสรรให้ชิปเล็ตแบบ round-robin ทำให้ ลำดับของ grid ส่งผลโดยตรงต่อประสิทธิภาพการใช้แคชซ้ำ
  • การจัดวางแบบ row-major อย่างง่ายทำให้อัตราการใช้แคช L2 ซ้ำต่ำและเกิด การสูญเสียแบนด์วิดท์
    • ตัวอย่าง: L2 55%, LLC 95%, 15.1 TB/s, 1113 TFLOPs
  • HK จึงนำ การจัดตารางแบบรับรู้ชิปเล็ต (grid) มาใช้ เพื่อใช้ประโยชน์จาก locality ของแคช L2 และ LLC พร้อมกัน
    • โดยจัดกลุ่ม thread block ตาม พื้นที่ใกล้เคียงกันของเมทริกซ์ผลลัพธ์ เพื่อเพิ่มการนำข้อมูลอินพุตกลับมาใช้ซ้ำให้สูงสุด

ตัวอย่างเคอร์เนลจริง

  • hot loop ของเคอร์เนล Attention Forward และ BF16 GEMM ใช้ ตาราง 8-wave ping-pong ของ HK
    • แต่ละลูปจะสลับรัน คลัสเตอร์ Compute–Memory และซิงก์กันด้วย schedule barrier
    • ในตัวอย่างโค้ดมีการใช้โอเปอเรชันของ HK เช่น mma_AtB, load, exp2, col_sum ซ้ำ ๆ

บทสรุป: AMD ในยุค Multi-silicon AI

  • HipKittens ทำประสิทธิภาพได้ แข่งขันได้ บน AMD CDNA3 และ CDNA4
    • มีแกนหลัก 3 อย่างคือ การเข้าถึงหน่วยความจำที่ปรับแต่งแล้ว, การจัดตาราง wave ที่ออกแบบเพื่อ AMD, และ การจัดตาราง grid แบบรับรู้ชิปเล็ต
  • เคอร์เนลของ HK ทำ ประสิทธิภาพสูงสุดในฝั่ง AMD และยังแข่งขันได้กับ เคอร์เนลของ NVIDIA Blackwell
  • เพื่อความหลากหลายของ AI computing จำเป็นต้อง เพิ่มการเข้าถึง AMD GPU และ HipKittens ก็เป็น ฐานซอฟต์แวร์สำคัญ สำหรับเป้าหมายนั้น
  • การปรับปรุง HIPCC register scheduling ของ AMD ถูกชี้ว่าเป็นพื้นที่พัฒนาสำคัญในอนาคต

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

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

 
GN⁺ 2025-11-16
ความคิดเห็นจาก Hacker News
  • แนะนำให้อ่านการสนทนาเกี่ยวกับ HipKittens
  • ยังมีโพสต์ HipKittens: Fast and furious AMD kernels ที่พูดถึงงานวิจัยเดียวกันด้วย โดยมีคอมเมนต์จาก George Hotz และพนักงานของ AMD
  • ยินดีที่วงการวิชาการมาจัดการปัญหาแบบนี้ แต่สุดท้ายผมคิดว่านี่คือ ปัญหาที่ AMD ต้องแก้เองภายในบริษัท
    • ผมมองว่าบริษัทฮาร์ดแวร์ควรทำฮาร์ดแวร์อย่างเดียว แบบนั้นจะทำให้ แรงจูงใจยังคงบริสุทธิ์ แม้ประสิทธิภาพจะลดลง 20% ผมก็คิดว่ายังดีกว่า
    • เห็นด้วยอย่างยิ่ง AMD ผัดปัญหานี้มา 10 ปีแล้ว และเพิ่งจะมาพยายามไล่ตามตอนนี้ ฮาร์ดแวร์ยอดเยี่ยม แต่ ขาดความสามารถในการพัฒนาเฟิร์มแวร์ เลยดึงศักยภาพออกมาไม่ได้
    • แต่ทีมวิจัยนี้ก็เคยสร้างซอฟต์แวร์คล้ายกันสำหรับ GPU ของ Nvidia มาก่อน ดูเหมือนเป็นนักวิจัยฝีมือดีที่กำลังใช้ความเชี่ยวชาญของตัวเอง
    • เท่าที่ผมรู้ AMD ก็กำลังจัดการปัญหานี้อยู่แล้วในหลายระดับ และกำลังร่วมงานกับ tinycorp ด้วย
  • จากบทความ ทำให้รู้สึกว่าการปรับแต่งทำได้ยากเพราะ ความซับซ้อนเชิงสถาปัตยกรรม ของ AMD GPU แต่ในระยะยาวแนวทางของ AMD อาจขยายต่อได้ดีกว่า ตอนนี้ Nvidia ใช้ 2 chiplet ในขณะที่ AMD มีโครงสร้าง 8 chiplet เลยมี ปัญหาเรื่อง memory locality ในอนาคตจำนวน chiplet ก็น่าจะเพิ่มขึ้นอีก ดังนั้นประสบการณ์ในการรับมือกับความซับซ้อนตอนนี้อาจช่วยได้ในระยะยาว
    • AMD ไม่ต้องใช้ warp specialization เพื่อให้ได้ประสิทธิภาพสูง ดังนั้นการเขียนโปรแกรมจึงง่ายกว่า
  • มีนักพัฒนาจำนวนมากพยายามทำให้ AMD GPU ‘go brrr’ สำหรับนักพัฒนาทั่วไป แต่ก็ล้มเหลว ผมไม่เข้าใจว่าทำไม AMD ถึงไม่ แก้ปัญหาซอฟต์แวร์ด้วยตัวเอง ตอนนี้ก็มีเงินมากพอแล้ว การไม่ยอมจ้างนักพัฒนาจึงไม่ใช่ข้ออ้าง GPU สำหรับดาต้าเซ็นเตอร์ก็ไม่ได้แย่ แต่ถ้าคนทั่วไปจะทดลอง ML·AI ก็ยังรู้สึกว่า Nvidia ดีกว่ามาก RTX 3090 อายุ 5 ปีของผมยังรู้สึกว่าดีกว่า AMD consumer GPU ทุกตัวที่ออกมาจนถึงตอนนี้
    • ประสบการณ์นักพัฒนาของ AMD แย่มาก ขนาดรายงานบั๊กไดรเวอร์ที่ทำให้ล่มยังไม่รับเลย
    • เมื่อไม่นานมานี้ผมเปลี่ยนเซิร์ฟเวอร์ inference จาก NVidia 5090 ไปเป็น AMD R9700 32GB สองใบ และเป็นประสบการณ์ที่ดีทั้งหมด ใช้กับเคอร์เนล Fedora ได้ทันทีโดยไม่ต้องตั้งค่า DKMS และเชื่อมคอนเทนเนอร์ผ่าน ROCm ก็ง่าย แค่เปลี่ยนการตั้งค่าของ Ollama กับ Storyteller ก็จบ เป็น ประสบการณ์ที่สบายกว่า CUDA มาก
    • Nvidia ถึงขั้นดูแล Unreal Engine fork ด้วยตัวเอง AMD ยังเทียบไม่ติดเลย
    • ในบรรดาบริษัทฮาร์ดแวร์ Nvidia เป็นรายเดียวที่ให้ ค่าตอบแทนแข่งขันได้กับวิศวกรซอฟต์แวร์ AMD ยังมีวัฒนธรรมที่มองซอฟต์แวร์ไม่ใช่ ‘งานจริง’ อยู่ และความเคยชินแบบนี้เปลี่ยนยาก
  • Mojo เคยมีไอเดียที่จะปรับปรุงประสบการณ์นักพัฒนา (devX) บน AMD GPU และผมสงสัยว่าตอนนี้ไปถึงไหนแล้ว
  • ผมไม่เข้าใจจริง ๆ ว่าทำไม AMD ถึง ไม่ลงทุนหลายพันล้านดอลลาร์เพื่อพัฒนาซอฟต์แวร์ Nvidia เป็นบริษัทที่มีมูลค่าสูงที่สุดในโลก และ AMD ก็เป็นคู่แข่งเพียงรายเดียว
    • AMD ก็พยายามอยู่ แต่ผมคิดว่าการเปลี่ยนวัฒนธรรมองค์กรที่อัปเดตฮาร์ดแวร์ทุกปีให้กลายเป็น วัฒนธรรมที่เน้นซอฟต์แวร์ เป็นเรื่องยาก ซอฟต์แวร์ไม่ได้สร้างรายได้ให้เห็นทันทีแบบฮาร์ดแวร์ ผู้บริหารจึงมักจัดลำดับความสำคัญไว้ต่ำกว่า อีกทั้งการที่ผู้ขายภายนอกส่งโค้ดแบบโอเพนซอร์สมาก็ดูดีในระยะสั้น แต่ส่งผลเสียต่อคุณภาพระยะยาว ถ้าพลาดเทรนด์ฮาร์ดแวร์เพียงครั้งเดียว ก็เสี่ยงจะถูกคู่แข่งทิ้งห่าง
    • ผมเคยทำงานกับผู้ผลิต GPU หลายเจ้า มีแค่ Nvidia ที่มองซอฟต์แวร์เป็น ทรัพย์สิน (asset) และยอมลงทุน บริษัทอื่นมองเป็นแค่ต้นทุน
  • ส่วนตัวผมไม่ได้ชอบมีม “go brr” เท่าไร แต่พอเห็นถูกใช้ในที่อย่าง Stanford ก็รู้สึกขำดี
    • จริง ๆ ใช้ “go brr” กันมาตั้งแต่ตอนประกาศ ThunderKittens เมื่อปีที่แล้วแล้ว
    • ถ้ามีมแบบนี้ไปโผล่ในช่องทางทางการของมหาวิทยาลัย ก็อาจเป็น สัญญาณว่ากระแสจบแล้ว
  • ตัวโปรเจกต์ยอดเยี่ยม แต่ก็สงสัยว่าทำไม AMD ไม่ทำเอง AMD ดูเหมือนยังไม่เข้าใจความสำคัญของ software stack ที่สุกงอม ยังต้องมี stack แบบรวมศูนย์ที่ใช้ได้กับทุกการ์ดเหมือน CUDA เมื่อก่อนผมเคยเชื่อว่า AMD จะตามทันสักวัน แต่ตอนนี้แทบจะถอดใจแล้ว
  • โปรเจกต์ดี แต่ตัวบทความเองให้ความรู้สึกว่า เขียนออกมาแปลก ๆ
    • บทความดูประหลาดเกินไป เหมือนพึ่ง AI มากเกินไป หรือพยายามเลียนแบบสำนวนของ AI มีประโยคอย่าง “ไปดู part one” หรือ “วิธีทำให้ AMD GPU go brr” ซ้ำไปซ้ำมา ส่วนที่น่าเสียดายเป็นพิเศษคือประเด็นทางเทคนิคที่ควรอธิบายด้วยกราฟ กลับถูก เขียนออกมาเป็นโค้ดยาว 100 บรรทัด แทน