Muse Spark Open: 1.109 Token pro Sekunde ist eine Kernel-Änderung, keine Modelländerung

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.

Die Baseline, die niemand benchmarkt: 1,450 Starts pro 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.

Das ist der Preis, den der Megakernel angreift, und die Autoren greifen zu einer Analogie aus den 1960er-Jahren, um ihn zu erklären: Frühe Maschinen führten Batch-Jobs aus, wobei ein Job die gesamte Maschine belegte, und der Prozessor lag brach, während ein Band rotierte. Multiprogramming behob das, indem mehrere Jobs resident blieben und zu einem anderen gewechselt wurde, wann immer der aktuelle blockierte. Ein Megakernel ist diese Idee, übertragen auf die GPU – ein Kernel bleibt auf jedem SM resident und führt viele Jobs gleichzeitig aus, sodass Warps mit der Arbeit für den nächsten Op beginnen können, bevor der vorherige Op abgeschlossen ist.

Der gemessene Effekt – aus der Abbildung im Abschnitt zur Leistung – ist ein Dekodierschritt, der von 12,9 ms auf 5,9 ms geht. Diese Abbildung teilt den Schritt in seine Draft- und Verify-Hälften auf: 1,16 ms und 10,9 ms mit 1.450 Kernen in der Baseline, 0,44 ms und 5,3 ms mit 97 Kernen im Megakernel. Zwölf Komma neun auf fünf Komma neun sind etwa 2,2×, und 1.109 geteilt durch 517 sind etwa 2,1×, sodass die beiden zentralen Verhältnisse zueinander passen. Es gibt eine dritte Prüfung, die jedem offensteht, der bereit ist, die Division durchzuführen: Der Schritt verifiziert 16 Token und akzeptiert 6,6, sodass 6,6 Token in 5,9 ms auf rund 1.120 Token pro Sekunde kommen. Die Arithmetik ist in sich stimmig, was das Erste ist, was bei einer solchen Zahl zu klären ist, und nicht das Letzte.

Eines verraten die Zahlen nicht: wie hoch die Geschwindigkeit pro GPU ausfällt. Der Beitrag berichtet über den Gesamtdurchsatz für das gesamte Acht-GPU-Deployment. Durch acht geteilt sind das etwa 139 Token pro Sekunde pro GPU – meine Rechnung auf Basis ihrer Zahl, keine im Beitrag veröffentlichte Zahl, und nicht vergleichbar mit einer Pro-GPU-Zahl, die Sie für ein anderes Modell auf einer anderen Maschine zitiert sehen könnten.

Wo die Kernel-Grenze verläuft und warum sie dort endet

Der Megakernel verschluckt nicht das gesamte Modell. Jeder Launch umfasst die Attention einer Decoder-Schicht bis hin zur QKV-Projektion und zum KV-Schreibvorgang der nächsten Schicht, was eine bewusste Grenze und keine Einschränkung von nerve ist. Dass es funktioniert, liegt daran, dass das meiste, was die Attention einer Schicht liest – die KV-Historie –, in früheren Schritten geschrieben wurde und von nichts abhängt, das sich noch in Ausführung befindet, sodass ein neuer Launch mit dem Laden beginnen kann, sobald es verfügbar ist, noch bevor die Ausgabe der vorherigen Schicht existiert. Die erste und die letzte Schicht haben etwas andere Operationen, und der Beitrag erwähnt das ebenfalls.

Bei einem einzelnen Launch ist die Struktur fest vorgegeben: ein Threadblock aus zwölf Warps auf jeder der 144 SMs, gestartet in Clustern von vier, wobei verteilter Shared Memory genutzt wird, wo eine Reduktion mit geringerer Latenz verfügbar ist. Die zwölf Warps sind nicht austauschbar. Warp 0 lädt KV und Gewichte und verwendet seine freie Zeit darauf, die Gewichte der nächsten Operation vorab zu laden. Warps 1 und 2 übernehmen Aktivierungen und Tensor-Core-Arbeit. Warps 3 bis 11 teilen sich Attention, die dichten GEMMs, das Experten-GEMM sowie die Empfänge für All-Reduce und All-Gather. In den Sliding-Window-Schichten laden Warps 4 bis 11 ihr eigenes KV, statt auf Warp 0 zu warten.

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.

