Muse Spark Open: 1.109 token al secondo sono un cambiamento di kernel, non un cambiamento di modello

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.

Il baseline che nessuno prende come benchmark: 1.450 lanci 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.

È questo il costo che il megakernel attacca, e gli autori ricorrono a un'analogia degli anni '60 per spiegarlo: le prime macchine eseguivano job in batch, un job occupava l'intera macchina e il processore restava inattivo mentre un nastro girava. Il multiprogramming lo risolse mantenendo diversi job residenti e passando a un altro ogni volta che quello corrente si bloccava. Un megakernel è quell'idea trasferita alla GPU: un kernel resta residente su ogni SM ed esegue molti job contemporaneamente, così i warp possono iniziare il lavoro dell'operazione successiva prima che l'operazione precedente sia terminata.

L'effetto misurato, dalla figura che accompagna la sezione sulle prestazioni, è un passo di decodifica che passa da 12,9 ms a 5,9 ms. La figura suddivide il passo nelle sue due metà, draft e verifica: 1,16 ms e 10,9 ms con 1.450 kernel nella baseline, 0,44 ms e 5,3 ms con 97 kernel nel megakernel. Da 12,9 a 5,9 è circa 2,2×, e 1.109 su 517 è circa 2,1×, quindi i due rapporti principali concordano tra loro. C'è un terzo controllo a disposizione di chiunque sia disposto a fare la divisione: il passo verifica 16 token e ne accetta 6,6, quindi 6,6 token in 5,9 ms corrispondono a circa 1.120 token al secondo. L'aritmetica è coerente al suo interno, che è la prima cosa da stabilire su un numero come questo, e non l'ultima.

Una cosa che i numeri non dicono è a quanto ammonta la velocità per GPU. Il post riporta il throughput aggregato per l’intera installazione a otto GPU. Diviso per otto, sono circa 139 token al secondo per GPU — è il mio calcolo sul loro dato, non un numero pubblicato dal post, e non è confrontabile con un numero per GPU che potresti vedere citato per un altro modello su un’altra macchina.

Dove si colloca il confine del kernel, e perché si ferma proprio lì

Il megakernel non ingoia l'intero modello. Ogni lancio copre l'attenzione di un layer del decoder fino alla proiezione qkv e alla scrittura KV del layer successivo, un confine deliberato anziché un limite di nerve. Funziona perché gran parte di ciò che l'attenzione di un layer legge — la cronologia KV — è stata scritta in passaggi precedenti e non dipende da nulla in corso, quindi un nuovo lancio può iniziare a caricarla non appena è disponibile, prima che esista l'output del layer precedente. Il primo e l'ultimo layer hanno operazioni leggermente diverse, e anche il post lo nota.

All'interno di un singolo lancio la forma è fissa: un thread block di dodici warp su ciascuna delle 144 SM, avviato in cluster di quattro, usando la memoria condivisa distribuita dove è disponibile una riduzione a latenza inferiore. I dodici warp non sono intercambiabili. Il warp 0 carica KV e pesi e usa il proprio margine per il prefetch dei pesi dell'operazione successiva. I warp 1 e 2 eseguono attivazioni e tensor core. I warp da 3 a 11 suddividono l'attenzione, i GEMM densi, il GEMM esperto e le ricezioni all-reduce e all-gather. Nei layer a finestra scorrevole, i warp da 4 a 11 caricano i propri KV invece di attendere il 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.

Quel dettaglio conta più di quanto sembri. I team di inferenza di solito accettano che un kernel più veloce produca numeri leggermente diversi e ritoccano silenziosamente le proprie soglie di valutazione. Un kernel bitwise identico alla parallelizzazione lato addestramento elimina del tutto quella discussione, ed è il tipo di proprietà che fa sì che un'ottimizzazione venga adottata anziché soltanto ammirata.

Tre bollette che un kernel così grande deve pagare

Il megakernel non è gratis. Il post è insolitamente franco sui sovraccarichi, e sono la parte che vale la pena leggere due volte, perché sono anche la parte che ti dice quanto margine resta.

