Muse Spark Open: 1 109 tokens per sekund är en kerneländring, inte en modelländring

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.

Baslinjen som ingen benchmarkar: 1 450 lanseringar 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.

Det är kostnaden som megakerneln angriper, och författarna tar till en analogi från 1960-talet för att förklara den: tidiga maskiner körde batchjobb, där ett jobb ägde hela maskinen, och processorn satt overksam medan en bandrulle snurrade. Multiprogrammering löste det genom att hålla flera jobb residenta och växla till ett annat när det aktuella jobbet blockerades. En megakernel är den idén flyttad till GPU:n — en kernel förblir resident på varje SM och kör många jobb samtidigt, så att warps kan börja arbeta med nästa op innan den föregående op:n har slutförts.

Den uppmätta effekten, enligt figuren som åtföljer prestandaavsnittet, är att ett avkodningssteg går från 12,9 ms till 5,9 ms. Figuren delar upp steget i dess utkast- och verifieringshalvor: 1,16 ms och 10,9 ms med 1 450 kärnor i baslinjen, 0,44 ms och 5,3 ms med 97 kärnor i megakärnan. Tolv komma nio mot fem komma nio är ungefär 2,2×, och 1 109 delat med 517 är ungefär 2,1×, så de två övergripande kvoterna stämmer med varandra. Det finns en tredje kontroll för den som är villig att göra divisionen: steget verifierar 16 token och accepterar 6,6, så 6,6 token på 5,9 ms blir ungefär 1 120 token per sekund. Aritmetiken är internt konsistent, vilket är det första man bör fastställa om ett tal som detta, och inte det sista.

En sak som siffrorna inte säger är var hastigheten hamnar per GPU. Inlägget rapporterar aggregerad genomströmning för hela distributionen med åtta GPU:er. Dividerat med åtta blir det cirka 139 tokens per sekund per GPU – min egen beräkning utifrån deras siffra, inte ett tal som inlägget publicerar, och inte jämförbart med ett per-GPU-tal som du kan se angivet för en annan modell på en annan maskin.

Var kärnans gräns ligger, och varför den slutar där

Megakärnan slukar inte hela modellen. Varje start omfattar attentionen i ett avkodarlager fram till nästa lagers qkv-projektion och KV-skrivning, vilket är en avsiktlig gräns snarare än en begränsning i nerve. Anledningen till att det fungerar är att det mesta som ett lagers attention läser – KV-historiken – skrevs i tidigare steg och inte beror på något som pågår, så en ny start kan börja läsa in det i samma stund som det är tillgängligt, innan föregående lagers utdata finns. Det första och det sista lagret har något olika operationer, och inlägget noterar det också.

Inom en enda launch är formen fast: ett trådblock med tolv warps på var och en av 144 SM:er, startade i kluster om fyra, med distribuerat delat minne där en reduktion med lägre latens är tillgänglig. De tolv warpsen är inte utbytbara. Warp 0 läser in KV och vikter och använder sin lediga tid till att förhämta nästa operations vikter. Warps 1 och 2 kör aktiveringar och tensor-kärnor. Warps 3 till 11 delar upp attention, de täta GEMM:erna, expert-GEMM:en samt mottagningarna för all-reduce och all-gather. I lagren med glidande fönster läser warps 4 till 11 in sin egen KV i stället för att vänta på 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.

Den detaljen är viktigare än den verkar. Inferensteam brukar acceptera att en snabbare kernel ger något annorlunda siffror och i tysthet justerar sina utvärderingströsklar. En kernel som är bitvis identisk med träningssidans parallellisering tar bort det argumentet helt, och det är den sortens egenskap som gör att en optimering införs i stället för att bara beundras.

Tre räkningar en så här stor kärna måste betala

Megakärnan är inte gratis. Inlägget är ovanligt öppenhjärtigt om overheadkostnaderna, och de är den del som är värd att läsa två gånger, eftersom de också är den del som berättar hur mycket utrymme som finns kvar.

