Muse Spark Open: 1,109 โทเคนต่อวินาทีคือการเปลี่ยนแปลงของเคอร์เนล ไม่ใช่การเปลี่ยนแปลงของโมเดล

M​use Spark Open was released on October 15, 2026, and nothing about its weights changed this week. What changed is the software underneath them. On November 11, 2026 the SGLang team at LMSYS published an engineering post titled "M​use Spark Open Ultrafast: Pushing beyond 1,000 Tokens per second," and the headline number in it is decoding throughput on eight GB300s going from 517 to 1,109 tokens per second for the same model on the same hardware. The mechanism is a megakernel: one resident kernel per layer instead of roughly 1,450 separate kernel launches per decode step. So this is a serving story, not a model story, and the interesting question is not whether 1,109 is impressive — it is — but which of those numbers you are allowed to carry into a capacity plan. The short answer is fewer than you would like: the benchmark is one request, one dataset, batch size one, and no one outside the authors has reproduced it yet.

ค่าฐานที่ไม่มีใครนำมาใช้เป็นเกณฑ์เปรียบเทียบ: 1,450 การเปิดตัวต่อโทเคน

A normal GPU kernel does one thing. It runs one matrix multiplication or one attention pass as a grid of thread blocks spread across the streaming multiprocessors, then exits. A serving engine like SGLang strings those kernels together into a chain, and for M​use Spark Open that chain is long: the post puts the count at about 1,450 kernel launches for each decoding step. Between two consecutive kernels the GPU drains and refills. SMs that have finished their tile sit idle waiting for the next launch, and inside an SM individual warps idle too, because a warp that has issued its loads has nothing to do until the data lands.

นั่นคือต้นทุนที่เมกะเคอร์เนลเข้าโจมตี และผู้เขียนหยิบการเปรียบเทียบจากยุคทศวรรษ 1960 มาอธิบาย: เครื่องรุ่นแรก ๆ รันงานแบบแบตช์ โดยงานหนึ่งครอบครองเครื่องทั้งหมด และโปรเซสเซอร์ก็ว่างเปล่าระหว่างที่เทปหมุนอยู่ การทำงานหลายโปรแกรมพร้อมกันแก้ปัญหานี้ด้วยการเก็บงานหลายงานไว้ในหน่วยความจำและสลับไปยังอีกงานหนึ่งเมื่อใดก็ตามที่งานปัจจุบันถูกบล็อก เมกะเคอร์เนลคือแนวคิดนั้นที่ย้ายมาสู่ GPU — เคอร์เนลหนึ่งตัวคงอยู่ในทุก SM และรันงานหลายงานพร้อมกัน ดังนั้นวาร์ปจึงสามารถเริ่มงานของคำสั่งถัดไปได้ก่อนที่คำสั่งก่อนหน้าจะเสร็จสิ้น

ผลที่วัดได้จากรูปที่มาพร้อมกับส่วนประสิทธิภาพ คือขั้นตอนถอดรหัสที่เปลี่ยนจาก 12.9 ms เป็น 5.9 ms รูปนั้นแบ่งขั้นตอนออกเป็นสองครึ่ง คือ draft และ verify: 1.16 ms และ 10.9 ms โดยมี 1,450 เคอร์เนลใน baseline, 0.44 ms และ 5.3 ms โดยมี 97 เคอร์เนลใน megakernel สิบสองจุดเก้าไปห้าจุดเก้าคือประมาณ 2.2× และ 1,109 ส่วน 517 คือประมาณ 2.1× ดังนั้นอัตราส่วนพาดหัวสองค่าจึงสอดคล้องกัน มีการตรวจสอบที่สามให้ใครก็ตามที่ยินดีจะหาร: ขั้นตอนนี้ตรวจสอบ 16 โทเคนและยอมรับ 6.6 ดังนั้น 6.6 โทเคนใน 5.9 ms คำนวณออกมาได้ประมาณ 1,120 โทเคนต่อวินาที เลขคณิตสอดคล้องกันภายใน ซึ่งเป็นสิ่งแรกที่ต้องยืนยันเกี่ยวกับตัวเลขแบบนี้ และไม่ใช่สิ่งสุดท้าย

