Muse Spark Open: 1.109 tokens per seconde is een kernelwijziging, geen modelwijziging

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.

De baseline die niemand benchmarkt: 1.450 lanceringen per token

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.

Dat is de kosten die de megakernel bestrijdt, en de auteurs grijpen naar een analogie uit de jaren 1960 om die te verklaren: vroege machines draaiden batchjobs, waarbij één job de hele machine in beslag nam, en de processor bleef inactief terwijl een tape ronddraaide. Multiprogrammering loste dat op door meerdere jobs resident te houden en naar een andere over te schakelen zodra de huidige blokkeerde. Een megakernel is dat idee overgebracht naar de GPU — één kernel blijft resident op elke SM en draait vele jobs tegelijk, zodat warps kunnen beginnen aan werk voor de volgende operatie voordat de vorige operatie klaar is.

Het gemeten effect, uit de figuur bij de prestatiesectie, is dat een decodeerstap van 12,9 ms naar 5,9 ms gaat. Die figuur splitst de stap in de draft- en verificatiehelften: 1,16 ms en 10,9 ms met 1.450 kernels in de baseline, 0,44 ms en 5,3 ms met 97 kernels in de megakernel. Van 12,9 naar 5,9 is ongeveer 2,2×, en 1.109 gedeeld door 517 is ongeveer 2,1×, dus de twee belangrijkste ratio's kloppen met elkaar. Er is een derde controle beschikbaar voor iedereen die bereid is de deling uit te voeren: de stap verifieert 16 tokens en accepteert er 6,6, dus 6,6 tokens in 5,9 ms komt neer op ongeveer 1.120 tokens per seconde. De rekensom is intern consistent, en dat is het eerste dat je bij zo'n getal moet vaststellen, maar niet het laatste.

Wat de cijfers niet zeggen, is waar de snelheid per GPU op uitkomt. Het bericht rapporteert de totale doorvoer voor de hele implementatie met acht GPU's. Gedeeld door acht is dat ongeveer 139 tokens per seconde per GPU — mijn eigen berekening op hun getal, geen getal dat het bericht publiceert, en niet vergelijkbaar met een getal per GPU dat je misschien voor een ander model op een andere machine ziet genoemd.

Waar de kernelgrens ligt, en waarom die daar stopt

De megakernel slokt niet het hele model op. Elke launch bestrijkt de attention van één decoderlaag tot en met de qkv-projectie en KV-schrijfactie van de volgende laag, wat een bewuste grens is en geen beperking van nerve. Het werkt omdat het meeste van wat de attention van een laag leest — de KV-geschiedenis — in eerdere stappen is geschreven en van niets afhangt dat nog in uitvoering is, zodat een nieuwe launch het kan beginnen te laden zodra het beschikbaar is, nog voordat de output van de vorige laag bestaat. De eerste en laatste lagen hebben iets andere bewerkingen, en de post merkt dat ook op.

Binnen één launch ligt de opzet vast: één threadblok van twaalf warps op elk van de 144 SM's, gelanceerd in clusters van vier, met gebruik van gedistribueerd gedeeld geheugen waar een reductie met lagere latentie beschikbaar is. De twaalf warps zijn niet uitwisselbaar. Warp 0 laadt KV en gewichten en besteedt zijn speling aan het prefetchen van de gewichten van de volgende bewerking. Warps 1 en 2 draaien activaties en tensorcores. Warps 3 tot en met 11 verdelen onder elkaar: attention, de dense GEMM's, de expert-GEMM en de ontvangsten van all-reduce en all-gather. In de sliding-window-lagen laden warps 4 tot en met 11 hun eigen KV in plaats van op warp 0 te wachten.

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.

Dat detail is belangrijker dan het lijkt. Inferentieteams accepteren meestal dat een snellere kernel iets andere getallen oplevert en stellen hun evaluatiedrempels stilletjes bij. Een kernel die bitwise gelijk is aan de parallelisatie aan de trainingszijde haalt dat argument volledig weg, en dat is precies het soort eigenschap waardoor een optimalisatie wordt overgenomen in plaats van alleen bewonderd.

Drie rekeningen die een kernel van dit formaat moet betalen

De megakernel is niet gratis. De blogpost is ongewoon openhartig over de overheadkosten, en dat is het deel dat je twee keer moet lezen, want het is ook het deel dat je vertelt hoeveel ruimte er nog over is.