Den första posten är instruktionscachen. En SM cachar ungefär 128 KB instruktioner, och dessa kärnor består av 215 till 314 KB maskinkod vardera, eftersom var och en innehåller varje fas av ett lager. Vid en start hamnar bara de första 32 till 64 KB nära SM:erna; allt bortom detta fönster hämtas rad för rad vid första användningen. Författarna genomförde ett rent experiment på detta och rapporterade det: två block på 100 KB som var för sig tog 12,4 µs tog 19,5 µs när de sammanfogades efter varandra, och när man krympte det ena blocket till 20 KB återfick det andra sin solohastighet. Instruktioner färdas genom samma L2-cache som data, så en kärna som växer ur instruktionscachen hamnar i konkurrens med sina egna viktladdningar.

Den andra posten är förhämtning av vikter, och inlägget behandlar det som ett placeringsproblem med tre svar. Delat minne får en ring på 5 × 32 KB per SM, cirka 23 MB över hela GPU:n. L2 är där resten hamnar, och GB300:s L2 rymmer cirka 135 MB, även om allt i L2 behöver en andra, kortare resa in i delat minne och L2 delas med allt annat som är igång. De två staplas: nära slutet av ett lager innehåller ringen redan nästa lagers första qkv-vikter. Dessutom läser warp 3 vilka experter föregående steg routade till och hämtar den här GPU:ns delar av de 14 mest använda till L2, cirka 50 MB per GPU, i ett vad om att efterföljande steg routar likadant. Det vadet är värt cirka 0,2 ms av ett steg på 5,3 ms och kan kosta bandbredd när det är fel.

Den tredje kostnadsposten är synkronisering, och det är den som saknar ett entydigt svar. När en warp väntar på en annan GPU över NVLink finns det inget avbrott och ingen wait-on-address i hårdvaran; warpen pollar en flagga tills den ändras. Inläggets mikrobenchmarks uppmäter en acquire-poll med GPU-scope till 141–145 ns, en enkelriktad flagga mellan två GPU:er på samma nod till cirka 1,9 µs och en poll som köats bakom masskopieringar till 1,0–2,6 µs. Fences kostar 0,4 µs när de är inaktiva och 3,6–4,0 µs med pågående TMA-laster. Även med perfekt distribuerat arbete kan GPU:er hamna efter varandra med upp till 4 µs, och varje stopp i lagret väntar på den långsammaste. Megakärnan ersätter fences med sentinelvärden som skrivs till vanliga lagringsplatser – en datafylld flagga, +0/−0 för all-reduce, NaN med alla ettor för bf16-gathers – och återställer den inaktiva kopian av varje plats medan den andra används.

Det här är siffrorna som får resultatet att kännas som ingenjörskonst snarare än magi: 4 µs GPU-skevhet mot ett steg på 5,9 ms är litet, men fyra synkroniseringsstopp per lager multiplicerat över 62 lager är precis den typ av overhead som äter upp en 2× vinst om man gör det fel.

Vad 5,9 ms fortfarande går åt till

Inlägget innehåller sin egen kritik, och den mest användbara raden i det handlar om gapet efter den andra GEMM:en. Block slutför den GEMM:en någonstans mellan 55 och 64 µs in i lagret, 10:e–90:e percentilen, men kombineringen väntar på varje GPU:s utdata, så den kan inte starta förrän vid 65 µs. Från slutet av GEMM 2 till slutet av lagret sitter ett block utan arbete i cirka 12,5 µs i genomsnitt – nästan hälften av dess totala väntetid. Det är den kedjan, inte overheaden för kernel-start som megakerneln redan har eliminerat, som författarna säger att de tittar på härnäst.