สิ่งหนึ่งที่ตัวเลขไม่ได้บอกคือความเร็วนั้นอยู่ที่เท่าใดต่อ GPU โพสต์รายงานปริมาณงานรวมสำหรับการติดตั้งใช้งาน GPU แปดตัวทั้งชุด เมื่อหารด้วยแปด จะได้ประมาณ 139 โทเคนต่อวินาทีต่อ GPU — นี่เป็นการคำนวณของผมจากตัวเลขของพวกเขา ไม่ใช่ตัวเลขที่โพสต์เผยแพร่ และเทียบไม่ได้กับตัวเลขต่อ GPU ที่คุณอาจเห็นอ้างอิงสำหรับโมเดลอื่นบนเครื่องอื่น

ขอบเขตของเคอร์เนลอยู่ตรงไหน และเหตุใดจึงสิ้นสุดตรงนั้น

เมกะเคอร์เนลไม่ได้กลืนโมเดลทั้งตัว การเรียกใช้แต่ละครั้งครอบคลุมตั้งแต่ attention ของเลเยอร์ decoder หนึ่ง ไปจนถึง qkv projection และการเขียน KV ของเลเยอร์ถัดไป ซึ่งเป็นขอบเขตที่ตั้งใจไว้ ไม่ใช่ข้อจำกัดของ nerve ที่มันทำงานได้เพราะส่วนใหญ่ของสิ่งที่ attention ของเลเยอร์อ่าน — ประวัติ KV — ถูกเขียนไว้ในขั้นก่อนหน้า และไม่ขึ้นกับสิ่งใดที่กำลังดำเนินอยู่ ดังนั้นการเรียกใช้ครั้งใหม่จึงเริ่มโหลดได้ทันทีที่ข้อมูลพร้อม ก่อนที่เอาต์พุตของเลเยอร์ก่อนหน้าจะมีอยู่ เลเยอร์แรกและเลเยอร์สุดท้ายมีการดำเนินการที่แตกต่างกันเล็กน้อย และโพสต์ก็ระบุไว้เช่นกัน

ภายในการ launch หนึ่งครั้ง โครงสร้างถูกกำหนดตายตัว: หนึ่ง thread block ที่มีสิบสอง warp บนแต่ละ SM จากทั้งหมด 144 SM เปิดตัวเป็นคลัสเตอร์ขนาดสี่ โดยใช้ distributed shared memory ในจุดที่สามารถทำ reduction ได้ด้วยความหน่วงต่ำ สิบสอง warp นี้ไม่สามารถใช้แทนกันได้ Warp 0 โหลด KV และ weights และใช้ช่วงเวลาว่างในการ prefetch weights ของ operation ถัดไป Warp 1 และ 2 รัน activations และ tensor cores Warp 3 ถึง 11 แบ่งกันทำ attention, dense GEMMs, expert GEMM และการรับ all-reduce กับ all-gather ในเลเยอร์ sliding-window warp 4 ถึง 11 โหลด KV ของตัวเองแทนที่จะรอ warp 0

Getting there required a change in how the model is split. M​use Spark Open is a mixture-of-experts model, routing every token to 8 of 256 experts, with experts stored in 4-bit NVFP4 and some layers left in bf16. It has 62 decoder layers, of which 46 use sliding-window attention over the last 2,048 tokens and 16 use global attention, plus a pooler projection that shrinks each token vector before the experts and an unpooler that restores it afterwards. The megakernel shards both MoE GEMMs on the N dimension, which forces an all-gather after each GEMM — an fp4 all-gather after the first, a bf16 one after the second. That is one more communication step than vanilla tensor parallelism, which shards on the intermediate dimension and needs a single all-reduce, but the authors argue it moves less data overall, scales better to the batch sizes they want next, and can be made bitwise identical to expert parallelism so it matches training numerics.

