การพัฒนาสมรรถนะของ GPU แบบก้าวกระโดด

 (hazyresearch.stanford.edu)
1 คะแนน โดย GN⁺ 2024-05-13 | 1 ความคิดเห็น | แชร์ทาง WhatsApp
  • ในสถานการณ์ที่ต้นทุนการประมวลผล AI สูงขึ้น Hazy Research สรุปว่าหัวใจของการปรับแต่งสมรรถนะ GPU คือการ ทำให้ tensor core ของ NVIDIA H100 ทำงานต่อเนื่องโดยไม่ว่าง
  • H100 ทำได้ 989 TFLOPs สำหรับการคูณเมทริกซ์แบบ half-precision แต่การคำนวณทั่วไปทำได้เพียงราว 60 TFLOPs ดังนั้นทันทีที่ tensor core หยุด อัตราการใช้งานก็จะลดลงอย่างมาก
  • หากต้องการเข้าใกล้สมรรถนะสูงสุด จำเป็นต้องจัดการทั้ง WGMMA, การจัดวาง shared memory, การสร้าง address และ occupancy ร่วมกัน และหากไม่มี wgmma.mma_async ในไมโครบेंชมาร์กจะทำได้เพียงราว 63% ของค่าสูงสุด
  • CUDA embedded DSL ที่เปิดเผยสู่สาธารณะชื่อ ThunderKittens ใช้ abstraction ของ tile และ vector เพื่อซ่อนความซับซ้อนอย่าง swizzling และ register layout ทำให้การเขียนเคอร์เนลตระกูล FlashAttention ง่ายขึ้น
  • เคอร์เนล forward ของ FlashAttention-2 สำหรับ H100 เขียนได้ด้วยโค้ดเพียงราว 100 บรรทัด และเร็วกว่า FlashAttention-2 ประมาณ 30% ขณะที่เคอร์เนล Based linear attention ทำงานที่ 215 TFLOPs

เงื่อนไขที่กำหนดสมรรถนะของ H100

  • AI ใช้ คอมพิวต์ จำนวนมาก และ Hazy Research ได้ทำงานในช่วงไม่กี่ปีที่ผ่านมาเพื่อให้ AI ใช้คอมพิวต์น้อยลง หรือทำงานได้มีประสิทธิภาพมากขึ้นภายใต้คอมพิวต์ที่มีอยู่
  • เป้าหมายเชิงปฏิบัติคือการสรุปสิ่งที่เรียนรู้จากการทำให้ GPU เร็วขึ้น และเปิดเผย CUDA embedded DSL ThunderKittens ที่ช่วยให้เขียนเคอร์เนลเร็วได้ง่ายขึ้น
  • ในภาพกว้างกว่านั้น ยังกล่าวถึงว่าความเข้าใจฮาร์ดแวร์ได้เปลี่ยนวิธีมอง AI compute อย่างไร

โครงสร้างของ NVIDIA H100 และคอขวด

  • การอธิบายนี้อ้างอิงจากการตั้งค่าของ GPU H100 SXM ดังนี้
    • 80GB HBM3, แบนด์วิดท์ 3TB/s
    • แคช L2 50MB, แบนด์วิดท์ 12TB/s แบ่งเป็น 2 ส่วน ส่วนละ 25MB ครอบคลุมทั้ง GPU และเชื่อมกันด้วย crossbar

  • 132 SM

    • แต่ละ SM มีแคช L1 ขนาด 256KB ที่รวม shared memory สูงสุด 227KB และรวมกันให้แบนด์วิดท์ราว 33TB/s
    • ฮาร์ดแวร์ใหม่ของ Hopper อย่าง Tensor Memory Accelerator(TMA) รับหน้าที่สร้าง address แบบอะซิงก์และดึงข้อมูลจากหน่วยความจำ
    • แต่ละ SM ประกอบด้วย 4 quadrant และแต่ละ quadrant มี warp scheduler, vector register 512 ตัว, tensor core สำหรับการคูณเมทริกซ์ และคำสั่งภายในแบบขนาน
    • การคำนวณทั้งหมดเกิดขึ้นบน SM และส่วนใหญ่ประมวลผลใน register
    • กุญแจสำคัญของการดึงสมรรถนะบน H100 คือการทำให้ tensor core มีข้อมูลป้อนไหลเข้าอย่างต่อเนื่อง
    • H100 ให้สมรรถนะ 989 TFLOPs สำหรับการคูณเมทริกซ์แบบ half-precision และราว 60 TFLOPs สำหรับการคำนวณ “อื่น ๆ”
    • ใน cycle ที่มีการใช้ tensor core จะไปถึงอัตราการใช้ฮาร์ดแวร์อย่างน้อย 94%
    • ใน cycle ที่ไม่ได้ใช้ tensor core จะอยู่ได้สูงสุดเพียง 6%