De förklarar också varför nästa lagers attention inte helt enkelt kan börja tidigare. PDL startar den vid 79 µs, cirka 4 µs innan det första blocket avslutas, men dess block landar först när de föregående lämnar sina SM:er vid 84 µs, så att starta tidigare skulle inte få något att börja tidigare. Lösningen är att prefetchar dess gamla KV-rader till L2 vid 65 µs, där de väntar. Det är samma mönster som vid viktprefetchningen: när man inte kan ta bort en väntan flyttar man in arbete i den.

Profilering är det andra strukturella problemet, och det är ett bra sådant: kärnan ser ut som en enda lång stapel per lager för alla externa profilerare, så den mäter sin egen tid. Lane 0 i varje block skriver en tidsstämpel till en global array via globaltimer-registret, och författarna mäter kostnaden för den instrumenteringen till ungefär 2,5 µs per lagerstart, 2,8 % av ett steg, med utdata lämnade bitidentiska och en byggflagga för att kompilera bort den. En kärna som måste bära sin egen profilerare är en kärna som har blivit ett operativsystem, vilket är exakt den inramning inlägget använder: en process är en kärna som håller alla GPU-resurser, trådar är warps utan preemption, att vänta på I/O är att busy-waits på en flagga, DMA är TMA som spåras av mbarrier, meddelanden är vanliga lagringar över NVLink.

Utkastmodellen blev också snabbare, och det är halva vinsten.

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.

Teamet tillämpade samma megakernel-behandling på draften: fem draft-lager körs vart och ett som en enda launch, plus tre andra, vilket tar draften från 175 kernels till 8 och från 1,16 ms per steg till 440 µs, en 2,6× förbättring på den halvan av steget. Acceptanssiffrorna är värda att ställa bredvid den närmaste offentliga jämförelsen som inlägget erbjuder.

• 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

• Varför 16 – sexton tokens kostar mer att verifiera eftersom de aktiverar fler experter, men teamet tyckte att det var lättare att nå 6,6 med ett bredare träd än att nå samma längd med en smalare kedja

• Utkaststack — 8000 rader CUDA och PTX för den stora modellen, med koden publicerad på facebookresearch/sgl-muse-mega

Läs jämförelsen noggrant, eftersom det inte är en benchmark mellan två modeller. Det är en jämförelse av två spekulativa avkodningskonfigurationer, en publicerad av Xiaomi och en av SGLang-teamet, på olika hårdvara och olika arbetsbelastningar. Det enda påståendet den stöder är att ett bredare verifieringsträd är en framkomlig väg till samma acceptanslängd, vilket är ett designpåstående, inte en rangordning.

Vad detta nummer är, och vad det inte är

Det här är den del som en kapacitetsplan måste få rätt på, och det är den del som inlägget självt är uppriktigt om. Siffran 1 109 är medianen för 20 avkodningssteg på en enda GSM8K-förfrågan, vid tensorparallellitetsgrad 8, på åtta GB300:or, med batchstorlek ett och spekulativ avkodning aktiverad. Inlägget säger rakt ut att megakärnan i dag bara stöder batchstorlek ett, och att utöka den till större batchar är nästa arbetsmoment.

Dessa begränsningar tar var och en bort ett annat påstående. Batchstorlek ett innebär att detta inte är ett genomströmningsresultat i den mening ett servingteam använder ordet: det finns ingen batchvinst i det, ingen kö, ingen blandning av promptlängder. En enda datamängd innebär att siffran är knuten till GSM8K:s prompt- och utdatadistribution, och till ett uppmätt routingmönster med 49 experter per lager som författarna sedan använde som syntetisk routing i sin egen profilering. Tjugo steg för en enda förfrågan innebär att variansen begränsas av antalet steg, inte av antalet förfrågningar.

Och en enskild förfrågan är det värsta fallet för överlappning, vilket författarna själva säger: "vid en förfrågan exponeras varje väntan; med fler finns det nästan alltid färdigt arbete." Så megakernelns 2,1× är en mätning av regimen där latensdöljning är svårast, och inlägget är tydligt med att större batchar borde vara lättare snarare än svårare. Det är ett argument till resultatets fördel, och det är fortfarande ett argument snarare än en mätning.