Dieses Detail ist wichtiger, als es aussieht. Inference-Teams akzeptieren üblicherweise, dass ein schnellerer Kernel leicht abweichende Zahlen liefert, und justieren stillschweigend ihre Bewertungsschwellen neu. Ein Kernel, der bitgenau mit der Parallelisierung auf der Trainingsseite übereinstimmt, beseitigt dieses Argument vollständig – und das ist genau die Art von Eigenschaft, die dafür sorgt, dass eine Optimierung übernommen und nicht bloß bewundert wird.

Drei Rechnungen, die ein so großer Kern bezahlen muss

Der Megakernel ist nicht kostenlos. Der Beitrag ist ungewöhnlich offen, was die Overheadkosten angeht, und diese sind der Teil, den man zweimal lesen sollte, denn sie sind auch der Teil, der verrät, wie viel Spielraum noch bleibt.

Der erste Posten ist der Instruction-Cache. Ein SM hält etwa 128 KB an Instruktionen im Cache, und diese Kernel umfassen jeweils 215 bis 314 KB Maschinencode, weil jeder von ihnen jede Phase einer Schicht enthält. Ein Launch bringt nur die ersten 32 bis 64 KB in die Nähe der SMs; alles jenseits dieses Fensters wird bei der ersten Verwendung Zeile für Zeile abgerufen. Die Autoren führten dazu ein sauberes Experiment durch und berichteten darüber: Zwei 100-KB-Blöcke, die allein jeweils 12,4 µs benötigten, brauchten 19,5 µs, wenn sie hintereinandergehängt wurden, und die Verkleinerung eines Blocks auf 20 KB brachte den anderen wieder auf seine Sologeschwindigkeit. Instruktionen laufen durch denselben L2-Cache wie Daten, sodass ein Kernel, der den Instruction-Cache übersteigt, schließlich mit seinen eigenen Weight-Loads konkurriert.

Die zweite Rechnung ist das Vorabrufen von Gewichten, und der Beitrag behandelt es als Platzierungsproblem mit drei Antworten. Shared Memory bekommt einen Ring von 5 × 32 KB pro SM, etwa 23 MB über die GPU hinweg. Der Rest landet im L2, und der L2 von GB300 fasst etwa 135 MB, wobei alles im L2 einen zweiten, kürzeren Weg in den Shared Memory benötigt und der L2 mit allem anderen, was gerade läuft, geteilt wird. Die beiden stapeln sich: Gegen Ende einer Schicht hält der Ring bereits die ersten qkv-Gewichte der nächsten Schicht. Darüber hinaus liest Warp 3, zu welchen Experten der vorherige Schritt geroutet hat, und zieht die Anteile dieser GPU an den 14 am häufigsten genutzten in den L2, etwa 50 MB pro GPU, in der Annahme, dass aufeinanderfolgende Schritte ähnlich routen. Diese Wette ist etwa 0,2 ms eines 5,3-ms-Schritts wert und kann Bandbreite kosten, wenn sie falsch liegt.

Der dritte Kostenpunkt ist die Synchronisation, und für sie gibt es keine saubere Antwort. Wenn ein Warp über NVLink auf eine andere GPU wartet, gibt es in der Hardware keinen Interrupt und kein Wait-on-Address; der Warp pollt ein Flag, bis es sich ändert. Die Mikrobenchmarks des Beitrags ergeben einen GPU-Scope-Acquire-Poll bei 141 bis 145 ns, ein Einweg-Flag zwischen zwei GPUs auf demselben Knoten bei etwa 1,9 µs und einen Poll, der hinter Bulk-Kopien in der Warteschlange steht, bei 1,0 bis 2,6 µs. Fences kosten 0,4 µs im Leerlauf und 3,6 bis 4,0 µs, wenn TMA-Ladevorgänge unterwegs sind. Selbst bei perfekt verteilter Arbeit können GPUs bis zu 4 µs hintereinander zurückliegen, und jede Barriere in der Layer wartet auf die langsamste. Der Megakernel ersetzt Fences durch Sentinel-Werte, die in gewöhnliche Store-Slots geschrieben werden – ein mit Daten gefülltes Flag, +0/−0 für das All-Reduce, All-Ones-NaN für die bf16-Gathers – und setzt die im Leerlauf befindliche Kopie jedes Slots zurück, während die andere in Benutzung ist.