WGMMA: จำเป็นแต่ใช้งานยาก

  • H100 มีคำสั่ง warp group matrix multiply accumulate ชื่อ wgmma.mma_async
    • ใน PTX คือ wgmma.mma_async
    • ใน SASS คือ HGMMA / IGMMA / QGMMA / BGMMA
  • คำสั่ง wmma.mma.sync, mma.sync บน GPU รุ่นก่อนหน้าเป็นแบบซิงก์ โดย warp หนึ่งชุดที่มี 32 thread จะป้อนข้อมูลเข้า tensor core และรอผลลัพธ์
  • wgmma.mma_async ทำให้ 128 thread ที่ต่อเนื่องกันซิงก์ร่วมกันข้ามทุก quadrant ของ SM และเริ่มการคูณเมทริกซ์แบบอะซิงก์จาก shared memory ได้โดยตรง
    • ระหว่างที่การคูณเมทริกซ์กำลังทำงาน warp สามารถไปทำงานอื่นใน register ได้
    • และสามารถรอผลลัพธ์ในเวลาที่ต้องการได้
  • ในไมโครบेंชมาร์ก คำสั่งนี้จำเป็นต่อการดึง compute ทั้งหมดของ H100 ออกมา
    • หากไม่ใช้ จะพบว่า GPU อยู่ที่เพียงราว 63% ของการใช้สมรรถนะสูงสุด
    • อาจเป็นเพราะ tensor core ต้องการฮาร์ดแวร์ pipeline ที่ลึก แม้กระทั่งกับทรัพยากรภายในเครื่องเอง
  • ปัญหาใหญ่ที่สุดคือความซับซ้อนของ memory layout
    • shared memory layout แบบไม่ swizzle มี coalescing แย่มากและต้องใช้แบนด์วิดท์ L2 สูง
    • เอกสารของ layout แบบ swizzle มีข้อผิดพลาด ทำให้ต้องใช้เวลาพอสมควรในการทำความเข้าใจ
    • layout แบบ swizzle ดูเหมือนจะทำงานได้เฉพาะกับเมทริกซ์บาง shape และเข้ากันได้ไม่ดีกับความสามารถอื่นของ wgmma.mma_async
    • ฮาร์ดแวร์สามารถทำ sub-matrix transpose ระหว่างทางไปยัง tensor core ได้ แต่ทำได้เฉพาะตอนที่ layout ไม่ได้เป็นแบบ swizzle
  • ในเคอร์เนลอย่าง FlashAttention, TMA และแคช L2 เร็วพอที่จะซ่อนปัญหานี้ได้บางส่วน
  • หากต้องการใช้ฮาร์ดแวร์ได้เต็มที่ ต้องทำให้ memory request coalescing และหลีกเลี่ยง bank conflict จึงทำให้การควบคุม layout เป็นเรื่องสำคัญ

Shared memory และ bank conflict

  • single-access latency ของ shared memory ดูเหมือนจะอยู่ที่ราว 30 cycles ซึ่งในช่วงเวลานี้ tensor core ของ SM สามารถทำการคูณเมทริกซ์จัตุรัส 32x32 ได้เกือบสองครั้ง
  • งานก่อนหน้าอย่าง FlashAttention มุ่งเน้นไปที่คอขวด HBM-SRAM เป็นหลัก และในอดีตคอขวดนี้ก็สำคัญจริง
  • เมื่อ HBM เร็วขึ้น และ tensor core เติบโตเร็วกว่าองค์ประกอบอื่นของชิป แม้ latency เล็กน้อยของ shared memory ก็กลายเป็นสิ่งที่ต้องกำจัดหรือซ่อน
  • shared memory ถูกแบ่งเป็น 32 bank ดังนั้นหากไม่ระวังจะเกิด bank conflict
    • หากมีการร้องขอหลายชิ้นข้อมูลคนละชิ้นไปยัง memory bank เดียวกันพร้อมกัน คำขอจะถูกทำให้เป็นลำดับ
    • จากประสบการณ์ เคอร์เนลอาจช้าลงอย่างไม่สมดุล
    • register layout ที่คำสั่ง WGMMA และ MMA ต้องการ อาจเจอ bank conflict ได้หากใช้อย่างตรงไปตรงมา
  • วิธีแก้คือการจัดวาง shared memory ใหม่ด้วยรูปแบบ swizzling หลายแบบเพื่อหลีกเลี่ยง conflict
  • หากเป็นไปได้ ควรหลีกเลี่ยงการย้ายข้อมูลระหว่าง register กับ shared memory และเมื่อจำเป็น ก็ควรใช้ฮาร์ดแวร์ในตัวอย่าง WGMMA และ TMA เพื่อย้ายข้อมูลแบบอะซิงก์
  • การย้ายข้อมูลแบบซิงก์ด้วย warp จริงเป็นวิธีที่พบได้ทั่วไปที่สุด แต่ก็ใกล้เคียงกับ fallback ที่แย่ที่สุด