De eerste kostenpost is de instructiecache. Een SM cachet ongeveer 128 KB aan instructies, en deze kernels zijn elk 215 tot 314 KB aan machinecode, omdat elk ervan elke fase van een laag bevat. Een launch brengt slechts de eerste 32 tot 64 KB dicht bij de SM's; alles voorbij dat venster wordt bij eerste gebruik regel voor regel opgehaald. De auteurs voerden hierop een zuiver experiment uit en rapporteerden het: twee blokken van 100 KB die elk afzonderlijk 12,4 µs kostten, kostten 19,5 µs toen ze achter elkaar werden geplakt, en het verkleinen van één blok tot 20 KB bracht het andere terug op zijn solosnelheid. Instructies reizen door dezelfde L2-cache als data, dus een kernel die de instructiecache ontgroeit, gaat uiteindelijk concurreren met het laden van zijn eigen gewichten.

De tweede kostenpost is gewichten-prefetch, en de post behandelt dit als een plaatsingsprobleem met drie antwoorden. Shared memory krijgt een ring van 5 × 32 KB per SM, ongeveer 23 MB over de hele GPU. L2 is waar de rest naartoe gaat, en de L2 van GB300 bevat ongeveer 135 MB, hoewel alles in L2 een tweede, kortere reis naar shared memory nodig heeft en L2 wordt gedeeld met al het andere dat in uitvoering is. De twee stapelen zich op: tegen het einde van een laag bevat de ring al de eerste qkv-gewichten van de volgende laag. Daarbovenop leest warp 3 naar welke experts de vorige stap heeft gerouteerd en haalt de slices van deze GPU van de 14 meest gebruikte experts naar L2, ongeveer 50 MB per GPU, in de verwachting dat opeenvolgende stappen op dezelfde manier routeren. Die gok levert ongeveer 0,2 ms op van een stap van 5,3 ms, en kan bandbreedte kosten als hij verkeerd uitpakt.

De derde rekening is synchronisatie, en dat is degene zonder eenduidig antwoord. Wanneer een warp op een andere GPU wacht via NVLink, is er geen interrupt en geen wait-on-address in hardware; de warp pollt een vlag tot die verandert. De microbenchmarks van de post zetten een GPU-scope acquire-poll op 141 tot 145 ns, een eenrichtingsvlag tussen twee GPU's op hetzelfde knooppunt op ongeveer 1,9 µs, en een poll die achter bulkkopieën in de wachtrij staat op 1,0 tot 2,6 µs. Fences kosten 0,4 µs bij inactiviteit en 3,6 tot 4,0 µs met TMA-loads in uitvoering. Zelfs met perfect verdeeld werk kunnen GPU's tot 4 µs op elkaar achterlopen, en elke stop in de laag wacht op de traagste. De megakernel vervangt fences door sentinelwaarden die in gewone store-slots worden geschreven — een met data gevulde vlag, +0/−0 voor de all-reduce, all-ones NaN voor de bf16-gathers — en reset de inactieve kopie van elk slot terwijl de andere in gebruik is.

Dit zijn de cijfers waardoor het resultaat als engineering aanvoelt in plaats van magie: 4 µs GPU-scheefstand tegenover een stap van 5,9 ms is klein, maar vier synchronisatiestops per laag, vermenigvuldigd over 62 lagen, is precies het soort overhead dat een 2×-winst opeet als je het verkeerd doet.

Waar de 5,9 ms nog aan wordt besteed

Het bericht bevat zijn eigen kritiek, en de nuttigste regel daarin gaat over de kloof na de tweede GEMM. Blokken ronden die GEMM ergens tussen 55 en 64 µs in de laag af, van het 10e tot het 90e percentiel, maar de combine wacht op de output van elke GPU en kan dus pas beginnen om 65 µs. Van het einde van GEMM 2 tot het einde van de laag zit een blok gemiddeld ongeveer 12,5 µs zonder werk — bijna de helft van zijn totale idle tijd. Die keten, niet de kernel launch-overhead die de megakernel al heeft geëlimineerd, is waar de auteurs zeggen dat ze nu naar kijken.