Das sind die Zahlen, die das Ergebnis eher nach Ingenieursarbeit als nach Magie wirken lassen: 4 µs GPU-Skew gegenüber einem 5,9-ms-Schritt sind klein, aber vier Synchronisationsstopps pro Layer, multipliziert über 62 Layer, sind genau die Art von Overhead, die einen 2×-Gewinn auffrisst, wenn man es falsch macht.

Wofür die 5,9 ms noch aufgewendet werden

Der Beitrag enthält seine eigene Kritik, und die nützlichste Aussage darin betrifft die Lücke nach dem zweiten GEMM. Blöcke beenden dieses GEMM irgendwann zwischen 55 und 64 µs innerhalb des Layers, 10. bis 90. Perzentil, aber die Combine-Phase wartet auf die Ausgabe jeder GPU, sodass sie nicht vor 65 µs beginnen kann. Vom Ende von GEMM 2 bis zum Ende des Layers sitzt ein Block im Durchschnitt etwa 12,5 µs ohne Arbeit da — fast die Hälfte seiner gesamten Leerlaufzeit. Diese Kette, nicht der Kernel-Launch-Overhead, den der Megakernel bereits beseitigt hat, ist das, was die Autoren als Nächstes untersuchen wollen.

Sie erklären auch, warum die Attention der nächsten Schicht nicht einfach früher beginnen kann. PDL startet sie bei 79 µs, etwa 4 µs bevor der erste Block endet, doch seine Blöcke treffen erst ein, wenn die vorherigen ihre SMs bei 84 µs verlassen; ein früherer Start würde also nichts früher beginnen. Die Behelfslösung besteht darin, seine alten KV-Zeilen bei 65 µs in den L2 vorab zu laden, wo sie warten. Das ist dasselbe Muster wie beim Weight-Prefetch: Wenn man eine Wartezeit nicht beseitigen kann, verlagert man Arbeit hinein.

Profiling ist das andere strukturelle Problem – und ein gutes: Der Kernel erscheint jedem externen Profiler wie ein langer Balken pro Layer, also misst er sich selbst. Lane 0 jedes Blocks schreibt über das globaltimer-Register einen Zeitstempel in ein globales Array, und die Autoren beziffern die Kosten dieser Instrumentierung auf etwa 2,5 µs pro Layer-Launch, 2,8 % eines Schritts, wobei die Ausgaben bitidentisch bleiben und ein Build-Flag sie herauskompilieren kann. Ein Kernel, der seinen eigenen Profiler mitführen muss, ist ein Kernel, der zu einem Betriebssystem geworden ist – genau die Einordnung, die der Beitrag verwendet: Ein Prozess ist ein Kernel, der alle GPU-Ressourcen hält, Threads sind Warps ohne Preemption, das Warten auf E/A ist Busy-Waiting auf ein Flag, DMA ist TMA, verfolgt durch mbarrier, Nachrichten sind einfache Stores über NVLink.

Auch das Draft-Modell wurde schneller, und das ist die halbe Miete.

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.

Das Team wandte dieselbe Megakernel-Behandlung auf den Draft an: fünf Draft-Schichten werden jeweils als ein Launch ausgeführt, plus drei weitere, wodurch der Draft von 175 Kernels auf 8 und von 1,16 ms pro Schritt auf 440 µs reduziert wird – eine 2,6-fache Verbesserung für diese Hälfte des Schritts. Es lohnt sich, die Akzeptanzzahlen neben den nächstliegenden öffentlichen Vergleich zu stellen, den der Beitrag bietet.