การสร้าง address และ TMA

  • H100 มีทั้ง tensor core และหน่วยความจำที่เร็วมาก จนการสร้าง memory address สำหรับดึงข้อมูลเองก็ใช้ทรัพยากรของชิปไปไม่น้อย
    • และจะยิ่งชัดเจนเมื่อมีรูปแบบ interleaved หรือ swizzling ที่ซับซ้อนเพิ่มเข้ามา
  • Tensor Memory Accelerator(TMA) ของ NVIDIA เปิดให้ระบุ layout แบบหลายมิติของ tensor ใน global/shared memory แล้วดึง subtile ของ tensor นั้นแบบอะซิงก์ พร้อม trigger barrier เมื่อเสร็จสิ้น
  • TMA ช่วยลดต้นทุนของการสร้าง address และทำให้จัด pipeline ได้ง่ายขึ้น
  • TMA ถูกมองว่าเป็นองค์ประกอบจำเป็นต่อการดึงศักยภาพของ H100 ออกมา เช่นเดียวกับ wgmma.mma_async
    • จากประสบการณ์ มันอาจสำคัญยิ่งกว่า WGMMA ด้วยซ้ำ
    • มันช่วยประหยัดทั้งทรัพยากร register และ instruction dispatch
    • และยังมีความสามารถทำ asynchronous reduction ไปยัง global memory ซึ่งมีประโยชน์ใน backward kernel ที่ซับซ้อน
  • TMA เองก็ต้องอาศัยการ reverse engineering บางส่วนเพื่อเข้าใจโหมด swizzling แต่ก็ยังเจ็บปวดน้อยกว่า WGMMA

Occupancy กับต้นทุนที่มันช่วยซ่อน

  • ใน CUDA, occupancy หมายถึงจำนวน thread ที่ถูก co-schedule บนฮาร์ดแวร์ประมวลผลเดียวกัน
  • warp scheduler ของแต่ละ SM quadrant จะพยายาม issue instruction ให้ warp ที่พร้อมรับคำสั่งในทุก cycle
  • H100 มีลักษณะที่พึ่งพา occupancy น้อยกว่ารุ่นก่อนหน้า
    • ด้วยความสามารถแบบอะซิงก์ แม้ instruction stream เดียวก็ทำให้ memory fetch, matrix multiply, shared memory reduction และ register math ทำงานพร้อมกันได้
  • แต่ occupancy ยังมีประโยชน์มากในการซ่อนข้อผิดพลาดและต้นทุนการซิงก์
    • pipeline ที่ออกแบบได้สมบูรณ์อาจเร็วได้โดยไม่ต้องเพิ่ม occupancy
    • จากสิ่งที่สังเกต NVIDIA GPU ดูเหมือนถูกออกแบบโดยคำนึงถึง occupancy
    • เพราะมีทั้งการซิงก์และโอกาสผิดพลาดจำนวนมาก การเพิ่ม occupancy จึงมักช่วยให้อัตราการใช้ฮาร์ดแวร์จริงดีขึ้น
  • บน H100 occupancy ยังมีประโยชน์ในระดับหนึ่ง แต่บน A100 และ RTX 4090 ถือว่ายิ่งสำคัญกว่า
    • มีการตั้งข้อสังเกตว่าอาจเป็นเพราะพึ่งพา instruction dispatch แบบซิงก์มากกว่า H100

ThunderKittens: DSL ขนาดเล็กภายใน CUDA

  • ThunderKittens คือ CUDA embedded DSL ที่สร้างขึ้นเพื่อให้เขียนเคอร์เนลเร็วบน H100 ได้ง่ายขึ้น
  • เดิมทีสร้างขึ้นเพื่อใช้งานภายในห้องวิจัย ก่อนจะเปิดเผยสู่สาธารณะในภายหลัง
  • ชื่อนี้มาจากความคิดที่ว่า kittens น่ารัก และการได้พิมพ์ kittens:: ในโค้ดก็สนุกดี
  • ThunderKittens มุ่งเน้นความเรียบง่าย และให้ type แบบ template อยู่ 4 แบบ
    • Register tiles: tensor 2 มิติบน register file
    • Register vectors: tensor 1 มิติบน register file
    • Shared tiles: tensor 2 มิติใน shared memory
    • Shared vectors: tensor 1 มิติใน shared memory
  • Tile ถูก parameterize ด้วยความสูง ความกว้าง และ layout
  • Register vector ถูก parameterize ด้วยความยาวและ layout ส่วน shared vector ใช้แค่ความยาว
    • โดยทั่วไป shared vector จะไม่เจอ bank conflict
  • operation ที่มีให้จะจัดการ tile และ vector ในระดับ warp หรือระดับ warp group แบบร่วมมือกัน
    • initializer: เช่นการทำ shared vector ให้เป็นศูนย์
    • unary op: เช่น exp
    • binary op: เช่น mul
    • row/column op: เช่น row_sum
  • ThunderKittens ฝังอยู่ใน CUDA ดังนั้นไม่เหมือนไลบรารีอย่าง Triton ตรงที่ abstraction สามารถ “ล้มเหลวอย่างนุ่มนวล” ได้
    • หากมีความสามารถที่ขาดไป ก็สามารถขยายเพิ่มในแบบที่ต้องการได้