La prima voce di costo è la cache delle istruzioni. Una SM mette in cache circa 128 KB di istruzioni, e questi kernel occupano da 215 a 314 KB di codice macchina ciascuno, perché ognuno contiene ogni fase di un layer. Un lancio porta solo i primi 32-64 KB vicino alle SM; tutto ciò che va oltre quella finestra viene recuperato linea per linea al primo utilizzo. Gli autori hanno eseguito un esperimento controllato su questo e lo hanno riportato: due blocchi da 100 KB che da soli impiegavano 12,4 µs ciascuno impiegavano 19,5 µs quando messi in sequenza, e riducendo un blocco a 20 KB l'altro tornava alla sua velocità da solo. Le istruzioni passano attraverso la stessa cache L2 dei dati, quindi un kernel che supera la cache delle istruzioni finisce per competere con i propri caricamenti dei pesi.

La seconda voce è il prefetch dei pesi, e il post la tratta come un problema di collocazione con tre risposte. Alla memoria condivisa viene assegnato un ring di 5 × 32 KB per SM, circa 23 MB sull'intera GPU. L2 è dove va il resto, e la L2 di GB300 contiene circa 135 MB, anche se qualsiasi cosa in L2 richiede un secondo viaggio, più breve, verso la memoria condivisa e la L2 è condivisa con tutto il resto in esecuzione. I due si sommano: verso la fine di un layer il ring contiene già i primi pesi qkv del layer successivo. Inoltre, il warp 3 legge verso quali esperti ha indirizzato il passo precedente e porta in L2 le slice di questa GPU dei 14 più usati, circa 50 MB per GPU, scommettendo che passi consecutivi instradino in modo simile. Quella scommessa vale circa 0,2 ms su un passo da 5,3 ms e può costare larghezza di banda quando è sbagliata.

Il terzo conto è la sincronizzazione, ed è quello che non ha una risposta netta. Quando un warp attende un'altra GPU tramite NVLink non esistono interrupt né wait-on-address nell'hardware; il warp fa polling su un flag finché non cambia. I microbenchmark del post collocano un acquire poll con scope GPU a 141–145 ns, un flag unidirezionale tra due GPU sullo stesso nodo a circa 1,9 µs e un poll accodato dietro copie bulk a 1,0–2,6 µs. Le fence costano 0,4 µs da inattive e 3,6–4,0 µs con caricamenti TMA in volo. Anche con un lavoro perfettamente distribuito, le GPU possono restare indietro l'una rispetto all'altra fino a 4 µs, e ogni stop nel layer attende la più lenta. Il megakernel sostituisce le fence con valori sentinella scritti in normali slot di store — un flag riempito di dati, +0/−0 per l'all-reduce, NaN con tutti i bit a uno per i gather bf16 — e resetta la copia inattiva di ciascuno slot mentre l'altra è in uso.

Questi sono i numeri che fanno sembrare il risultato più ingegneria che magia: 4 µs di skew della GPU rispetto a un passo di 5,9 ms è poco, ma quattro fermate di sincronizzazione per layer moltiplicate su 62 layer sono esattamente il tipo di overhead che divora un vantaggio di 2× se si sbaglia.

Su cosa vengono ancora spesi i 5,9 ms

Il post include una sua critica, e l'osservazione più utile riguarda il divario dopo il secondo GEMM. I blocchi terminano quel GEMM tra 55 e 64 µs dall'inizio del layer, dal 10° al 90° percentile, ma il combine attende l'output di ogni GPU, quindi non può iniziare prima di 65 µs. Dalla fine del GEMM 2 alla fine del layer, un blocco resta senza lavoro per circa 12,5 µs in media — quasi metà del suo tempo di inattività totale. È quella catena, non l'overhead di lancio dei kernel che il megakernel ha già eliminato, che gli autori dicono di stare esaminando adesso.