• 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

• Warum 16 — sechzehn Tokens teurer zu verifizieren sind, weil sie mehr Experten aktivieren, das Team es jedoch einfacher fand, mit einem breiteren Baum 6.6 zu erreichen als mit einer schmaleren Kette dieselbe Länge zu erreichen

• Draft-Stack — 8000 Zeilen CUDA- und PTX-Code für das große Modell, wobei der Code unter facebookresearch/sgl-muse-mega veröffentlicht wurde

Lies den Vergleich sorgfältig, denn er ist kein Benchmark zwischen zwei Modellen. Er ist ein Vergleich zweier Speculative-Decoding-Konfigurationen, eine von Xiaomi und eine vom SGLang-Team veröffentlicht, auf unterschiedlicher Hardware und mit unterschiedlichen Workloads. Die einzige Aussage, die er stützt, ist, dass ein breiterer Verifikationsbaum ein gangbarer Weg zur selben Acceptance-Länge ist, was eine Design-Aussage und kein Ranking ist.

Was diese Zahl ist und was sie nicht ist

Der Teil, den ein Kapazitätsplan richtig hinbekommen muss, und genau der Teil, bei dem der Beitrag selbst ehrlich ist. Die Zahl 1,109 ist der Median von 20 Decode-Schritten bei einer einzelnen GSM8K-Anfrage, bei Tensor-Parallelitätsgrad 8, auf acht GB300s, mit Batchgröße eins und aktivierter spekulativer Dekodierung. Der Beitrag sagt klar, dass der Megakernel heute nur Batchgröße eins unterstützt und dass die Erweiterung auf größere Batches der nächste Arbeitsschritt ist.

Jede dieser Einschränkungen entkräftet einen anderen Anspruch. Batch-Größe eins bedeutet, dass dies kein Durchsatzergebnis in dem Sinne ist, wie ein Serving-Team den Begriff verwendet: Es steckt kein Batching-Gewinn darin, keine Warteschlange, keine Mischung aus Prompt-Längen. Ein einziger Datensatz bedeutet, dass die Zahl an die Prompt- und Ausgabeverteilung von GSM8K gebunden ist und an ein gemessenes Routing-Muster von 49 Experten pro Schicht, das die Autoren dann als synthetisches Routing in ihrem eigenen Profiling verwendeten. Zwanzig Schritte einer einzigen Anfrage bedeuten, dass die Varianz durch die Schrittanzahl begrenzt ist, nicht durch die Anzahl der Anfragen.

Und eine einzelne Anfrage ist der Worst Case für Overlap, was die Autoren selbst sagen: „Bei einer Anfrage wird jede Wartezeit offengelegt; bei mehr gibt es fast immer bereite Arbeit.“ Also ist das 2,1× des Megakernels eine Messung des Regimes, in dem Latenzverdeckung am schwierigsten ist, und der Beitrag stellt ausdrücklich klar, dass größere Batches eher einfacher als schwieriger sein sollten. Das ist ein Argument zugunsten des Ergebnisses, und es ist immer noch ein Argument statt einer Messung.

Was sich heute wissen lässt, ist also eng und real: Der Code existiert in einem öffentlichen Repository, der Beitrag beschreibt den Mechanismus ausreichend detailliert, um überprüfbar zu sein, die internen Berechnungen stimmen über drei unabhängige Verhältnisse hinweg überein, und die Autoren haben die drei Overheads dokumentiert, die weitere Gewinne begrenzen. Nicht wissen lässt sich, ob ein unabhängiges Team auf seinen eigenen acht GB300s 1.109 erreicht, ob der Ansatz einen echten Batch übersteht und ob er es in einen veröffentlichten SGLang-Build schafft — der Beitrag behauptet nicht, dass dies der Fall ist, und das Repository ist ein Forschungsartefakt und kein Paket, das man installieren kann. Der Vergleich, der am meisten klären würde, ist der gegen eine optimierte Baseline auf derselben Hardware, da „SGLang's default chain“ ein bewegliches Ziel ist und das Team, das den Megakernel geschrieben hat, auch die Baseline optimiert, die er schlägt.