รายละเอียดนั้นสำคัญกว่าที่ดูเหมือน ทีม inference มักยอมรับว่า kernel ที่เร็วกว่าจะให้ตัวเลขที่ต่างออกไปเล็กน้อย และแอบปรับเกณฑ์การประเมินของตนใหม่แบบเงียบ ๆ ส่วน kernel ที่ให้ผลเท่ากันทุกบิตกับการทำ parallelisation ฝั่ง training จะตัดข้อโต้แย้งนั้นออกไปทั้งหมด ซึ่งเป็นคุณสมบัติแบบที่ทำให้ optimisation ถูกนำไปใช้จริง มากกว่าจะเป็นเพียงสิ่งที่ถูกชื่นชม

เคอร์เนลตัวใหญ่ขนาดนี้ต้องจ่ายบิลสามใบ

megakernel ไม่ได้ฟรี โพสต์นี้เปิดเผยเกี่ยวกับภาระค่าใช้จ่ายส่วนเกินอย่างตรงไปตรงมาอย่างผิดปกติ และนั่นคือส่วนที่คุ้มค่าจะอ่านซ้ำ เพราะมันเป็นส่วนที่บอกคุณด้วยว่าเหลือพื้นที่อีกมากแค่ไหน

ค่าใช้จ่ายแรกอยู่ที่แคชคำสั่ง SM หนึ่งตัวมีแคชคำสั่งประมาณ 128 KB และเคอร์เนลเหล่านี้มีโค้ดเครื่องขนาด 215 ถึง 314 KB ต่อตัว เพราะแต่ละตัวเก็บทุกเฟสของเลเยอร์ การเปิดใช้งานจะนำเฉพาะ 32 ถึง 64 KB แรกเข้ามาใกล้ SM เท่านั้น ส่วนที่เกินหน้าต่างนั้นจะถูกดึงทีละบรรทัดเมื่อใช้งานครั้งแรก ผู้เขียนได้ทำการทดลองที่ชัดเจนในเรื่องนี้และรายงานว่า บล็อกขนาด 100 KB สองบล็อกที่แต่ละบล็อกใช้เวลา 12.4 µs เมื่ออยู่ลำพัง ใช้เวลา 19.5 µs เมื่อต่อเชื่อมติดกัน และเมื่อย่อบล็อกหนึ่งเหลือ 20 KB บล็อกอีกอันก็กลับมาที่ความเร็วเมื่ออยู่ลำพัง คำสั่งเดินทางผ่านแคช L2 เดียวกันกับข้อมูล ดังนั้นเคอร์เนลที่มีขนาดโตเกินแคชคำสั่งจึงต้องแย่งชิงกับการโหลดเวตของตัวเอง

ค่าใช้จ่ายก้อนที่สองคือการ prefetch น้ำหนัก และโพสต์นี้มองมันเป็นปัญหาการจัดวางที่มีสามคำตอบ หน่วยความจำร่วม (shared memory) ได้วงแหวนขนาด 5 × 32 KB ต่อ SM รวมราว 23 MB ทั่วทั้ง GPU ส่วนที่เหลือไปอยู่ที่ L2 และ L2 ของ GB300 จุได้ราว 135 MB แม้ว่าอะไรก็ตามใน L2 จะต้องเดินทางสั้น ๆ อีกรอบเข้าไปใน shared memory และ L2 ก็ถูกใช้ร่วมกับทุกอย่างอื่นที่กำลังทำงานอยู่ ทั้งสองอย่างซ้อนกัน โดยช่วงใกล้ท้ายของเลเยอร์ วงแหวนมีน้ำหนัก qkv ของเลเยอร์ถัดไปอยู่แล้ว และเหนือสิ่งนั้นขึ้นไป warp 3 จะอ่านว่าขั้นตอนก่อนหน้าส่งไปยัง expert ใด แล้วดึงส่วนข้อมูลของ GPU เครื่องนี้จาก 14 ตัวที่ถูกใช้มากที่สุดเข้าไปใน L2 ประมาณ 50 MB ต่อ GPU โดยเดิมพันว่าขั้นตอนที่ต่อเนื่องกันจะส่งไปคล้าย ๆ กัน การเดิมพันนี้มีมูลค่าราว 0.2 ms จากขั้นตอนที่ใช้เวลา 5.3 ms และอาจทำให้สูญเสียแบนด์วิดท์เมื่อคาดผิด

