คอมไพล์ LLM เป็น MegaKernel เพื่อทำให้การอนุมานมีความหน่วงต่ำ
(zhihaojia.medium.com)- มีการพัฒนาคอมไพเลอร์ที่แปลงการอนุมานของ LLM ให้เป็น เมกะเคอร์เนล เดียวโดยอัตโนมัติ
- แนวทาง MegaKernel(Persistent kernel) ทำให้สามารถรวมการคำนวณและการสื่อสารทั้งหมดในการอนุมาน LLM เข้าไว้ใน GPU kernel เดียว จึงทำให้เกิด ความหน่วงต่ำมาก
- มีข้อจำกัดว่าโครงสร้างแบบกระจายของเฟรมเวิร์ก ML หรือไลบรารีเคอร์เนลแบบเดิม ทำให้ การทำทั้ง pipeline ให้เป็นเคอร์เนลเดียว ทำได้ยากมาก
- Mirage Persistent Kernel(MPK) ใช้ คอมไพเลอร์และระบบรันไทม์ เพื่อแปลงการอนุมาน LLM แบบหลาย GPU ให้เป็น megakernel ประสิทธิภาพสูงโดยอัตโนมัติ
- MPK แปลง computation graph ให้เป็น task graph แบบละเอียด เพื่อเพิ่ม software pipelining และการซ้อนทับระหว่างการคำนวณกับการสื่อสารให้มากที่สุด
- เมื่อใช้ MPK จะ ลดความหน่วงในการสร้างโทเคน เมื่อเทียบกับระบบเดิม และยิ่งมีจำนวน GPU มากเท่าไร การเพิ่มขึ้นของประสิทธิภาพก็ยิ่งชัดเจน
ภาพรวมและข้อดีของแนวทาง MegaKernel
- ในการอนุมานของโมเดลภาษาขนาดใหญ่ (LLM) หนึ่งในวิธีที่มีประสิทธิภาพในการลด เวลาแฝง คือการหลอมรวมทุกขั้นตอนของการคำนวณและการสื่อสารให้เป็น megakernel(เคอร์เนลแบบต่อเนื่อง) เดียว
- วิธีนี้ทำให้การประมวลผลทั้งหมด ตั้งแต่การคำนวณรายเลเยอร์ของโมเดลไปจนถึงการสื่อสารระหว่าง GPU ดำเนินการอย่างต่อเนื่องโดย GPU kernel เดียว
- ข้อดีหลักมีดังนี้
- ตัดการเรียก kernel ซ้ำ ๆ ออกไป จึง กำจัด kernel launch overhead
- ทำ software pipelining ได้ตลอดทั้งเลเยอร์
- ดำเนินการคำนวณและการสื่อสารพร้อมกัน เพื่อซ่อนเวลาแฝง
ข้อจำกัดเดิมและการมาของ MPK
- เฟรมเวิร์ก ML เดิมอย่าง PyTorch, Triton, TVM ไม่ได้รองรับการสร้าง end-to-end megakernel แบบอัตโนมัติในระดับแก่นแท้
- ระบบ LLM จริงประกอบด้วยการผสมไลบรารีเคอร์เนลหลากหลาย เช่น NCCL/NVSHMEM(การสื่อสาร), FlashInfer/FlashAttention(แอ็ตเทนชัน), CUDA/Triton(การคำนวณแบบกำหนดเอง) ทำให้ การรวมเป็นเคอร์เนลเดียวทำได้ยาก
- จากฉากหลังนี้ นักวิจัยจาก CMU, UW, Berkeley, NVIDIA และ Tsinghua ได้พัฒนา Mirage Persistent Kernel(MPK) ขึ้น
- MPK ผสาน คอมไพเลอร์และรันไทม์ เพื่อแปลง pipeline การอนุมาน LLM ทั้งหมดให้เป็น megakernel ประสิทธิภาพสูงโดยอัตโนมัติ
คุณค่าหลักของ MPK
- MPK กำจัด kernel launch overhead ได้อย่างสมบูรณ์ และเพิ่มการซ้อนทับระหว่างการคำนวณ การโหลดข้อมูล และการสื่อสารข้ามชั้นให้สูงที่สุด จึงสร้างสภาพแวดล้อมการอนุมาน LLM ที่มีเวลาแฝงต่ำมาก
- ในการทดสอบจริง (พรอมป์ต์ 39 โทเคน, สร้าง 512 โทเคน, ไม่ใช้ speculative decoding)
- บน GPU NVIDIA A100 40GB เดี่ยว เมื่อเทียบกับเวลาแฝงการดีโค้ดต่อโทเคนของระบบที่ปรับแต่งมาแล้วอย่าง vLLM/SGLang ที่ 14.5ms, MPK ลดลงได้ถึง 12.5ms
- ตัวเลขนี้เข้าใกล้ขีดล่างเชิงทฤษฎีที่ 10ms (อ้างอิงแบนด์วิดท์หน่วยความจำ 1.6TB/s และการโหลดน้ำหนัก 16GB)
- ในสภาพแวดล้อม หลาย GPU การรวมการคำนวณและการสื่อสารเข้าด้วยกันอย่างสมบูรณ์ทำให้ความได้เปรียบด้านประสิทธิภาพของ MPK เด่นชัดยิ่งขึ้นเมื่อจำนวน GPU เพิ่มขึ้น
โครงสร้างการทำงานของ MPK แบบละเอียด
Part 1. คอมไพเลอร์ – แปลง LLM computation graph → task graph
- โดยทั่วไป การคำนวณของ LLM แสดงในรูป computation graph ที่แต่ละโหนดคือการคำนวณ (เช่น matrix multiplication, attention) หรือการสื่อสาร (เช่น all-reduce) และเส้นเชื่อมคือการพึ่งพาของข้อมูล
- ในการออกแบบแบบเดิม มักใช้วิธีรันเคอร์เนลแยกตามโอเปอเรเตอร์ แต่แนวทางนี้สะท้อนเพียง การพึ่งพาระดับเคอร์เนล ไม่ใช่ระดับข้อมูลที่ต้องพึ่งพาจริง จึง จำกัดโอกาสในการทำ pipelining
- ตัวอย่าง: หากมี allreduce ต่อจาก matrix multiplication จะต้องรอให้ matrix multiplication ทั้งหมดเสร็จก่อน allreduce จึงเริ่มได้ ทั้งที่จริงสามารถแบ่งข้อมูลย่อยออกมาเพื่อใช้การประมวลผลบางส่วนและความสัมพันธ์แบบบางส่วนได้
- คอมไพเลอร์ MPK จะแบ่ง computation graph ให้ละเอียดขึ้น และแปลงเป็น fine-grained task graph ที่เหมาะกับหน่วยข้อมูลจริงโดยอัตโนมัติ
- แต่ละ task (สี่เหลี่ยม) คือหน่วยการคำนวณ/การสื่อสารที่ถูกกำหนดให้กับ GPU SM แต่ละตัว
- แต่ละ event (วงกลม) คือจุดซิงโครไนซ์ระหว่าง task
- ใช้เส้นเชื่อมระหว่าง task และ event เพื่อแทนความสัมพันธ์ด้านข้อมูล/การควบคุมได้อย่างมีประสิทธิภาพ
- ด้วย task graph นี้ MPK จึงทำให้การคำนวณและการสื่อสารซ้อนทับกันได้มากขึ้น ทั้งแบบบางส่วนและแบบขนาน
- ยังใช้ Mirage kernel superoptimizer เพื่อสร้าง implementation แบบ CUDA ประสิทธิภาพสูงที่เหมาะกับแต่ละ task โดยอัตโนมัติ
Part 2. รันไทม์ – รัน task graph ภายในเมกะเคอร์เนล
- รันไทม์ของ MPK ใช้วิธี รัน task graph ทั้งหมดภายใน GPU kernel เดียว (megakernel)
- มีการแบ่ง SM(Streaming Multiprocessors) ทั้งหมดของ GPU แบบคงที่ให้ทำหน้าที่เป็น worker และ scheduler
Worker
- worker แต่ละตัวทำงานในระดับ SM และจัดการ task queue ของตนเองโดยเฉพาะ
- ทำงานแบบลูปดังนี้
- ดึง task ถัดไปจากคิว
- ดำเนินการ (เช่น matmul, attention, การส่งข้อมูล)
- เมื่อเสร็จแล้วแจ้ง event
- ทำซ้ำ
- วิธีนี้ช่วย เพิ่มประสิทธิภาพการใช้ทรัพยากรของแต่ละ worker และรองรับการคำนวณข้ามเลเยอร์แบบอะซิงโครนัส
Scheduler
- distributed scheduler ทำงานที่ระดับ warp เดียวภายในแต่ละ SM และสามารถรัน scheduler พร้อมกันได้สูงสุด 4 ตัว
- scheduler แต่ละตัวจะจัดการคิวของ event ที่ถูกเปิดใช้งาน และมอบหมาย task ที่มีเงื่อนไขพร้อมแล้วให้กับ worker
- ด้วยวิธีนี้จึงสามารถกระจายและประมวลผล task จำนวนมากได้โดย ไม่มี overhead จากการซิงโครไนซ์แบบรวมศูนย์
วิธีรันแบบ event-based
- เมื่อ task เสร็จสิ้น จะเพิ่มค่า event counter ที่เกี่ยวข้อง เมื่อเคาน์เตอร์ถึงค่าที่กำหนด event จะถูกเปิดใช้งานและถูกใส่เข้าไปในคิวของ scheduler
- scheduler จะรัน task ถัดไปที่มีความสัมพันธ์พึ่งพากับ event นั้น
- ด้วยเหตุนี้ fine-grained software pipelining และ การซ้อนทับระหว่างการคำนวณกับการสื่อสาร จึงเกิดขึ้นได้อย่างเป็นธรรมชาติ
- ตัวอย่าง: matmul ของเลเยอร์หนึ่งและ attention ของอีกเลเยอร์หนึ่งสามารถรันพร้อมกันได้
- สามารถเริ่มการสื่อสาร allreduce ได้ทันทีที่ได้ผลลัพธ์บางส่วนจาก matmul
- เนื่องจากการจัดตารางและการสลับ task ทั้งหมดเกิดขึ้นภายในคอนเท็กซ์ของ kernel เดียว overhead ระหว่าง task จึงต่ำมาก อยู่ที่ระดับ 1–2 ไมโครวินาที(μs)
ทิศทางในอนาคต
-
เป้าหมายของ MPK: ช่วยให้นักพัฒนาเพียงเขียนโค้ด Python จำนวนน้อย (ระดับไม่กี่สิบบรรทัด) ก็สามารถคอมไพล์ LLM ให้เป็น megakernel ได้ง่ายและดึงประสิทธิภาพสูงสุดออกมาได้
-
ทิศทางการพัฒนาหลัก
- รองรับสถาปัตยกรรม GPU รุ่นใหม่: เช่น NVIDIA Blackwell รวมถึงแนวทางที่ปรับจูนเฉพาะระดับ warp
- รองรับ workload แบบไดนามิก: วิจัยกลยุทธ์การคอมไพล์สำหรับโมเดลที่ต้องการ dynamic control flow เช่น mixture-of-experts(MoE)
- การจัดตาราง task ขั้นสูง: ศึกษาและผลักดันความเป็นไปได้ของนโยบายสมัยใหม่ เช่น แบบอิงลำดับความสำคัญ หรือการเพิ่มประสิทธิภาพ throughput
-
MPK ชี้ให้เห็น จุดเปลี่ยนเชิงรากฐานของวิธีคอมไพล์และรันงานอนุมาน LLM บน GPU และหวังขยายความร่วมมือกับชุมชนต่อไป
เอกสารเพิ่มเติม
- สามารถดูโค้ดและเอกสารของ MPK(Mirage Persistent Kernel) รวมถึงผลงานวิจัยล่าสุดได้ที่ GitHub(https://github.com/mirage-project/mirage)
1 ความคิดเห็น
ความคิดเห็นจาก Hacker News
ถึงผู้เขียน แนวทาง interpreter บน GPU โดยตรงดูเป็นทิศทางอนาคตที่น่าสนใจมาก มีงานวิจัยอีกชิ้นที่ใช้แนวทางแทบจะเหมือนกัน จึงแนะนำให้อ่านโพสต์ที่เกี่ยวข้อง ด้วย โมเดลการเขียนโปรแกรมพื้นฐานของ CUDA (เช่น kernel launch) กำลังถูกเลี่ยงเพื่อรองรับการขนานงานย่อยแบบละเอียด และผมก็ได้เห็นด้วยตาตัวเองว่าวิธีนี้ใช้ฮาร์ดแวร์ได้คุ้มค่ากว่า เลยอดสงสัยไม่ได้ว่า CUDA กำลังฉุดเราไว้ในหลายด้านหรือเปล่า หวังว่างานของผู้เขียนจะมีโอกาสเข้าไปเป็น experimental backend ของ PyTorch และขอทักเรื่องพิมพ์ผิดเล็กน้อยว่า สองย่อหน้าในช่วงแรกแทบจะเหมือนกันเลย
ผมทำงานใกล้ชิดกับ vLLM และ SGLang มาระยะหนึ่ง และมั่นใจว่าโปรเจกต์นี้คือภาพในอุดมคติของงานลำดับถัดไปพอดี การวิเคราะห์กราฟ dependency ของการคำนวณ รวมถึงการ fuse operation และจัดตาราง task ให้ฉลาดขึ้นนั้นน่าประทับใจมาก ขอแสดงความยินดีกับทีม
ผมลองอ่านทั้งบทความและ README บน github แล้ว รู้สึกว่าเป็นโปรเจกต์ที่ยอดเยี่ยมมาก เลยสงสัยว่าแนวทางการปรับแต่งแบบนี้จะนำไปใช้กับขั้นตอนฝึกได้ด้วยหรือไม่ โดยเฉพาะการ fuse งาน backward กับการสื่อสาร gradient ซึ่งน่าจะเป็นโจทย์ท้าทาย ผมเข้าใจว่าตอนนี้ยังไม่รองรับ dynamic workload (เช่น MoE) แต่ขอพูดถึงงานวิจัยล่าสุด FlashDMoE: Fast Distributed MoE in a Single Kernel ที่ประมวลผล MoE ในเคอร์เนลเดียว
ขอบคุณที่อ่านทั้งบทความและ README การรองรับขั้นตอนฝึกก็เป็นไปได้ แต่โดยทั่วไป training kernel มักมีขนาดใหญ่กว่า จึงทำให้ kernel launch overhead ไม่ใช่ปัญหาใหญ่เท่ากับฝั่ง inference ดังนั้นผู้ที่ได้ประโยชน์มากกว่าคือ inference โดยเฉพาะแบบ latency ต่ำ เราได้อ่านงาน FlashDMoE ที่คุณแชร์แล้วและคิดว่าน่าสนใจมาก พร้อมทั้งย้ำว่าการรองรับโมเดล MoE คือเป้าหมายถัดไปของเรา
โดยส่วนตัวผมค่อนข้างกังขากับการทุ่มเวลาไปกับการเพิ่มประสิทธิภาพการฝึกแบบอิง gradient เพราะงานฝึกจำนวนมากในโลกจริงมีลักษณะเป็นค่าดิสครีต ซึ่งผมคิดว่าแนวทางแบบ gradient-based จัดการได้ไม่ดีนัก
ขั้นต่อไปคงเป็นการคอมไพล์ตรงไปเป็น Verilog แล้วไปซื้อฮาร์ดแวร์ LLM จาก aliexpress มาใช้เอง นี่คือความฝันเลย
ขอแชร์บทความแนะนำเทคโนโลยีฮาร์ดแวร์อย่าง Chisel ก่อนยุค AI และ GPU ไอเดียการแปลงจากซอฟต์แวร์ไปเป็นฮาร์ดแวร์โดยตรงแบบนี้เคยเป็นแนวทางที่มีอนาคต การพัฒนา CPU เริ่มชะงัก และความต้องการจะปรับชั้นกลางระหว่างซอฟต์แวร์กับฮาร์ดแวร์ให้เหมาะขึ้นก็ยังมีต่อเนื่อง แต่เป็นไปได้สูงว่าการประมวลผลขนานสไตล์ GPU จะยังเป็นแนวทางเร่งความเร็วหลักต่อไป CPU ทั่วไปก็น่าจะเหลือบทบาทเป็นสมองเล็ก ๆ ที่คอยจัดการ GPU อย่างไรก็ตาม คาดว่าวิธีแปลงจากซอฟต์แวร์ไปเป็นฮาร์ดแวร์โดยตรงคงยากที่จะกลายเป็นกระแสหลัก
มีการคาดการณ์ว่าอีก 5~10 ปีข้างหน้า เมื่อโครงสร้างของ LLM มีเสถียรมากขึ้น การแมปลงฮาร์ดแวร์โดยตรงอาจใช้งานได้จริง ด้วยเทคโนโลยีปัจจุบัน มีโอกาสที่โมเดลระดับหลายหมื่นล้านพารามิเตอร์จะใส่ลงบนเวเฟอร์เดียวได้โดยใช้เพียง logic gate ความละเอียดต่ำมากระดับใกล้ 1.5 บิต ยิ่งต้องการความแม่นยำสูง จำนวน gate ก็ยิ่งเพิ่มแบบทวีคูณ ดังนั้นในตอนนี้การคงน้ำหนักไว้ในหน่วยความจำและแชร์หน่วยคำนวณยังมีประสิทธิภาพกว่า ในอนาคตการพัฒนา LLM ความละเอียดต่ำมากอาจกลายเป็นงานสำคัญที่ขาดไม่ได้
มุกตลกว่าต้นทุนการฝึกก็สูงอยู่แล้ว ถ้ายังต้องเพิ่มต้นทุนของ mask เข้าไปอีกคงยิ่งแย่ และก็มีข้อสังเกตแบบตรงไปตรงมาว่าในทางปฏิบัติ สตาร์ตอัปด้านฮาร์ดแวร์ AI ก็ทดลองแนวนี้กันมานานแล้ว
ถ้ามี LLM-in-a-box อยู่จริงก็คงน่าสนใจมาก อีกไม่นานผมอาจได้ทำงานในสภาพแวดล้อมแบบออฟไลน์เต็มรูปแบบ (air-gap) และโซลูชันแบบนั้นน่าจะมีประโยชน์มาก
ผมลองรันโค้ดในสภาพแวดล้อม GPU ของ Modal ด้วยตัวเองแล้ว และตัวเลขประสิทธิภาพที่งานวิจัยอ้างก็นำมาทำซ้ำได้จริง ขอแชร์โค้ดผลลัพธ์ของโปรเจกต์ mirage ในชุด Triton + FlashInfer ได้ latency ต่อ 1 token ราว 19.2ms ขณะที่ MPK ภายใต้เงื่อนไขเดียวกันลดลงมาเหลือ 7.7ms ซึ่งดีขึ้นมาก
เมื่อก่อนผมเคยแข่ง CUDA เล็ก ๆ ครั้งหนึ่ง เป็นอัลกอริทึมขนานสายภาพหรือวิชัน แล้วผมก็คิดว่าตัวเองฉลาดเลย cache ผลลัพธ์กลางไว้ในหน่วยความจำ พอประกาศผลออกมากลับพบว่าคนอื่นส่งโค้ดที่เร็วกว่าผมมาก พอไปดูเหตุผลก็พบว่า พวกเขาไม่ cache ผลลัพธ์กลางเลย แต่เลือกคำนวณซ้ำต่อไปเรื่อย ๆ เพราะต้นทุนการคำนวณถูกกว่าการวิ่งไปกลับหน่วยความจำมาก โปรเจกต์นี้ก็น่าจะคล้ายกัน การคอมไพล์เป็น megakernel ทำให้ขอบเขตระหว่าง layer หายไป การแชร์ผลลัพธ์กลางลดลงแต่ปริมาณการคำนวณเพิ่มขึ้น อย่างไรก็ตามโดยรวมแล้วการลดการวิ่งไปกลับหน่วยความจำให้กำไรมหาศาล โดยเฉพาะใน convolution network น่าจะมี sweet spot อยู่ เพียงแต่ผมยังไม่แน่ใจว่า megakernel จัดการจุดนี้อย่างไร
ตอนนี้ก็ยังมีอุปมาใหม่ ๆ สำหรับ LLM โผล่มาเรื่อย ๆ ผมเลยนึกเล่น ๆ ว่าเราจะมอง LLM เป็นเหมือนทรานซิสเตอร์ได้ไหม ตอนนี้มันชวนให้นึกถึงยุคคอมพิวเตอร์ขนาดเท่าห้องที่ใช้บัตรเจาะรูและทำได้แค่การคูณ ลองจินตนาการดูว่าถ้ารัน o3-pro ได้พร้อมกันหนึ่งล้านคำขอ จะเกิดอะไรขึ้นบ้าง
โปรเจกต์นี้มาจาก CMU (Carnegie Mellon) และทาง Hazy Research ของ Stanford ก็มีบล็อกเกี่ยวกับ megakernel ชื่อ No Bubbles เช่นกัน น่าประทับใจที่เห็นการแข่งขันในด้านนี้คึกคักมาก (เพิ่มเติม) ยังมีงานวิจัยที่พูดถึงภาพใหญ่ของโปรเจกต์ "mirage" ด้วย แต่ไม่ได้พูดถึงแนวทาง megakernel ลิงก์งานวิจัย
ผู้เขียนโพสต์เข้ามาตอบด้วยตัวเอง เห็นด้วยว่างานวิจัยกับ Stanford กำลังดำเนินไปแบบขนานกัน ความแตกต่างสำคัญคือเราโฟกัสที่คอมไพเลอร์สำหรับสร้าง megakernel แบบอัตโนมัติ
และ ThunderKittens ของ Hazy Research ก็เป็นไลบรารีที่เจ๋งมากเช่นกัน ช่วงนี้มีความพยายามอย่างมากในการทำให้เป็นระบบมากขึ้น สร้าง pipeline ใช้การแบ่งแยกแล้วพิชิต เพิ่มประสิทธิภาพให้สุดทาง และพัฒนาคอมไพเลอร์/DSL เฉพาะทาง เพื่อดึงศักยภาพของ NVIDIA GPU รุ่นใหม่ออกมาให้มากที่สุด
ตัวเลขประสิทธิภาพของ Qwen 8B ถ้าตรวจสอบแล้วเป็นจริงจะน่าประทับใจมาก ดูใช้งานได้จริงกว่าแนวทาง megakernel รุ่นก่อน ๆ เคอร์เนลแบบที่คงอยู่หนึ่งตัวต่อ SM นี้ทำให้นึกถึง Larrabee ถ้าโลกในวันนั้นเลือกเส้นทาง process-thread-SIMD แบบดั้งเดิมแทน CUDA เราคงได้เห็นโลกอีกแบบหนึ่ง
มีแนวคิดเรื่องการสร้าง LLM แบบคงรูปด้วย ASIC ล้วน แทนการทำ inference ด้วยซอฟต์แวร์ ข้อได้เปรียบด้านต้นทุนจะเป็นอย่างไร จะมีความเป็นไปได้ไหมที่จะมีชั้นเพิ่มเติมที่จัดการหรือปรับจูนจากฝั่งซอฟต์แวร์ได้ ตอนนี้เราอาจใกล้ถึงจุดที่ถือว่า “ดีพอแล้ว” มาก ๆ และในอีก 2~4 ปีข้างหน้าอาจมีคนตัดสินใจตรึงไว้บนชิปเฉพาะทาง คำถามคือข้อได้เปรียบของฮาร์ดแวร์ที่เฉพาะทางสุด ๆ จะเริ่มฉายชัดเมื่อไร