Muse Spark Open: 1,109 Tokens a Second Is a Kernel Change, Not a Model Change

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.

The baseline nobody benchmarks: 1,450 launches 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.

That is the cost the megakernel attacks, and the authors reach for a 1960s analogy to explain it: early machines ran batch jobs, one job owning the whole machine, and the processor sat idle while a tape spun. Multiprogramming fixed it by keeping several jobs resident and switching to another whenever the current one blocked. A megakernel is that idea moved to the GPU — one kernel stays resident on every SM and runs many jobs at once, so warps can start next-op work before the previous op has finished.

The measured effect, from the figure accompanying the performance section, is a decode step going from 12.9 ms to 5.9 ms. That figure breaks the step into its draft and verify halves: 1.16 ms and 10.9 ms with 1,450 kernels in the baseline, 0.44 ms and 5.3 ms with 97 kernels in the megakernel. Twelve-point-nine to five-point-nine is about 2.2×, and 1,109 over 517 is about 2.1×, so the two headline ratios agree with each other. There is a third check available to anyone willing to do the division: the step verifies 16 tokens and accepts 6.6, so 6.6 tokens in 5.9 ms works out to roughly 1,120 tokens per second. The arithmetic is internally consistent, which is the first thing to establish about a number like this and not the last.

One thing the numbers do not say is where the speed lands per GPU. The post reports aggregate throughput for the whole eight-GPU deployment. Divided by eight, that is about 139 tokens per second per GPU — my arithmetic on their figure, not a number the post publishes, and not comparable to a per-GPU number you might see quoted for another model on another machine.

Where the kernel boundary sits, and why it stops there

The megakernel does not swallow the whole model. Each launch covers one decoder layer's attention through the next layer's qkv projection and KV write, which is a deliberate boundary rather than a limitation of nerve. The reason it works is that most of what a layer's attention reads — the KV history — was written in earlier steps and depends on nothing in flight, so a new launch can begin loading it the instant it is available, before the previous layer's output exists. The first and last layers have slightly different operations, and the post notes that too.

Inside one launch the shape is fixed: one thread block of twelve warps on each of 144 SMs, launched in clusters of four, using distributed shared memory where a lower-latency reduction is available. The twelve warps are not interchangeable. Warp 0 loads KV and weights and spends its slack prefetching the next operation's weights. Warps 1 and 2 run activations and tensor cores. Warps 3 through 11 split attention, the dense GEMMs, the expert GEMM and the all-reduce and all-gather receives. In the sliding-window layers, warps 4 through 11 load their own KV rather than waiting on 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.

That detail matters more than it looks. Inference teams usually accept that a faster kernel produces slightly different numbers and quietly retune their evaluation thresholds. A kernel that is bitwise equal to the training-side parallelisation removes that argument entirely, which is the sort of property that gets an optimisation adopted rather than merely admired.

Three bills a kernel this big has to pay

The megakernel is not free. The post is unusually candid about the overheads, and they are the part worth reading twice, because they are also the part that tells you how much room is left.

The first bill is instruction cache. An SM caches roughly 128 KB of instructions, and these kernels are 215 to 314 KB of machine code each, because each one holds every phase of a layer. A launch brings only the first 32 to 64 KB close to the SMs; everything past that window is fetched line by line on first use. The authors ran a clean experiment on this and reported it: two 100 KB blocks that each took 12.4 µs alone took 19.5 µs when stitched back to back, and shrinking one block to 20 KB put the other back at its solo speed. Instructions travel through the same L2 cache as data, so a kernel that outgrows the instruction cache ends up competing with its own weight loads.

The second bill is weight prefetch, and the post treats it as a placement problem with three answers. Shared memory gets a ring of 5 × 32 KB per SM, about 23 MB across the GPU. L2 is where the rest goes, and GB300's L2 holds about 135 MB, though anything in L2 needs a second, shorter trip into shared memory and L2 is shared with everything else in flight. The two stack: near the end of a layer the ring already holds the next layer's first qkv weights. On top of that, warp 3 reads which experts the previous step routed to and pulls this GPU's slices of the 14 most-used ones into L2, about 50 MB per GPU, betting that consecutive steps route alike. That bet is worth about 0.2 ms out of a 5.3 ms step, and can cost bandwidth when it is wrong.

The third bill is synchronisation, and it is the one with no clean answer. When a warp waits on another GPU over NVLink there is no interrupt and no wait-on-address in hardware; the warp polls a flag until it changes. The post's microbenchmarks put a GPU-scope acquire poll at 141 to 145 ns, a one-way flag between two GPUs on the same node at about 1.9 µs, and a poll queued behind bulk copies at 1.0 to 2.6 µs. Fences cost 0.4 µs when idle and 3.6 to 4.0 µs with TMA loads in flight. Even with perfectly distributed work, GPUs can lag each other by up to 4 µs, and every stop in the layer waits for the slowest one. The megakernel replaces fences with sentinel values written into ordinary store slots — a data-filled flag, +0/−0 for the all-reduce, all-ones NaN for the bf16 gathers — and resets the idle copy of each slot while the other is in use.

These are the numbers that make the result feel like engineering rather than magic: 4 µs of GPU skew against a 5.9 ms step is small, but four synchronisation stops per layer multiplied across 62 layers is exactly the kind of overhead that eats a 2× win if you get it wrong.

What the 5.9 ms is still spent on

The post includes its own critique, and the most useful line in it is about the gap after the second GEMM. Blocks finish that GEMM anywhere between 55 and 64 µs into the layer, 10th to 90th percentile, but the combine waits for every GPU's output, so it cannot start until 65 µs. From the end of GEMM 2 to the end of the layer a block sits with no work for about 12.5 µs on average — almost half of its total idle time. That chain, not the kernel launch overhead the megakernel already eliminated, is what the authors say they are looking at next.