ตัวอย่าง FlashAttention และสมรรถนะ

  • มีการยกตัวอย่าง ThunderKittens ด้วยเคอร์เนล forward ของ FlashAttention แบบง่ายสำหรับ RTX 4090
    • รองรับเฉพาะ headdim=64
    • n ต้องเป็นพหุคูณของ 256
    • เขียนด้วยโค้ด CUDA เพียงราว 60 บรรทัด
    • อัตราการใช้ฮาร์ดแวร์อยู่ที่ 75%
    • ความซับซ้อนส่วนใหญ่ไม่ได้อยู่ที่รูปแบบ swizzling หรือ register layout แต่อยู่ที่ตัวอัลกอริทึมเอง
  • forward pass ของ FlashAttention-2 สำหรับ H100 ก็ถูกเขียนด้วย ThunderKittens เช่นกัน
    • ThunderKittens ช่วยซ่อนความซับซ้อนของ TMA, WGMMA, โหมด swizzling และ descriptor
    • เคอร์เนลมีความยาวราว 100 บรรทัด
    • บน H100 เร็วกว่า FlashAttention-2 ราว 30%
  • ThunderKittens ถูกอธิบายว่าเหมือน “mini-pytorch” สำหรับ GPU ที่ห่อหุ้ม layout และ instruction พร้อมให้ primitive ต่าง ๆ
  • ยังมีการเปิดเผยเคอร์เนลสำหรับ Based linear attention และสถาปัตยกรรมอื่นที่จะตามมา
    • เคอร์เนล Based linear attention ทำงานที่ 215 TFLOPs
    • และหากคำนึงถึง recompute ของตัวอัลกอริทึมเอง จะเกิน 300 TFLOPs
    • ในทางทฤษฎี linear attention มีประสิทธิภาพสูงกว่า แต่ในฮาร์ดแวร์จริงในอดีตกลับมีประสิทธิภาพต่ำกว่ามาก
    • ผลลัพธ์นี้อาจช่วยขยายขอบเขตของแอปพลิเคชันที่ต้องการ throughput สูง

การคิดแบบยึด tile เป็นศูนย์กลาง

  • มีมุมมองว่าเหตุผลที่ ThunderKittens ทำงานได้ดี เพราะมันไม่ได้พยายามทำทุกอย่าง
    • CUDA มีพลังในการแสดงออกสูงกว่า ThunderKittens มาก
    • ThunderKittens เป็น DSL ขนาดเล็กและเรียบง่าย
  • abstraction หลักคือ small tile ซึ่งถูกมองว่าสอดคล้องกับทิศทางที่ AI และฮาร์ดแวร์กำลังมุ่งไป
  • ThunderKittens ไม่รองรับมิติที่เล็กกว่า 16
    • เพราะมองว่าฮาร์ดแวร์เองก็ไม่ได้ต้องการมิติเล็กขนาดนั้นเป็นพิเศษ
    • มีการตั้งคำถามในทำนองว่า “ถ้า matrix multiply เล็กกว่า 16x16 จะยังมั่นใจได้หรือว่านั่นคือ AI”
  • มุมมองแบบยุค CPU ที่เห็น 32-bit word เป็น register นั้นไม่เหมาะกับฮาร์ดแวร์ AI
    • vector register ขนาด 1024-bit ของ CUDA ถูกมองว่าเป็นอีกก้าวในทิศทางที่ถูกต้อง
    • ในบริบทนี้ register คือข้อมูลของ tile ขนาด 16x16
  • AI ยังคงมี matrix multiply, reduction และ reshape เป็นแกนกลาง ดังนั้น abstraction แบบ tile จึงเหมาะทั้งกับ AI และฮาร์ดแวร์
  • ต่อจากนี้จำเป็นต้องจัดเรียงไอเดีย AI ใหม่ให้แมปกับฮาร์ดแวร์ได้ดี
    • recurrent state ควรมีขนาดใหญ่พอที่จะใส่ใน SM ได้
    • compute density ต้องไม่ต่ำกว่าระดับที่ฮาร์ดแวร์ต้องการ
    • การนำสิ่งที่เรียนรู้จากฮาร์ดแวร์มาปรับเข้ากับการออกแบบ AI คือทิศทางสำคัญในอนาคต