บิลใบที่สามคือการซิงโครไนซ์ และเป็นใบที่ไม่มีคำตอบที่ชัดเจน เมื่อ warp รอ GPU อีกตัวผ่าน NVLink จะไม่มีการขัดจังหวะและไม่มีการรอที่แอดเดรสในฮาร์ดแวร์ warp จะ poll ธงจนกว่ามันจะเปลี่ยน ไมโครเบนช์มาร์กในโพสต์ระบุว่าการ poll แบบ acquire ที่ขอบเขต GPU อยู่ที่ 141 ถึง 145 ns ธงทางเดียวระหว่าง GPU สองตัวบนโหนดเดียวกันอยู่ที่ประมาณ 1.9 µs และการ poll ที่ต่อคิวอยู่หลังการคัดลอกจำนวนมากอยู่ที่ 1.0 ถึง 2.6 µs เฟนซ์มีค่าใช้จ่าย 0.4 µs เมื่อไม่ได้ใช้งาน และ 3.6 ถึง 4.0 µs เมื่อมี TMA load กำลังทำงานอยู่ แม้งานจะกระจายอย่างสมบูรณ์แบบ GPU ก็ยังสามารถตามกันไม่ทันได้ถึง 4 µs และทุกจุดหยุดในเลเยอร์จะรอตัวที่ช้าที่สุด เมกะเคอร์เนลแทนที่เฟนซ์ด้วยค่าซึ่งเป็น sentinel ที่เขียนลงในสล็อตสโตร์ธรรมดา — ธงที่บรรจุข้อมูล, +0/−0 สำหรับ all-reduce, NaN ที่มีบิตเป็นหนึ่งทั้งหมดสำหรับการ gather แบบ bf16 — และรีเซ็ตสำเนาที่ไม่ได้ใช้งานของแต่ละสล็อตในขณะที่อีกสำเนาถูกใช้งาน

ตัวเลขเหล่านี้คือสิ่งที่ทำให้ผลลัพธ์รู้สึกเหมือนงานวิศวกรรมมากกว่าเวทมนตร์: ความคลาดเคลื่อนของ GPU 4 µs เทียบกับสเต็ป 5.9 ms นั้นถือว่าเล็กน้อย แต่จุดหยุดซิงโครไนซ์สี่จุดต่อเลเยอร์คูณกระจายไปทั่ว 62 เลเยอร์ ก็เป็นค่าโอเวอร์เฮดแบบที่กลืนกินชัยชนะ 2× ได้พอดี ถ้าคุณทำพลาด

5.9 ms ยังคงถูกใช้ไปกับอะไรอยู่

โพสต์นี้มีบทวิจารณ์ของตัวเองรวมอยู่ด้วย และบรรทัดที่เป็นประโยชน์ที่สุดในนั้นพูดถึงช่องว่างหลัง GEMM ครั้งที่สอง บล็อกต่าง ๆ ทำ GEMM นั้นเสร็จภายในเวลาใดก็ได้ระหว่าง 55 ถึง 64 µs นับจากจุดเริ่มต้นของเลเยอร์ (เปอร์เซ็นไทล์ที่ 10 ถึง 90) แต่ขั้นตอน combine รอเอาต์พุตจากทุก GPU ดังนั้นจึงเริ่มไม่ได้จนกว่าจะถึง 65 µs ตั้งแต่จุดสิ้นสุดของ GEMM 2 จนถึงจุดสิ้นสุดของเลเยอร์ บล็อกหนึ่งจะอยู่ในสภาพที่ไม่มีงานทำประมาณ 12.5 µs โดยเฉลี่ย — เกือบครึ่งหนึ่งของเวลาว่างทั้งหมดของมัน ห่วงโซ่นั้น ไม่ใช่โอเวอร์เฮดจากการ launch เคอร์เนลที่ megakernel กำจัดไปแล้ว คือสิ่งที่ผู้เขียนบอกว่าพวกเขากำลังพิจารณาต่อไป