Wie es mit den 169 Tokens pro Sekunde zusammenhängt, die Sie heute mieten können

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.

Die Lücke ist das Nützlichste in dieser Geschichte für alle, die die Hardware noch nicht besitzen. Wenn ein Endpoint mit 169 Tokens pro Sekunde antwortet und ein selbst gebauter Acht-GPU-Stack mit 1.109, dann ist die Differenz kein Grund, den Stack zu bauen. Acht GB300s sind eine Investitionsausgabe in Millionenhöhe, der Kernel ist Forschungscode, und die Einschränkung auf Batch-Größe eins bedeutet, dass der Stack derzeit genau für die Arbeitslast gut ist, die ein gehosteter Endpoint bereits gut bewältigt: eine Anfrage nach der anderen. Wo der Megakernel die Rechnung wirklich ändert, ist bei einem Team, das die GPUs bereits für Training gekauft hat oder für eine Arbeitslast mit Auflagen zur Datenresidenz, und nun fragt, welche Auslastung es bei niedriger Batch-Größe aus ihnen herausholen kann. Für dieses Team sind 2,1× auf demselben Silizium mit derselben Numerik eine echte Zahl mit einer echten Rechnung dahinter, und die Rechnung ist bereits bezahlt.

Für alle anderen ist die ehrliche Einschätzung, dass dieser Beitrag ein Signal dafür ist, wohin sich die Inferenzinfrastruktur entwickelt, und kein Produkt, das Sie diese Woche einführen können. Die Richtung deckt sich mit dem, was dasselbe Labor das ganze Jahr über veröffentlicht hat: weniger, größere Kernel, ein größerer Teil der Scheduling-Entscheidung in den Kernel verlagert und die GPU als eine Maschine behandelt, die programmiert werden muss, statt als eine Reihe von Bibliotheksaufrufen, die sequenziert werden müssen.

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.

Was würde dieses Stück verändern?

Vier Dinge, in absteigender Reihenfolge danach, wie stark sie das Bild verändern würden.

• Ein Batch größer als eins. Die Autoren sagen, größere Batches seien als Nächstes dran, und argumentieren, sie seien einfacher, weil bei mehr Tokens in Bearbeitung einige der Ausgaben der vorherigen Operation bereits fertig sind, während der Rest noch eintrifft. Sobald eine gebatchte Zahl existiert, ist der Vergleich nicht mehr „ein Stream gegen einen Stream“, sondern beginnt, eine Durchsatzaussage zu sein, die ein Serving-Team verwenden kann.

Eine unabhängige Reproduktion. Jede Zahl hier stammt aus dem eigenen Messaufbau der Autoren, auf ihrer eigenen Hardware, mit ihrem eigenen synthetischen Routing für den Profiling-Durchlauf. Das ist bei Kernel-Arbeit normal, und es ist trotzdem nur eine einzige Quelle. Ein zweites Team, das 1.109 ebenfalls erreicht – oder es nicht erreicht –, wäre die Tatsache, die das entscheiden würde.

• Ein veröffentlichter SGLang-Build, der ihn enthält. Der Beitrag veröffentlicht den Code und beschreibt das Design; er sagt nicht, dass der Megakernel den Standard-Decode-Pfad für dieses Modell in einer ausgelieferten Version ersetzt hat. Bis dahin bedeutet die Übernahme, ein Forschungs-Repository zu übernehmen.

• Das Schließen der GEMM-2-to-combine-Lücke. Die Autoren haben die 12.5 µs Leerlaufzeit pro Block dort bereits als fast die Hälfte des Verbleibenden bezeichnet. Wenn diese Kette behoben wird, erzeugt dieselbe Hardware eine deutlich bessere Schrittzeit, und die 2.1× wird zur Untergrenze einer größeren Zahl statt zur Schlagzeile.

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.