แผนรองรับ AMD

  • การรองรับ AMD hardware ของ ThunderKittens จะมาในเร็ว ๆ นี้

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

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

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

  • คำถามที่ว่า “ถ้าการคูณเมทริกซ์เล็กกว่า 16x16 แล้วแน่ใจหรือว่านั่นคือ AI จริง ๆ?” น่าสนใจ
    ข้อกำหนดของฮาร์ดแวร์ AI กำลังชัดเจนขึ้นเรื่อย ๆ เดิมที GPU ถูกออกแบบมาเพื่อการใช้งานที่ต่างไปโดยสิ้นเชิง แต่ถูกนำมาใช้กับ AI เพราะฮาร์ดแวร์สำหรับการคูณเมทริกซ์ทำได้ดี และ “AI GPU” ก็อาจตัดฟังก์ชันบางอย่างที่มีอยู่ใน GPU จริงออกได้
    รูปแบบการแทนค่าตัวเลขก็มีแนวโน้มสั้นลงเรื่อย ๆ เช่น เลขทศนิยม 16 บิต, 8 บิต, 2 บิต, 1 บิต และสักวันหนึ่งก็คงจะมีจุดที่เหมาะสมถูกกำหนดขึ้น บทความนี้แสดงให้เห็นว่าฮาร์ดแวร์ที่ชอบ ไทล์ 16x16 นั้นค่อนข้างสมเหตุสมผล ตอนนี้คงมีใครบางคนกำลังเขียนสิ่งแบบนี้ด้วย VHDL อยู่แล้ว หรือมีแนวโน้มว่าจะทำในไม่ช้า
    ท้ายที่สุดดูเหมือนว่าจะมีอุปกรณ์ที่เรียบง่ายกว่า อเนกประสงค์น้อยกว่า และราคาถูกกว่า ซึ่งทำเฉพาะการคำนวณ “AI” ได้มากที่สุดโดยไม่ต้องแบกภาระฮาร์ดแวร์ที่ไม่จำเป็น

    • GPU ได้วิวัฒนาการมาเป็น เครื่อง AI ที่ไร้ส่วนเกินเท่าที่เป็นไปได้อยู่แล้ว อย่างน้อยตั้งแต่ Nervana ก่อตั้งในปี 2014 ก็มีคนพูดว่า GPU เป็นเทคโนโลยีเก่าและไม่เหมาะกับ AI แต่ดูเหมือนพวกเขาไม่ได้คาดว่า GPU จะวิวัฒนาการเป็นเครื่อง AI ได้รวดเร็วขนาดนี้
    • Apple เดินไปในทิศทางนี้มาหลายปีแล้ว NPU บนไดนั้นแตกต่างจาก GPU หรือ CPU อย่างสิ้นเชิง[1]
      Nvidia ก็คงกำลังทำอยู่เช่นกัน แต่ในเชิงธุรกิจ การคงรูปแบบอุปกรณ์ที่รวมเกม/ความบันเทิง/คริปโตเคอร์เรนซี/AI ไว้ด้วยกัน หรือก็คือการ์ดจอ อาจเป็นทางเลือกที่ดีกว่า
      [1] https://github.com/hollance/neural-engine/blob/master/docs/a...
    • ตอนที่ว่า “คำโกหกของ NVIDIA” แสดงให้เห็นความลึกของการแข่งขัน ข้อผิดพลาดในเอกสารคงไม่ใช่เรื่องบังเอิญล้วน ๆ และไดอะแกรมก็ขโมยหรือคัดลอกได้ง่าย ดังนั้น Nvidia อาจมีประโยชน์บางอย่างจากการจงใจปล่อยไว้
      ทำให้นึกถึงช่วงที่ Nervana ของ Naveen Rao สร้าง ไดรเวอร์ Nvidia Maxwell ที่เร็วกว่าไดรเวอร์ของ Nvidia เอง ไม่ใช่ว่าความผิดพลาดในเอกสารของผลิตภัณฑ์ที่เติบโตอย่างรวดเร็วทั้งหมดจะเป็นมาตรการตอบโต้คู่แข่ง แต่เมื่อพิจารณาว่านักวิจัยใช้เวลานานในการทำวิศวกรรมย้อนกลับ wgmma และยังมีสถานการณ์การเมืองสหรัฐฯ-จีนรอบ H100 ด้วย ก็ทำให้ดูเหมือน Nvidia กำลังใช้วิธีเก่า ๆ เพื่อรักษาคูเมืองของตน
      ดังนั้นแทนที่จะเจาะลึกความเฉพาะตัวของ H100 มากเกินไป ควรมองว่า “ฮาร์ดแวร์ที่ AI ต้องการคืออะไร” นั้นรวมถึงสถานการณ์เชิงพาณิชย์ด้วย
    • AMD มาถึงรุ่นที่ 2 ของไลน์ Versal แล้ว
      https://www.amd.com/en/products/accelerators/alveo/v80.html
      XDNA Architecture
      https://www.amd.com/en/technologies/xdna.html
    • Google ก็ทำอุปกรณ์แบบนี้มาเกือบ 10 ปีแล้วไม่ใช่หรือ?

  • ตอนที่ว่า “คำโกหกของ NVIDIA นี่เป็นคำอธิบายที่ทำให้เข้าใจผิดอย่างมหาศาลเกี่ยวกับ layout แบบ 128b swizzled wgmma ของจริง เพราะไดอะแกรมนี้ทำให้ผมเสียเวลาในชีวิตไป 3 สัปดาห์แบบเอาคืนไม่ได้ เลยขอประจานต่อสาธารณะ” น่าประทับใจ
    ผมสงสัยว่าจะมีใครแปลกใจไหมที่ความก้าวหน้าของ AI ส่วนใหญ่อย่างมหาศาลอยู่ในงานวิศวกรรมอย่าง การปรับแต่งการคูณเมทริกซ์ให้เหมาะสม และงานวิศวกรรมจำนวนมากนั้นก็คือ การทำวิศวกรรมย้อนกลับชิป NVIDIA

    • ตัวสถาปัตยกรรมเองไม่ได้สร้างความแตกต่างมากนัก ถ้าฝึกโมเดลที่ใหญ่พอด้วยข้อมูลที่ใหญ่พอ ผลลัพธ์มักจะออกมาคล้ายกันไม่ว่าสถาปัตยกรรมจะเป็นแบบไหน ดังนั้นพัฒนาการส่วนใหญ่ของ AI ตอนนี้จึงมองได้ว่าเป็นผลจากการที่เราสามารถคูณเมทริกซ์ได้เร็วมาก

  • warp scheduler, 4 quadrant, tensor memory accelerator, layout แบบ unswizzled wgmma…
    เส้นแบ่งระหว่างศัพท์ GPU กับ technobabble สไตล์ Star Trek เริ่มพร่าเลือนขึ้นเรื่อย ๆ

    • ตอนอ่านบทความก็พอรู้อยู่บ้าง แต่ถ้าพูดว่า “กำลัง warp quadrant ด้วย tensor accelerator” มันให้ความรู้สึก Star Trek จริง ๆ
      เวลาอ่านบทความอื่น ๆ ก็เคยคิดแบบนี้เป็นครั้งคราวเหมือนกัน สงสัยว่าถ้ามีใครได้รับลิงก์บทความจากที่นี่แล้วเข้าไปอ่านจะรู้สึกยังไง น่าจะเหมือนเดินเข้าไปในงานรวมตัวแฟน Trek ที่กำลังถกกันเรื่อง warp core
    • พอเห็นประโยคนี้ก็ถอยออกมามองศัพท์เหล่านี้ด้วยสายตาใหม่ และมันก็จริงมากจนขำ

  • หากต้องการลดการใช้พลังงานของการอนุมาน AI และเพิ่มความเร็ว น่าจะดีที่สุดคือเปลี่ยนไปใช้ วงจรประมาณค่าแบบแอนะล็อก
    ไม่ได้ต้องการการคูณและบวกเลขทศนิยมลอยตัวที่สมบูรณ์แบบ แค่ต้องการอุปกรณ์ที่รับแรงดันอินพุตสองค่า แล้วให้แรงดันเอาต์พุตที่ใกล้เคียงผลลัพธ์ของการคูณมากพอเท่านั้น

    • ผมรู้จักคนที่ทำงานในทิศทางนี้ และได้ยินมาว่าอุปสรรคใหญ่คือจะสร้างสิ่งที่ทำ ตรรกะแบบแอนะล็อก ได้ด้วยเทคโนโลยีการผลิตชิปที่มีอยู่ได้อย่างไร, การออกแบบที่ไม่ทำตัวเหมือนเสาอากาศ, และความเป็นไปได้ที่ต้องปรับจูนโมเดลที่จะรันแยกตามชิปแต่ละตัว เพราะค่าคลาดเคลื่อนจากการผลิตต่างกันในชิปจริงแต่ละชิ้น
      ข้อดีใหญ่คือแทนที่จะแทน float16 ด้วยสาย 16 เส้น ก็แทนตัวเลขนั้นด้วยแรงดันบนสายเส้นเดียว ตามทฤษฎีอาจให้ความแม่นยำสูงกว่า float32 มากได้ด้วย อีกทั้งยังเชื่อมต่อค่าโดยตรงได้โดยไม่ต้องโหลดเข้า arithmetic logic unit ทำให้มีศักยภาพลดพื้นที่ไดและพลังงานได้หลายลำดับขั้น
    • ผมมองว่าวงจรแอนะล็อกยังอีกไกลกว่าจะมีประโยชน์ในทางปฏิบัติ แต่ที่ที่ยอมรับความไม่ถูกต้องได้ อาจเป็น วงจรดิจิทัลที่มีสัญญาณรบกวน
      เช่น ยอมรับว่าบิตเอาต์พุตหนึ่งในล้านบิตอาจกลับค่า เพื่อแลกกับการปรับปรุงอัตราส่วนประสิทธิภาพ/พลังงาน เรื่องนี้อาจยากใน float32 ที่ค่า infinity เพียงค่าเดียวอาจทำให้ทุกอย่างพัง แต่ใน int8 ถ้าอยากได้ 0 แล้วบางครั้งได้ 128 ก็น่าจะพอทนได้
      [1] ไม่แน่ใจจริง ๆ ว่ายูนิตเมทริกซ์เลขทศนิยมลอยตัวของ H100 ปฏิบัติตาม IEEE 754 หรือไม่
    • ถ้าก้าวไปอีกขั้น ผมคิดว่าเราต้องมีอะไรบางอย่างที่คล้ายกับวิธีทำงานของ สมองชีวภาพ จริง ๆ แต่ผลิตได้ง่าย
      โครงข่ายประสาทชีวภาพไม่ได้ใกล้เคียงกับการเชื่อมต่อแบบเต็มเหมือนโครงข่ายประสาทเทียมทั่วไป และจำนวนการเชื่อมต่ออินพุต/เอาต์พุตของนิวรอนมีค่าน้อยกว่า 10 จึงมีความเป็นท้องถิ่นสูงมาก เท่าที่เรารู้ ในชีววิทยาไม่มี backpropagation แต่มี feedback และการวนกลับแทน
      อาจยังมีเซลล์หรือกระบวนการสนับสนุนที่จำเป็นต่อการทำงานของระบบประสาทส่วนกลางซึ่งเรายังไม่รู้จักอยู่ก็ได้ แม้ในระดับสูงก็มีความเป็นไปได้ว่าจะมีการเชื่อมต่อที่ “ฮาร์ดโค้ด” อยู่จำนวนมาก และบางส่วนก็เป็นที่รู้กันแล้ว ตัวอย่างเช่น นิวรอนการได้ยินในหูเชื่อมต่อกัน และมีสิ่งที่คล้าย convolution เกิดขึ้นเพื่อระบุตำแหน่งของเสียง นี่ไม่ใช่ปรากฏการณ์อุบัติใหม่ แต่เป็นฟังก์ชันที่ทำได้แม้ไม่ผ่านการฝึก
      ไม่แปลก เพราะสิ่งมีชีวิตค้นพบสิ่งนี้ผ่านเวลาหลายพันล้านปีและจำนวนรุ่นที่ใกล้เคียงกัน ในทางทฤษฎีซอฟต์แวร์ก็น่าจะทำได้ แต่เมื่อพิจารณานิวรอนมากกว่า 1 ล้านล้านตัวในสมองไพรเมต/มนุษย์ แม้เครื่องระดับพันคอร์ในปัจจุบันก็ยังยากอย่างยิ่ง แม้จะเป็น “คลาวด์” ก็คงไม่สามารถตอบสนองการเชื่อมต่อและ latency ที่ต้องการได้
      ถ้าแนวทางแบบนี้สามารถจำลองได้สำเร็จในระดับหนอนหรือแมลงก็น่าจะยอดเยี่ยม
    • การตอบโจทย์ทั้ง ช่วงค่าและความแม่นยำ ให้เพียงพอพร้อมกันดูแทบเป็นไปไม่ได้
    • พูดตรง ๆ ดูเหมือนฝันร้ายสำหรับการดีบัก

  • บทความนี้ทำให้นึกถึงความสนุกที่เคยรู้สึกในวิชา CS 149 การเขียนโปรแกรมแบบขนาน อีกครั้ง

    • Kayvon กับ Kunle สุดยอดมาก ผมลง CS149 Parallel Programming เมื่อสองเทอมก่อน และมันดีจริง ๆ :)

  • สไตล์การเขียนของบทความนี้น่าประทับใจมาก และตั้งตารอที่จะได้เห็นสิ่งนี้บน AMD MI300x ถ้าอยากใช้เวลากับเครื่องของผมก็บอกได้

    • อยากรู้ว่าเคยทำงาน AI ด้วยผลิตภัณฑ์ AMD มาเยอะไหม ผมไม่อยากจ่ายมากกว่า 2500 ดอลลาร์ให้ RTX 4090 เลยกำลังพิจารณา RX 7900XTX สำหรับทดลองหรือเริ่มต้น
      อยากรู้ว่าจริง ๆ แล้วมันทำงานได้ดีแค่ไหน หรือควรเก็บเงินเพิ่มอีกหน่อยแล้วซื้อ XTX แทน 7900 XT หรือถ้า VRAM ลดลงจะกระทบการใช้งานจริงมากแค่ไหน
    • งานเขียนที่ดีควรชัดเจนและไม่กำกวม เวลาพูดสามารถขัดกลางคันเพื่อขอ clarification ได้ แต่ตัวหนังสือมีโอกาสส่งสารแค่ครั้งเดียว
      ผู้อ่านไม่ควรต้องไปเปิด knowyourmeme.com เพื่อเดาว่าผู้เขียนพยายามจะพูดอะไร ผมไม่รู้ด้วยซ้ำว่าชื่อนี้หมายความว่าอะไร และมองว่ามันหลุดเป้าหมายไปไกลขนาดนั้น
    • จริงเหรอ? ทำให้นึกถึง PTSD สมัย Wallstreetbets เลย

  • ถ้าจะเข้าใจบทความแบบนี้ให้ครบถ้วน ควรเริ่มจากตรงไหน และควรตาม โรดแมป แบบไหนดี

    • มีคอร์สที่ดีสำหรับเรียนการเขียนโปรแกรม GPU ประมาณบทเรียน 4.0 ก็น่าจะได้พื้นฐานที่จำเป็นแล้ว: https://youtube.com/playlist?list=PLzn6LN6WhlN06hIOA_ge6Srgd...
      แล้วก็ควรลองเขียน CUDA kernel สำหรับทำการคูณเวกเตอร์-เมทริกซ์ด้วยตัวเอง ถ้าใช้ pycuda ก็จะโฟกัสที่ kernel ได้ ส่วนที่เหลือเขียนด้วย Python ได้ บอก ChatGPT ว่าอยากสร้างการใช้งานที่คูณเวกเตอร์ 4000 องค์ประกอบกับเมทริกซ์ 4000x12000 ด้วยตัวเอง และให้ช่วยแนะนำตลอดกระบวนการก็ได้
      สำหรับการเช่า GPU Runpod ใช้ดี และตอนนี้มีตั้งแต่ GPU ราคาถูกไปจนถึง H100 เริ่มแรกใช้ GPU ระดับล่างก่อนก็พอ
    • ถ้าอยากลงลึก น่าจะดูเพลย์ลิสต์การคูณเมทริกซ์ของ Spiral: https://www.youtube.com/playlist?list=PL04PGV4cTuIWT_NXvvZsn...
      ผมใช้เวลา 2 เดือนในการทำและปรับแต่ง kernel การคูณเมทริกซ์ด้วย Spiral

  • กราฟใน GitHub README (https://github.com/HazyResearch/ThunderKittens/blob/main/att...) เวียนหัวมาก กราฟแท่งหยัก ๆ แบบนี้ถูกต้องตามกฎได้ด้วยเหรอ? :P

    • เห็นด้วย ดูเหมือนพยายามใส่ภาพลวงตาอะไรสักอย่าง ดูแค่ตัวเลขโดยไม่มีแท่งน่าจะดีกว่า
    • ดูเหมือน ธีม xkcd ของ matplotlib[1] ถึงอย่างนั้นก็เห็นด้วยว่าคลื่นมันเยอะเกินไป
      [1]: https://matplotlib.org/stable/gallery/showcase/xkcd.html#sph...

  • ชื่อ ThunderKittens ยอดเยี่ยมมาก อยากเห็น ThunderKittens จัดการกับ FlashAttention backward pass ซึ่งยากกว่า forward pass อยู่หนึ่งหลัก

  • งานวิจัยแบบนี้ไม่ใช่สิ่งที่ทีมที่สร้าง NPU ในปัจจุบันทำกันมาแล้วหรือ? เช่น ชิป Groq ทำประสิทธิภาพแบบตอนนี้ได้เพราะใช้สถาปัตยกรรมเฉพาะสำหรับ AI ฝั่งผู้บริโภค Apple Silicon ก็ถือว่าค่อนข้างมีความสามารถ
    ผมไม่ได้อยู่ในสายนี้ แต่คิดว่าน่าจะมีข้อจำกัดถ้าใช้แค่โปรเซสเซอร์ทั่วไปที่สื่อสารกันผ่านเส้นทางที่ค่อนข้างช้า การคิดการออกแบบใหม่ในระดับฮาร์ดแวร์ และท้ายที่สุดการลดราคาในตลาดผู้บริโภค น่าจะเป็นกลยุทธ์ระยะยาวที่ดีกว่า

    • ไม่ค่อยมั่นใจนักกับคำที่ว่า Apple Silicon ค่อนข้างมีความสามารถในฝั่งผู้บริโภค ถ้าดูซับเรดดิต localllama บน reddit จะเห็นโพสต์จาก ผู้ใช้ CPU จำนวนมากที่หงุดหงิดเพราะพยายามทำให้ได้ความเร็วที่ใช้งานได้จริง
      ด้วยเงินไม่กี่ร้อยดอลลาร์ก็ซื้อ Nvidia GPU ได้ หรือซื้อโน้ตบุ๊กเกมมิงที่มี 4050 VRAM 6GB ราคา 900 ดอลลาร์ได้ จึงยากจะเรียก AI บน CPU ว่ามีความสามารถ
      ที่ทำงานผมก็ไม่มี GPU เลยลองทำบน CPU แต่ในทางปฏิบัติแทบทำอะไรไม่ได้ นอกจากใช้โมเดลเล็ก ๆ แล้วรอ สุดท้ายก็ต้องขอคอมพิวเตอร์ที่มี GPU
      “ทำได้ในเชิงเทคนิค” กับ “ใช้งานจริงได้ดี” นั้นต่างกัน Nvidia ใช้งานดีมากจริง ๆ ส่วน CPU นั้นเจ็บปวดและน่าหงุดหงิด