พวกเขายังอธิบายด้วยว่าเหตุใด attention ของเลเยอร์ถัดไปจึงไม่สามารถเริ่มเร็วขึ้นได้ง่าย ๆ PDL เปิดตัวมันที่ 79 µs ซึ่งเร็วกว่าที่บล็อกแรกจะออกประมาณ 4 µs แต่บล็อกของมันจะมาถึงก็ต่อเมื่อบล็อกก่อนหน้าออกจาก SM ที่ 84 µs ดังนั้นการเปิดตัวเร็วขึ้นก็ไม่ได้ทำให้อะไรเริ่มเร็วขึ้น วิธีแก้คือ prefetch แถว KV เก่าของมันเข้า L2 ที่ 65 µs แล้วให้รออยู่ที่นั่น ซึ่งเป็นรูปแบบเดียวกับการ prefetch น้ำหนัก: เมื่อคุณกำจัดการรอไม่ได้ คุณก็ย้ายงานเข้าไปในช่วงรอนั้น

การทำโปรไฟล์เป็นปัญหาทางโครงสร้างอีกประการหนึ่ง และเป็นปัญหาที่ดี: เคอร์เนลดูเหมือนแท่งยาวหนึ่งแท่งต่อเลเยอร์สำหรับโปรไฟเลอร์ภายนอกใด ๆ ดังนั้นมันจึงจับเวลาเอง เลน 0 ของแต่ละบล็อกเขียนเวลาประทับลงในอาร์เรย์โกลบอลผ่านรีจิสเตอร์ globaltimer และผู้เขียนวัดต้นทุนของ instrumentation นั้นไว้ที่ประมาณ 2.5 µs ต่อการเรียกใช้งานเลเยอร์หนึ่งครั้ง หรือ 2.8% ของหนึ่งสเต็ป โดยเอาต์พุตยังคงเหมือนเดิมทุกบิตและมีแฟล็กการบิลด์เพื่อคอมไพล์มันออกไป เคอร์เนลที่ต้องพกโปรไฟเลอร์ของตัวเองไว้ด้วยคือเคอร์เนลที่กลายเป็นระบบปฏิบัติการไปแล้ว ซึ่งเป็นกรอบแนวคิดที่โพสต์ใช้อยู่พอดี: โปรเซสคือเคอร์เนลหนึ่งตัวที่ถือครองทรัพยากร GPU ทั้งหมด เธรดคือวาร์ปที่ไม่มีการ preemption การรอ I/O คือการ busy-wait บนแฟล็ก DMA คือ TMA ที่ถูกติดตามด้วย mbarrier ข้อความคือการ store ธรรมดา ๆ ข้าม NVLink

โมเดล Draft ก็เร็วขึ้นเช่นกัน และนั่นก็คือครึ่งหนึ่งของความสำเร็จ

M​use Spark Open's decode path is speculative: a draft model proposes tokens and the big model verifies them in one pass. The draft here is DFlash V2, which guesses the next 15 tokens in a single pass, and the big model checks all of them in a pass over 16 tokens. The verification uses a tree of candidate tokens generated from the draft rather than a linear chain, which the post says is slightly better.

ทีมใช้วิธีแบบ megakernel เดียวกันกับดราฟต์ โดยเลเยอร์ดราฟต์ห้าชั้นแต่ละชั้นทำงานเป็นหนึ่งครั้งในการเรียกใช้งาน บวกอีกสามรายการ ซึ่งทำให้ดราฟต์ลดจาก 175 เคอร์เนลเหลือ 8 และจาก 1.16 ms ต่อสเต็ปเหลือ 440 µs เป็นการปรับปรุง 2.6× ในครึ่งนั้นของสเต็ป ตัวเลขการยอมรับน่าจะนำไปวางเทียบกับข้อมูลเปรียบเทียบสาธารณะที่ใกล้เคียงที่สุดซึ่งโพสต์เสนอบอกไว้

• Acceptance length — M​use Spark Open verifies 16 tokens per step and accepts 6.6 on average; Xiaomi's MiMo-V2.5-Pro, which crossed 1,000 tokens per second earlier, verifies 8 and accepts 6.3