They also explain why the next layer's attention cannot simply start earlier. PDL launches it at 79 µs, about 4 µs before the first block exits, but its blocks only land once the previous ones leave their SMs at 84 µs, so launching sooner would start nothing sooner. The workaround is to prefetch its old KV rows into L2 at 65 µs, where they wait. That is the same pattern as the weight prefetch: when you cannot remove a wait, you move work into it.

Profiling is the other structural problem, and it is a good one: the kernel looks like one long bar per layer to any external profiler, so it times itself. Lane 0 of each block writes a timestamp to a global array through the globaltimer register, and the authors measure the cost of that instrumentation at about 2.5 µs per layer launch, 2.8% of a step, with outputs left bit-identical and a build flag to compile it out. A kernel that has to carry its own profiler is a kernel that has become an operating system, which is exactly the framing the post uses: a process is one kernel holding all GPU resources, threads are warps with no preemption, waiting for I/O is busy-waiting on a flag, DMA is TMA tracked by mbarrier, messages are plain stores across NVLink.

The draft model got faster too, and that is half the win

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.

The team applied the same megakernel treatment to the draft: five draft layers each run as one launch, plus three others, which takes the draft from 175 kernels to 8 and from 1.16 ms per step to 440 µs, a 2.6× improvement on that half of the step. The acceptance numbers are worth putting next to the closest public comparison the post offers.

• 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

• Why 16 — sixteen tokens cost more to verify because they activate more experts, but the team found it easier to reach 6.6 with a wider tree than to reach the same length with a narrower chain

• Draft stack — 8000 lines of CUDA and PTX for the big model, with the code published at facebookresearch/sgl-muse-mega

Read the comparison carefully, because it is not a benchmark between two models. It is a comparison of two speculative-decoding configurations, one published by Xiaomi and one by the SGLang team, on different hardware and different workloads. The only claim it supports is that a wider verification tree is a viable route to the same acceptance length, which is a design claim, not a ranking.

What this number is, and what it is not

Here is the part a capacity plan has to get right, and it is the part the post itself is honest about. The 1,109 figure is the median of 20 decode steps on a single GSM8K request, at tensor parallel degree 8, on eight GB300s, with batch size one and speculative decoding on. The post states plainly that the megakernel today supports only batch size one, and that extending it to larger batches is the next piece of work.

Those constraints each remove a different claim. Batch size one means this is not a throughput result in the sense a serving team uses the word: there is no batching gain in it, no queue, no mixture of prompt lengths. One dataset means the number is tied to GSM8K's prompt and output distribution, and to a measured routing pattern of 49 experts per layer that the authors then used as synthetic routing in their own profiling. Twenty steps of one request means the variance is bounded by step count, not by request count.

And a single request is the worst case for overlap, which the authors say themselves: "at one request every wait is exposed; with more, there is almost always ready work." So the megakernel's 2.1× is a measurement of the regime where latency hiding is hardest, and the post is explicit that larger batches should be easier rather than harder. That is an argument in the result's favour, and it is still an argument rather than a measurement.

What is knowable today, then, is narrow and real: the code exists in a public repository, the post describes the mechanism in enough detail to be checked, the internal arithmetic agrees across three independent ratios, and the authors have documented the three overheads that bound further gains. What is not knowable is whether an independent team gets 1,109 on its own eight GB300s, whether the approach survives a real batch, and whether it lands in a released SGLang build — the post does not claim it has, and the repository is a research artefact rather than a package you can install. The comparison that would settle the most is against a tuned baseline on the same hardware, since "SGLang's default chain" is a moving target and the team that wrote the megakernel also tunes the baseline it beats.

How it relates to the 169 tokens per second you can rent today

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.

The gap is the most useful thing in this story for anyone who does not already own the hardware. If an endpoint answers at 169 tokens per second and a hand-built eight-GPU stack answers at 1,109, the difference is not a reason to build the stack. Eight GB300s is a capital purchase measured in millions, the kernel is research code, and the batch-size-one constraint means the stack is currently good at exactly the workload a hosted endpoint already handles well: one request at a time. Where the megakernel genuinely changes the calculus is for a team that has already bought the GPUs for training or for a workload with data-residency constraints, and is now asking what utilisation it can get out of them at low batch. For that team, 2.1× on the same silicon with the same numerics is a real number with a real invoice attached, and the invoice is already paid.

For everyone else, the honest read is that this post is a signal about where inference infrastructure is going, not a product you can adopt this week. The direction is consistent with what the same lab has been publishing all year: fewer, larger kernels, more of the scheduling decision moved inside the kernel, and the GPU treated as a machine to be programmed rather than a set of library calls to be sequenced.

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.

What would change this piece

Four things, in descending order of how much they would move the picture.

• A batch greater than one. The authors say larger batches are next and argue they are easier, because with more tokens in flight some of the previous op's outputs are ready while the rest still arrive. The moment a batched number exists, the comparison stops being "one stream versus one stream" and starts being a throughput claim a serving team can use.

• An independent reproduction. Every figure here is from the authors' own harness, on their own hardware, with their own synthetic routing for the profiling pass. That is normal for kernel work and it is still a single source. A second team matching 1,109, or failing to, is the fact that would settle it.

• A released SGLang build that carries it. The post publishes code and describes the design; it does not say the megakernel has replaced the default decode path for this model in a shipped version. Until it does, adopting it means adopting a research repository.

• The GEMM-2-to-combine gap closing. The authors have already named 12.5 µs of per-block idle time there as almost half of what remains. If that chain gets fixed, the same hardware produces a meaningfully better step time, and the 2.1× becomes the floor of a larger number rather than the headline.

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.