Det som går att veta i dag är alltså begränsat och verkligt: koden finns i ett offentligt kodförråd, inlägget beskriver mekanismen i tillräcklig detalj för att kunna kontrolleras, den interna aritmetiken stämmer över tre oberoende kvoter och författarna har dokumenterat de tre omkostnader som begränsar ytterligare vinster. Det som inte går att veta är huruvida ett oberoende team får 1 109 på sina egna åtta GB300, om angreppssättet klarar en verklig batch och om det landar i en släppt SGLang-version – inlägget hävdar inte att så är fallet, och kodförrådet är en forskningsartefakt snarare än ett paket som går att installera. Den jämförelse som skulle avgöra mest är mot en optimerad baslinje på samma hårdvara, eftersom "SGLang:s standardkedja" är ett rörligt mål och teamet som skrev megakerneln också optimerar den baslinje den slår.

Hur det hänger samman med de 169 tokens per sekund som du kan hyra idag

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.

Gapet är det mest användbara i den här historien för alla som inte redan äger hårdvaran. Om en slutpunkt svarar med 169 token per sekund och en handbyggd stack med åtta GPU:er svarar med 1 109, är skillnaden inte ett skäl att bygga stacken. Åtta GB300 är ett kapitalköp som mäts i miljoner, kärnan är forskningskod, och begränsningen med batchstorlek ett innebär att stacken för närvarande är bra på exakt den arbetsbelastning som en hostad slutpunkt redan hanterar väl: en begäran i taget. Där megakärnan verkligen förändrar kalkylen är för ett team som redan har köpt GPU:erna för träning eller för en arbetsbelastning med krav på datalokalisering, och som nu frågar vilken utnyttjandegrad de kan få ut av dem vid låg batch. För det teamet är 2,1× på samma kisel med samma numerik ett verkligt tal med en verklig faktura kopplad till sig, och fakturan är redan betald.

För alla andra är den ärliga tolkningen att det här inlägget är en signal om vart inferensinfrastrukturen är på väg, inte en produkt som du kan ta i bruk den här veckan. Inriktningen stämmer överens med det som samma labb har publicerat hela året: färre, större kärnor, mer av schemaläggningsbeslutet flyttat in i kärnan, och GPU:n behandlad som en maskin som ska programmeras snarare än en uppsättning biblioteksanrop som ska sekvenseras.

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.

Vad skulle förändra det här stycket?

Fyra saker, i fallande ordning efter hur mycket de skulle förändra bilden.

• En batch större än ett. Författarna säger att större batchar är nästa steg och hävdar att de är enklare, eftersom det med fler tokens i omlopp finns en del av föregående ops utdata redan klara medan resten fortfarande anländer. I samma stund som en batchad siffra finns upphör jämförelsen att vara "en ström mot en ström" och blir i stället ett påstående om genomströmning som ett servingteam kan använda.

• En oberoende reproduktion. Varje siffra här kommer från författarnas egen testuppställning, på deras egen hårdvara, med deras egen syntetiska routning för profileringskörningen. Det är normalt för kernelarbete och det är fortfarande en enda källa. Att ett andra team matchar 1,109, eller misslyckas med det, är det faktum som skulle avgöra saken.

• En släppt SGLang-build som innehåller det. Inlägget publicerar kod och beskriver designen; det säger inte att megakernel har ersatt standardavkodningsvägen för den här modellen i en levererad version. Fram till dess innebär ett införande att man inför ett forskningsrepo.

• Att GEMM-2-till-combine-gapet stängs. Författarna har redan pekat ut 12.5 µs ledig tid per block där som nästan hälften av det som återstår. Om den kedjan åtgärdas ger samma hårdvara en meningsfullt bättre stegtid, och 2.1× blir golvet för ett större tal snarare än rubriken.

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.