• ทำไม 16 — โทเคนสิบหกตัวมีต้นทุนในการตรวจสอบสูงกว่าเพราะมันกระตุ้นผู้เชี่ยวชาญมากขึ้น แต่ทีมพบว่าการไปถึง 6.6 ด้วยต้นไม้ที่กว้างกว่านั้นทำได้ง่ายกว่าการไปถึงความยาวเดียวกันด้วยสายโซ่ที่แคบกว่า

• สแตกฉบับร่าง — โค้ด CUDA และ PTX จำนวน 8,000 บรรทัดสำหรับโมเดลขนาดใหญ่ โดยเผยแพร่โค้ดไว้ที่ facebookresearch/sgl-muse-mega

อ่านการเปรียบเทียบนี้อย่างละเอียด เพราะมันไม่ใช่การทดสอบประสิทธิภาพระหว่างสองโมเดล แต่เป็นการเปรียบเทียบการกำหนดค่าการถอดรหัสแบบคาดเดาสองแบบ โดยแบบหนึ่งเผยแพร่โดย Xiaomi และอีกแบบโดยทีม SGLang บนฮาร์ดแวร์ที่แตกต่างกันและภาระงานที่แตกต่างกัน ข้อสรุปเดียวที่การเปรียบเทียบนี้รองรับคือ ต้นไม้การตรวจสอบที่กว้างขึ้นเป็นเส้นทางที่เป็นไปได้สู่ความยาวการยอมรับที่เท่ากัน ซึ่งเป็นข้อสรุปเชิงการออกแบบ ไม่ใช่การจัดอันดับ

ตัวเลขนี้คืออะไร และไม่ใช่อะไร

นี่คือส่วนที่แผนความจุต้องทำให้ถูกต้อง และเป็นส่วนที่ตัวโพสต์เองก็ยอมรับอย่างตรงไปตรงมา ตัวเลข 1,109 คือค่ามัธยฐานของ 20 ขั้นตอนการถอดรหัส (decode steps) ในคำขอ GSM8K หนึ่งรายการ ที่ระดับ tensor parallel 8 บน GB300s แปดตัว โดยมี batch size เท่ากับหนึ่งและเปิด speculative decoding ไว้ โพสต์ระบุอย่างชัดเจนว่า megakernel ในปัจจุบันรองรับเฉพาะ batch size เท่ากับหนึ่งเท่านั้น และการขยายให้รองรับแบตช์ที่ใหญ่ขึ้นคืองานชิ้นถัดไป

ข้อจำกัดแต่ละข้อตัดข้ออ้างที่แตกต่างกันออกไป แบทช์ขนาดหนึ่งหมายความว่านี่ไม่ใช่ผลลัพธ์ด้านทรูพุตในความหมายที่ทีมให้บริการใช้คำนี้: ไม่มีกำไรจากการทำแบทช์อยู่ในนั้น ไม่มีคิว ไม่มีการผสมความยาวพรอมป์ต์ ชุดข้อมูลชุดเดียวหมายความว่าตัวเลขนั้นผูกกับพรอมป์ต์และรูปแบบผลลัพธ์ของ GSM8K และผูกกับรูปแบบการจัดเส้นทางที่วัดได้ของเอ็กซ์เพิร์ต 49 ตัวต่อเลเยอร์ ซึ่งผู้เขียนนำไปใช้เป็นการจัดเส้นทางแบบสังเคราะห์ในการโปรไฟล์ของตนเอง ยี่สิบสเต็ปของคำขอเดียวหมายความว่าความแปรปรวนถูกจำกัดด้วยจำนวนสเต็ป ไม่ใช่จำนวนคำขอ

และคำขอเดียวเป็นกรณีที่แย่ที่สุดสำหรับการซ้อนทับ ซึ่งผู้เขียนเองกล่าวไว้ว่า “เมื่อมีคำขอเพียงหนึ่ง ทุกการรอจะถูกเปิดเผย; เมื่อมีมากขึ้น มักจะมีงานที่พร้อมทำอยู่เกือบตลอดเวลา” ดังนั้น 2.1× ของเมกะเคอร์เนลจึงเป็นการวัดในสภาวะที่การซ่อนความหน่วงทำได้ยากที่สุด และโพสต์ระบุชัดว่าแบตช์ที่ใหญ่ขึ้นควรง่ายกว่า ไม่ใช่ยากกว่า นั่นเป็นข้อโต้แย้งที่สนับสนุนผลลัพธ์ และก็ยังเป็นข้อโต้แย้งมากกว่าการวัด