Ze leggen ook uit waarom de attention van de volgende laag niet simpelweg eerder kan beginnen. PDL start die op 79 µs, ongeveer 4 µs voordat het eerste blok vertrekt, maar de blokken ervan komen pas aan zodra de vorige om 84 µs hun SM's verlaten, dus eerder starten zou niets eerder laten beginnen. De oplossing is om de oude KV-rijen al om 65 µs naar L2 te prefetchen, waar ze wachten. Dat is hetzelfde patroon als bij het prefetchen van gewichten: als je een wachttijd niet kunt wegnemen, verplaats je werk ernaartoe.

Profilering is het andere structurele probleem, en een goed probleem: voor elke externe profiler ziet de kernel er per laag uit als één lange balk, dus meet hij zichzelf. Lane 0 van elk block schrijft een timestamp naar een globale array via het globaltimer-register, en de auteurs meten de kosten van die instrumentatie op ongeveer 2,5 µs per layer launch, 2,8% van een step, waarbij de outputs bit-identiek blijven en er een build flag is om het eruit te compileren. Een kernel die zijn eigen profiler moet meedragen, is een kernel die een besturingssysteem is geworden, en dat is precies de framing die de post gebruikt: een proces is één kernel die alle GPU-resources vasthoudt, threads zijn warps zonder preemption, wachten op I/O is busy-waiting op een flag, DMA is TMA bijgehouden door mbarrier, berichten zijn gewone stores over NVLink.

Het conceptmodel werd ook sneller, en dat is de helft van de winst.

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.

Het team paste dezelfde megakernel-aanpak toe op het draftmodel: vijf draftlagen draaien elk als één launch, plus drie andere, waardoor het draftmodel van 175 kernels naar 8 gaat en van 1,16 ms per stap naar 440 µs, een verbetering van 2,6× op die helft van de stap. Het loont de acceptatiecijfers naast de dichtstbijzijnde openbare vergelijking te zetten die het bericht biedt.

• 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

Waarom 16 — zestien tokens kosten meer om te verifiëren omdat ze meer experts activeren, maar het team vond het gemakkelijker om 6.6 te bereiken met een bredere boom dan om dezelfde lengte te bereiken met een smallere keten

• Concept-stack — 8000 regels CUDA en PTX voor het grote model, met de code gepubliceerd op facebookresearch/sgl-muse-mega

Lees de vergelijking zorgvuldig, want het is geen benchmark tussen twee modellen. Het is een vergelijking van twee configuraties voor speculatieve decoding, één gepubliceerd door Xiaomi en één door het SGLang-team, op verschillende hardware en verschillende workloads. De enige bewering die het ondersteunt, is dat een bredere verificatieboom een haalbare route is naar dezelfde acceptatielengte, wat een ontwerpclaim is, geen rangschikking.

Wat dit nummer is, en wat het niet is

Dit is het onderdeel dat een capaciteitsplan goed moet hebben, en het is het onderdeel waarover de post zelf eerlijk is. Het cijfer 1.109 is de mediaan van 20 decoderingsstappen op één enkele GSM8K-aanvraag, bij tensorparallelle graad 8, op acht GB300's, met batchgrootte één en speculatieve decoding aan. De post stelt duidelijk dat de megakernel vandaag alleen batchgrootte één ondersteunt, en dat uitbreiding naar grotere batches het volgende stuk werk is.

Die beperkingen nemen elk een andere bewering weg. Batchgrootte één betekent dat dit geen doorvoerresultaat is in de zin waarin een servingteam het woord gebruikt: er zit geen batchingwinst in, geen wachtrij, geen mengeling van promptlengtes. Eén dataset betekent dat het getal vastzit aan de prompt- en uitvoerverdeling van GSM8K, en aan een gemeten routeringspatroon van 49 experts per laag dat de auteurs vervolgens als synthetische routering in hun eigen profilering gebruikten. Twintig stappen van één verzoek betekent dat de variantie wordt begrensd door het aantal stappen, niet door het aantal verzoeken.

En een enkele aanvraag is het slechtste geval voor overlap, wat de auteurs zelf zeggen: "bij één aanvraag wordt elke wachttijd zichtbaar; bij meer is er bijna altijd gereed werk." Dus de 2,1× van de megakernel is een meting van het regime waarin het verbergen van latentie het moeilijkst is, en de post is expliciet dat grotere batches gemakkelijker zouden moeten zijn in plaats van moeilijker. Dat is een argument in het voordeel van het resultaat, en het is nog steeds een argument in plaats van een meting.