Spiegano anche perché l'attenzione del layer successivo non possa semplicemente iniziare prima. PDL la avvia a 79 µs, circa 4 µs prima che il primo blocco esca, ma i suoi blocchi si insediano solo quando i precedenti lasciano le loro SM a 84 µs, quindi avviarla prima non anticiperebbe nulla. La soluzione è precaricare le sue vecchie righe KV in L2 a 65 µs, dove restano in attesa. È lo stesso schema del prefetch dei pesi: quando non puoi eliminare un'attesa, ci sposti dentro il lavoro.

Il profiling è l'altro problema strutturale, ed è un problema interessante: a qualsiasi profiler esterno il kernel appare come un'unica lunga barra per layer, quindi si cronometra da solo. La corsia 0 di ogni blocco scrive un timestamp in un array globale tramite il registro globaltimer, e gli autori misurano il costo di questa strumentazione in circa 2,5 µs per lancio di layer, il 2,8% di uno step, con output lasciati identici bit per bit e un flag di build per escluderla in compilazione. Un kernel che deve portarsi dietro il proprio profiler è un kernel che è diventato un sistema operativo, che è esattamente l'inquadramento usato dal post: un processo è un kernel che detiene tutte le risorse GPU, i thread sono warp senza preemption, attendere I/O è busy-waiting su un flag, il DMA è TMA tracciato da mbarrier, i messaggi sono semplici store attraverso NVLink.

Anche il modello draft è diventato più veloce, e questo è metà del successo.

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.

Il team ha applicato lo stesso trattamento megakernel al draft: cinque layer del draft vengono eseguiti ciascuno come un unico lancio, più altri tre, il che porta il draft da 175 kernel a 8 e da 1,16 ms per step a 440 µs, un miglioramento di 2,6× su quella metà dello step. Vale la pena mettere i numeri di accettazione accanto al confronto pubblico più vicino offerto dal post.

• 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

• Perché 16 — sedici token costano di più da verificare perché attivano più esperti, ma il team ha trovato più facile raggiungere 6,6 con un albero più ampio che raggiungere la stessa lunghezza con una catena più stretta

• Stack di bozza — 8000 righe di CUDA e PTX per il modello grande, con il codice pubblicato su facebookresearch/sgl-muse-mega

Leggi attentamente il confronto, perché non è un benchmark tra due modelli. È un confronto tra due configurazioni di decodifica speculativa, una pubblicata da Xiaomi e una dal team SGLang, su hardware diversi e carichi di lavoro diversi. L'unica affermazione che supporta è che un albero di verifica più ampio è una via percorribile per raggiungere la stessa lunghezza di accettazione, il che è un'affermazione di progettazione, non una classifica.

Che cos'è questo numero e che cosa non è

Ecco la parte che un piano di capacità deve centrare, ed è la parte su cui il post stesso è onesto. La cifra di 1.109 è la mediana di 20 passi di decodifica su una singola richiesta GSM8K, con tensor parallel degree 8, su otto GB300, con batch size uno e decodifica speculativa attiva. Il post afferma chiaramente che oggi il megakernel supporta solo batch size uno, e che estenderlo a batch più grandi è il prossimo lavoro da fare.

Ciascuno di questi vincoli elimina un’affermazione diversa. Una dimensione del batch pari a uno significa che questo non è un risultato di throughput nel senso in cui lo intende un team di serving: non c’è alcun guadagno dal batching, nessuna coda, nessuna mescolanza di lunghezze dei prompt. Un solo dataset significa che il numero è legato alla distribuzione dei prompt e degli output di GSM8K, e a un pattern di routing misurato di 49 esperti per layer che gli autori hanno poi usato come routing sintetico nel proprio profiling. Venti step di una singola richiesta significano che la varianza è limitata dal numero di step, non dal numero di richieste.

E una singola richiesta è il caso peggiore per la sovrapposizione, come dicono gli stessi autori: «con una richiesta ogni attesa è esposta; con più richieste, c’è quasi sempre lavoro pronto». Quindi il 2,1× del megakernel è una misura del regime in cui nascondere la latenza è più difficile, e il post è esplicito sul fatto che batch più grandi dovrebbero essere più facili anziché più difficili. È un argomento a favore del risultato, ed è comunque un argomento e non una misura.