สิ่งที่รู้ได้ในวันนี้จึงแคบและเป็นจริง: โค้ดมีอยู่ในรีโพซิทอรีสาธารณะ โพสต์อธิบายกลไกไว้ละเอียดพอที่จะตรวจสอบได้ การคำนวณภายในสอดคล้องกันในอัตราส่วนอิสระสามค่า และผู้เขียนได้บันทึกค่าใช้จ่ายส่วนเกินสามประการที่เป็นขีดจำกัดของผลประโยชน์ที่จะเพิ่มขึ้นอีก สิ่งที่ยังรู้ไม่ได้คือทีมอิสระจะได้ 1,109 บน GB300 แปดเครื่องของตนเองหรือไม่ แนวทางนี้จะอยู่รอดในแบตช์จริงหรือไม่ และมันจะไปอยู่ในบิลด์ SGLang ที่เผยแพร่แล้วหรือไม่ — โพสต์ไม่ได้อ้างว่าได้ทำเช่นนั้น และรีโพซิทอรีก็เป็นสิ่งประดิษฐ์งานวิจัยมากกว่าจะเป็นแพ็กเกจที่ติดตั้งได้ การเปรียบเทียบที่จะตัดสินได้มากที่สุดคือการเปรียบเทียบกับ baseline ที่ปรับแต่งแล้วบนฮาร์ดแวร์เดียวกัน เนื่องจาก "default chain ของ SGLang" เป็นเป้าหมายที่เคลื่อนที่ และทีมที่เขียน megakernel ก็เป็นทีมเดียวกับที่ปรับแต่ง baseline ที่มันเอาชนะได้

สิ่งนี้เกี่ยวข้องอย่างไรกับ 169 โทเคนต่อวินาทีที่คุณเช่าได้ในวันนี้

Artificial Analysis records a very different number for M​use Spark Open: 169 tokens per second of output speed, measured end to end including time to first token, against a hosted endpoint, with a first-answer-token time of 13.81 seconds. That is not a contradiction and it is not a rebuke. It is a different measurement of a different deployment: one stream through somebody's managed API, versus one request through your own eight-GPU node with a research kernel.

ช่องว่างนี้คือสิ่งที่มีประโยชน์ที่สุดในเรื่องนี้สำหรับใครก็ตามที่ยังไม่ได้เป็นเจ้าของฮาร์ดแวร์นั้น หาก endpoint ตอบที่ 169 โทเคนต่อวินาที และสแต็ก GPU แปดตัวที่สร้างขึ้นเองตอบที่ 1,109 ความแตกต่างนั้นไม่ใช่เหตุผลที่จะสร้างสแต็กนั้น GB300 จำนวนแปดตัวเป็นการซื้อสินทรัพย์ทุนที่วัดกันเป็นหลักล้าน kernel เป็นโค้ดวิจัย และข้อจำกัด batch-size-one หมายความว่าปัจจุบันสแต็กนี้เหมาะสมพอดีกับงานประเภทที่ endpoint แบบโฮสต์จัดการได้ดีอยู่แล้ว นั่นคือหนึ่งคำขอในแต่ละครั้ง จุดที่ megakernel เปลี่ยนการคำนวณอย่างแท้จริงคือสำหรับทีมที่ซื้อ GPU ไปแล้วเพื่อการเทรน หรือสำหรับงานที่มีข้อจำกัดด้านที่ตั้งข้อมูล และตอนนี้กำลังถามว่าจะใช้ประโยชน์จาก GPU เหล่านั้นได้มากเพียงใดที่ batch ต่ำ สำหรับทีมนั้น 2.1× บนซิลิคอนเดียวกันด้วยการคำนวณเชิงตัวเลขแบบเดียวกันคือตัวเลขจริงที่มีใบแจ้งหนี้จริงแนบมา และใบแจ้งหนี้ก็ชำระแล้ว

