ประสิทธิภาพของ GPU ที่พุ่งทะยานแบบก้าวกระโดด

 (hazyresearch.stanford.edu)
1 คะแนน โดย GN⁺ 2024-05-13 | 1 ความคิดเห็น | แชร์ทาง WhatsApp

คุณลักษณะของ GPU H100

  • มาพร้อมหน่วยความจำ HBM3 ขนาด 80GB และแบนด์วิดท์ 3TB/s (ในทางปฏิบัติจะต่ำกว่านี้เล็กน้อย)
  • มีแคช L2 ขนาด 50MB และแบนด์วิดท์ 12TB/s โดยภายใน GPU ถูกแบ่งเป็น 2 ส่วน ๆ ละ 25MB และเชื่อมต่อกันด้วย crossbar (ซึ่ง crossbar เป็นปัจจัยที่ทำให้ประสิทธิภาพลดลง)
  • ประกอบด้วย 132 Streaming Multiprocessor (SM) และแต่ละ SM มีโครงสร้างดังนี้:
    • มี shared memory ได้สูงสุด 227KB ภายในแคช L1 ขนาด 256KB (รวมกันแล้วให้แบนด์วิดท์ราว 33TB/s)
    • มี Tensor Memory Accelerator (TMA) ที่รองรับการสร้างแอดเดรสแบบ asynchronous และการ fetching หน่วยความจำ แม้ยังมีความสามารถอย่างการรองรับเครือข่ายหน่วยความจำบนชิป แต่โพสต์นี้จะไม่ลงรายละเอียดในส่วนนั้น
    • แบ่งเป็น 4 quadrant โดยแต่ละ quadrant ประกอบด้วย warp scheduler, vector register 512 ตัว (แต่ละตัวมี 32 คำขนาด 4 ไบต์), tensor core สำหรับ matrix multiply และ instruction ภายในที่รองรับการคำนวณขนานอย่าง sum/multiply เป็นต้น

เคล็ดลับการปรับแต่งประสิทธิภาพ H100 GPU

  • สิ่งสำคัญคือการใช้ Tensor Core ให้ได้มากที่สุด H100 มีสมรรถนะการคูณเมทริกซ์ FP16 สูงถึง 989 TFLOPs ดังนั้นอัตราการใช้งาน Tensor Core จึงส่งผลอย่างมากต่ออัตราการใช้ GPU โดยรวม
  • อย่างไรก็ตาม หากต้องการใช้ Tensor Core ให้เต็มที่ ต้องคำนึงถึงประเด็นต่อไปนี้:
    • จำเป็นต้องใช้คำสั่ง Warp Group Matrix Multiply Accumulate (WGMMA) แต่ใช้งานค่อนข้างยาก
    • Shared Memory ไม่ได้เร็วอย่างที่คิด และต้องใช้อย่างระมัดระวัง
    • การสร้างแอดเดรสมีต้นทุนสูง จึงควรปรับแต่งด้วย Tensor Memory Accelerator (TMA) เป็นต้น
    • การเพิ่ม occupancy ยังคงช่วยได้ และ register คือทรัพยากรหลัก
  • คุณลักษณะเหล่านี้ไม่ได้มีผลแค่กับ H100 เท่านั้น แต่ใช้ได้กับ GPU รุ่นอื่นในระดับหนึ่งเช่นกัน อย่างไรก็ตาม บน H100 การใช้ Tensor Core มีความสำคัญและท้าทายเป็นพิเศษ

ThunderKittens : CUDA embedded DSL ที่ปรับแต่งมาสำหรับ H100

  • เป็น embedded DSL บนพื้นฐาน CUDA ที่พัฒนาขึ้นเพื่อดึงประสิทธิภาพของ GPU รุ่นใหม่อย่าง NVIDIA H100 ออกมาให้ได้มากที่สุด
  • มี template type แบบอิง tile อยู่ 4 แบบดังนี้:
    • Register Tile (เทนเซอร์ 2D ที่เก็บอยู่ใน register)
    • Register Vector (เทนเซอร์ 1D ที่เก็บอยู่ใน register)
    • Shared Tile (เทนเซอร์ 2D ที่เก็บอยู่ใน Shared Memory)
    • Shared Vector (เทนเซอร์ 1D ที่เก็บอยู่ใน Shared Memory)
  • นอกจากนี้ยังมีโอเปอเรเตอร์หลากหลายสำหรับจัดการ tile (เช่น exp, mul, sum ฯลฯ) ในระดับ warp หรือ warp group
  • เมื่อนำ ThunderKittens ไปใช้เขียน kernel ของ Flash Attention และ Flash Attention 2 พบว่าโค้ดกระชับขึ้นมาก และบน H100 ให้ประสิทธิภาพเพิ่มขึ้นสูงสุด 30%
  • ยังมีการใช้ ThunderKittens เขียน kernel ของ Based Linear Attention จนทำได้ถึง 215 TFLOPs (และหากรวม recompute ตามลักษณะของอัลกอริทึม จะเกิน 300 TFLOPs)