Ciò che è conoscibile oggi, dunque, è ristretto e reale: il codice esiste in un repository pubblico, il post descrive il meccanismo con dettaglio sufficiente per essere verificato, l’aritmetica interna è coerente tra tre rapporti indipendenti, e gli autori hanno documentato i tre overhead che limitano ulteriori guadagni. Ciò che non è conoscibile è se un team indipendente ottenga 1,109 sulle proprie otto GB300, se l’approccio sopravviva a un batch reale e se approdi in una build SGLang rilasciata — il post non afferma che sia accaduto, e il repository è un artefatto di ricerca più che un pacchetto installabile. Il confronto che risolverebbe più di ogni altro è contro una baseline ottimizzata sullo stesso hardware, poiché la "catena predefinita di SGLang" è un bersaglio mobile e il team che ha scritto il megakernel mette a punto anche la baseline che batte.

Come si rapporta ai 169 token al secondo che puoi noleggiare oggi

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.

Il divario è la cosa più utile in questa storia per chiunque non possieda già l'hardware. Se un endpoint risponde a 169 token al secondo e uno stack da otto GPU costruito a mano risponde a 1.109, la differenza non è un motivo per costruire lo stack. Otto GB300 sono un acquisto in conto capitale dell'ordine di milioni, il kernel è codice di ricerca, e il vincolo di batch size uno significa che allo stato attuale lo stack è efficiente esattamente sul carico di lavoro che un endpoint ospitato già gestisce bene: una richiesta alla volta. È per un team che ha già comprato le GPU per l'addestramento o per un carico di lavoro con vincoli di residenza dei dati, e che ora si chiede quale utilizzo possa ricavarne a batch basso, che il megakernel cambia davvero i calcoli. Per quel team, 2,1× sullo stesso silicio con la stessa numerica è un numero reale con una fattura reale allegata, e la fattura è già pagata.

Per tutti gli altri, la lettura onesta è che questo post è un segnale di dove sta andando l'infrastruttura di inferenza, non un prodotto che si possa adottare questa settimana. La direzione è coerente con ciò che lo stesso laboratorio ha pubblicato per tutto l'anno: kernel meno numerosi e più grandi, una maggiore parte della decisione di scheduling spostata all'interno del kernel e la GPU trattata come una macchina da programmare anziché come un insieme di chiamate di libreria da mettere in sequenza.

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.

Cosa cambierebbe questo pezzo

Quattro cose, in ordine decrescente di quanto cambierebbero il quadro.

• Un batch maggiore di uno. Gli autori affermano che i batch più grandi sono il prossimo passo e sostengono che siano più semplici, perché con più token in elaborazione alcuni output dell'operazione precedente sono pronti mentre gli altri stanno ancora arrivando. Nel momento in cui esiste un numero relativo a un batch, il confronto smette di essere "un flusso contro un flusso" e inizia a essere un'affermazione di throughput che un team addetto al serving può usare.

• Una riproduzione indipendente. Ogni cifra qui riportata proviene dall'harness degli stessi autori, sul loro hardware, con il loro routing sintetico per la passata di profilazione. È normale per il lavoro sui kernel ed è comunque un'unica fonte. Un secondo team che eguaglia 1,109, o che non vi riesce, è il dato che chiarirebbe la questione.

• Una build rilasciata di SGLang che lo include. Il post pubblica il codice e descrive il design; non afferma che il megakernel abbia sostituito il percorso di decodifica predefinito per questo modello in una versione distribuita. Finché ciò non avviene, adottarlo significa adottare un repository di ricerca.

• La chiusura del divario GEMM-2-to-combine. Gli autori hanno già indicato 12,5 µs di tempo di inattività per blocco in quel punto come quasi la metà di quanto resta. Se quella catena viene sistemata, lo stesso hardware produce un tempo per passo sensibilmente migliore, e il 2,1× diventa il limite inferiore di un numero più grande anziché il dato di punta.

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.