สำหรับคนอื่นๆ การตีความตามจริงก็คือ โพสต์นี้เป็นสัญญาณว่าโครงสร้างพื้นฐานด้าน inference กำลังมุ่งไปทางไหน ไม่ใช่ผลิตภัณฑ์ที่คุณจะนำมาใช้ได้ในสัปดาห์นี้ ทิศทางนี้สอดคล้องกับสิ่งที่แล็บเดิมเผยแพร่มาตลอดทั้งปี คือเคอร์เนลที่น้อยลงแต่ใหญ่ขึ้น มีการย้ายการตัดสินใจด้านการจัดตารางเข้าไปในเคอร์เนลมากขึ้น และมอง GPU เป็นเครื่องที่ต้องเขียนโปรแกรมให้ มากกว่าจะเป็นชุดของการเรียกใช้ไลบรารีที่ต้องเรียงลำดับ

That is also the right frame for our own position, which is worth stating plainly. M​use Spark Open is not on OrcaRouter; it is a download, a gated repository and a licence, not an endpoint we serve, and we are not going to describe it as one. Muse Spark 1.2 and Muse Spark 1.1 are the two Muse models OrcaRouter actually routes, and they are there at Meta's own list price because the provider's list price passes through rather than gets marked up. If your question is which model to put behind a product this quarter rather than which kernel to compile, that is the shorter path: one OpenAI-compatible endpoint across 200+ models, with the provider's list price passed through at 0% markup, plus automatic failover and a routing DSL that moves traffic between models by configuration rather than by redeploying code.

อะไรจะทำให้ชิ้นนี้เปลี่ยนไป

สี่สิ่ง โดยเรียงจากมากไปหาน้อยตามว่าสิ่งใดจะเปลี่ยนภาพรวมได้มากเพียงใด

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

• การทำซ้ำโดยอิสระ ตัวเลขทุกตัวที่นี่มาจากฮาร์เนสของผู้เขียนเอง บนฮาร์ดแวร์ของผู้เขียนเอง ด้วยการกำหนดเส้นทางสังเคราะห์ของผู้เขียนเองสำหรับรอบการโปรไฟล์ นั่นเป็นเรื่องปกติสำหรับงานเคอร์เนล และก็ยังคงเป็นแหล่งข้อมูลเพียงแหล่งเดียว ทีมที่สองที่ทำได้ 1,109 หรือทำไม่ได้ จะเป็นข้อเท็จจริงที่ชี้ขาดเรื่องนี้

• บิลด์ SGLang ที่เผยแพร่แล้วซึ่งมีสิ่งนี้อยู่ โพสต์ดังกล่าวเผยแพร่โค้ดและอธิบายการออกแบบ แต่ไม่ได้ระบุว่า megakernel ได้เข้ามาแทนที่เส้นทางการถอดรหัสเริ่มต้นสำหรับโมเดลนี้ในเวอร์ชันที่จัดส่งแล้ว จนกว่าจะเป็นเช่นนั้น การนำมาใช้จึงหมายถึงการนำรีพозиทอรีงานวิจัยมาใช้

• การปิดช่องว่าง GEMM-2-to-combine ผู้เขียนได้ระบุเวลา idle ต่อบล็อก 12.5 µs ตรงจุดนั้นว่าเป็นเกือบครึ่งหนึ่งของส่วนที่เหลือแล้ว หากสายงานนั้นได้รับการแก้ไข ฮาร์ดแวร์เดียวกันจะให้ step time ที่ดีขึ้นอย่างมีนัยสำคัญ และ 2.1× จะกลายเป็นค่าขั้นต่ำของตัวเลขที่ใหญ่กว่า แทนที่จะเป็นตัวเลขพาดหัว

Nothing in that list is a reason to discount the result. A team that writes 8,000 lines of CUDA and PTX to move a step from 12.9 ms to 5.9 ms, publishes the code, and then publishes the exact places where the remaining time goes is doing the thing the field claims to value and rarely does. The correct summary is not that M​use Spark Open got faster. It is that a specific serving stack, on specific hardware, at a specific batch size, got 2.1× faster, and the recipe for the next 2× is now public.