Wat vandaag dan kenbaar is, is beperkt en reëel: de code staat in een openbare repository, de post beschrijft het mechanisme in voldoende detail om het te kunnen controleren, de interne rekenkunde klopt over drie onafhankelijke verhoudingen, en de auteurs hebben de drie overheads gedocumenteerd die verdere winst begrenzen. Wat niet kenbaar is, is of een onafhankelijk team 1,109 haalt op zijn eigen acht GB300's, of de aanpak een echte batch doorstaat, en of het in een uitgebrachte SGLang-build terechtkomt — de post beweert niet dat dat is gebeurd, en de repository is een onderzoeksartefact in plaats van een pakket dat je kunt installeren. De vergelijking die het meest zou beslechten, is met een afgestemde baseline op dezelfde hardware, aangezien "SGLang's standaardketen" een bewegend doel is en het team dat de megakernel schreef ook de baseline afstemt die het verslaat.

Hoe het zich verhoudt tot de 169 tokens per seconde die je vandaag kunt huren

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.

Het gat is het nuttigste in dit verhaal voor iedereen die de hardware nog niet bezit. Als een endpoint 169 tokens per seconde antwoordt en een zelfgebouwde stack van acht GPU's 1.109, dan is dat verschil geen reden om de stack te bouwen. Acht GB300's is een kapitaalinvestering van miljoenen, de kernel is onderzoekscode, en de beperking van batchgrootte één betekent dat de stack op dit moment precies goed is voor de workload die een gehost endpoint al goed aankan: één verzoek tegelijk. Waar de megakernel de rekening echt verandert, is voor een team dat de GPU's al heeft gekocht voor training of voor een workload met beperkingen rond gegevensresidentie, en zich nu afvraagt welke benutting het er bij een lage batch uit kan halen. Voor dat team is 2,1× op dezelfde silicon met dezelfde numerieke eigenschappen een echt getal met een echte factuur eraan vast, en die factuur is al betaald.

Voor alle anderen is de eerlijke interpretatie dat deze post een signaal is over waar de inferentie-infrastructuur naartoe gaat, niet een product dat je deze week in gebruik kunt nemen. De richting komt overeen met wat hetzelfde lab het hele jaar al publiceert: minder, grotere kernels, meer van de planningsbeslissing naar binnen de kernel verplaatst, en de GPU behandeld als een machine die geprogrammeerd moet worden in plaats van als een reeks bibliotheekaanroepen die op volgorde moeten worden uitgevoerd.

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.

Wat zou dit stuk veranderen?

Vier dingen, in aflopende volgorde van hoeveel ze het beeld zouden veranderen.

• Een batch groter dan één. De auteurs zeggen dat grotere batches de volgende stap zijn en beargumenteren dat die eenvoudiger zijn, omdat er bij meer tokens in verwerking sommige outputs van de vorige bewerking al gereed zijn terwijl de rest nog binnenkomt. Zodra er een gebatcht getal bestaat, is de vergelijking niet langer "één stream versus één stream" en wordt het een doorvoerclaim die een servingteam kan gebruiken.

• Een onafhankelijke reproductie. Elk cijfer hier komt uit de eigen harness van de auteurs, op hun eigen hardware, met hun eigen synthetische routing voor de profileringspass. Dat is normaal voor kernelwerk en het blijft een enkele bron. Een tweede team dat 1,109 evenaart — of er niet in slaagt dat te doen — is het feit dat de knoop zou doorhakken.

• Een uitgebrachte SGLang-build die het bevat. De post publiceert code en beschrijft het ontwerp; er staat niet dat de megakernel het standaard decode-pad voor dit model heeft vervangen in een uitgebrachte versie. Totdat dat gebeurt, betekent het adopteren ervan het adopteren van een onderzoeksrepository.

• Het dichten van de kloof tussen GEMM-2 en combine. De auteurs hebben de 12,5 µs idle tijd per blok daar al bestempeld als bijna de helft van wat overblijft. Als die keten wordt opgelost, levert dezelfde hardware een beduidend betere steptijd op, en wordt de 2,1× de ondergrens van een groter getal in plaats van het hoofdpunt.

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.