ข้อพิจารณาในเชิงปรัชญา

  • เหตุผลที่ ThunderKittens ทำงานได้ดี คือมันไม่ได้พยายามรองรับทุกอย่าง แต่เลือกใช้ abstraction แบบอิง tile ที่เรียบง่ายและสอดรับกับสถาปัตยกรรม GPU อย่างมาก
  • การใช้ vector register ขนาด 1024-bit แทน 32-bit word แบบดั้งเดิมก็เป็นพัฒนาการหนึ่ง แต่สิ่งที่สำคัญกว่าคือการเปลี่ยนกรอบคิดให้มอง tile ขนาด 16x16 เป็นหน่วยพื้นฐานของ register
  • เมื่อพิจารณาว่า workload ด้าน AI ส่วนใหญ่ลงท้ายที่ matrix multiply, reduction และ reshape แนวทางแบบอิง tile จึงสมเหตุสมผล และในมุมฮาร์ดแวร์ก็มีแนวโน้มจะพัฒนาไปในทิศทางที่รองรับ matrix multiply ขนาดเล็กเพิ่มเติมนอกเหนือจาก systolic array
  • ยิ่งไปกว่านั้น เราจำเป็นต้องเปลี่ยนวิธีคิดไปสู่การออกแบบอัลกอริทึม AI ในรูปแบบที่เหมาะกับฮาร์ดแวร์โดยตรง เช่น จำกัดขนาด state ของ RNN ให้พอดีกับ SM และปรับความหนาแน่นของการคำนวณให้สอดคล้องกับระดับที่ฮาร์ดแวร์ต้องการ

ความเห็นของ GN⁺

  • ThunderKittens อาจเป็นตัวเลือกที่น่าสนใจสำหรับนักพัฒนาที่คุ้นเคยกับ CUDA เพราะสามารถยกระดับประสิทธิภาพได้ง่ายโดยแทบไม่ต้องแก้ kernel เดิมมากนัก
  • อย่างไรก็ตาม สำหรับผู้เริ่มต้น อุปสรรคในการเข้าถึงยังค่อนข้างสูง จึงน่าจะต้องมีตัวอย่างโค้ดและสื่อการเรียนรู้ที่หลากหลายกว่านี้ในอนาคต
  • แผนการขยายการรองรับ ThunderKittens ไปยัง GPU อื่นอย่าง AMD นอกเหนือจาก H100 ก็น่าสนใจ เพราะอาจช่วยลดการผูกติดกับผู้ผลิตรายเดียวได้
  • ท้ายที่สุดแล้ว การออกแบบตัวโมเดล/อัลกอริทึม AI เองให้เหมาะกับฮาร์ดแวร์เป็นประเด็นที่สำคัญมาก ซึ่งต้องอาศัยความเข้าใจเชิงลึกเกี่ยวกับคุณลักษณะของฮาร์ดแวร์เป็นพื้นฐาน และ ThunderKittens ก็น่าจะช่วยให้นักพัฒนาได้มุมมองเชิงลึกเหล่านั้น
  • คาดหวังได้ว่างานวิจัยและการมีส่วนร่วมกับโอเพนซอร์สอย่างต่อเนื่องของ Hazy Research จะช่วยกระตุ้น ecosystem ของ CUDA ได้มาก แต่ในระยะยาวก็ดูเหมือนว่ายังจำเป็นต้องมีเฟรมเวิร์กระดับ abstraction ที่สูงกว่านี้ด้วย

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

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

 
GN⁺ 2024-05-13
ความคิดเห็นจาก Hacker News

  • ความต้องการของฮาร์ดแวร์ AI กำลังชัดเจนขึ้นเรื่อย ๆ เดิมที GPU ถูกออกแบบมาเพื่อวัตถุประสงค์อื่น แต่ถูกนำมาใช้กับ AI เพราะมีฮาร์ดแวร์สำหรับการคูณเมทริกซ์ที่ดีอยู่แล้ว "AI GPU" อาจตัดความสามารถบางอย่างของ GPU แบบดั้งเดิมออกได้ และแนวโน้มกำลังมุ่งไปสู่การใช้ตัวเลขที่สั้นลงกว่าเดิม (16 บิต, 8 บิต, 2 บิต, 1 บิต floating point) งานวิจัยฉบับนี้ชี้ว่าฮาร์ดแวร์ที่ชอบไทล์ขนาด 16x16 มีข้อได้เปรียบมากมาย

  • จำเป็นต้องมีตัวประมวลผลเสริมเฉพาะทางสำหรับ AI (NPU) โดยเฉพาะในระบบเดสก์ท็อประดับ prosumer สำหรับนักพัฒนา ผู้เชี่ยวชาญ และเกมเมอร์ GPU ใช้งานได้ในฝั่งองค์กร แต่ไม่สะดวกเมื่อนำมาใช้กับ AI ในบริบทของคอมพิวเตอร์ส่วนบุคคล โดยเฉพาะข้อจำกัดของ VRAM และการขาด API แบบเปิดมาตรฐานนอกเหนือจาก Vulkan

  • หากต้องการผลักดันงานวิจัย AI ให้ก้าวหน้า เราจำเป็นต้องศึกษาประสาทวิทยาและจิตวิทยาให้ดีขึ้น นอกจากนี้ สิ่งที่เกี่ยวข้องกับกราฟโทโพโลยีของโครงข่ายประสาทก็อาจมีส่วนเกี่ยวข้องด้วย

  • นี่คือ CUTLASS เวอร์ชันที่ทำให้ใช้งานเป็นมิตรกับผู้ใช้มากขึ้นหรือ?

  • มาสคอตของ ThunderKittens ให้อารมณ์แบบแมว/ Sony Aibo ดูเหมือนสร้างด้วย AI ได้ดีทีเดียว

  • ถ้า Universal Basic Compute (UBC) ถูกพิจารณาให้เป็นสิ่งทดแทน Universal Basic Income อนาคตแบบนั้นคงดิสโทเปียมาก ลองนึกภาพว่ามีบริษัทเดียวอย่าง Nvidia เป็นผู้สร้างระบบคอมพิวต์ทั้งหมด