Muse Spark Open: 1.109 tokens por segundo é uma mudança de kernel, não uma mudança de modelo

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.

A linha de base que ninguém usa como referência: 1.450 lançamentos por 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.

Esse é o custo que o megakernel ataca, e os autores recorrem a uma analogia dos anos 1960 para explicá-lo: as primeiras máquinas executavam tarefas em lote, com uma única tarefa ocupando a máquina inteira, e o processador ficava ocioso enquanto uma fita girava. A multiprogramação resolveu isso mantendo várias tarefas residentes e alternando para outra sempre que a atual bloqueava. Um megakernel é essa ideia transportada para a GPU — um kernel permanece residente em cada SM e executa muitas tarefas ao mesmo tempo, de modo que as warps podem iniciar o trabalho da próxima operação antes que a operação anterior tenha terminado.

O efeito medido, a partir da figura que acompanha a seção de desempenho, é uma etapa de decodificação passando de 12,9 ms para 5,9 ms. Essa figura divide a etapa em suas metades de rascunho e verificação: 1,16 ms e 10,9 ms com 1.450 kernels na linha de base, 0,44 ms e 5,3 ms com 97 kernels no megakernel. Doze vírgula nove para cinco vírgula nove é cerca de 2,2×, e 1.109 sobre 517 é cerca de 2,1×, então as duas proporções principais concordam entre si. Há uma terceira verificação disponível para qualquer pessoa disposta a fazer a divisão: a etapa verifica 16 tokens e aceita 6,6, então 6,6 tokens em 5,9 ms resulta em aproximadamente 1.120 tokens por segundo. A aritmética é internamente consistente, o que é a primeira coisa a estabelecer sobre um número como esse e não a última.

Uma coisa que os números não dizem é qual é a velocidade resultante por GPU. O post informa a vazão agregada para a implantação inteira de oito GPUs. Dividido por oito, isso dá cerca de 139 tokens por segundo por GPU — uma conta minha em cima do número deles, não um número que o post publique, e não comparável a um número por GPU que você talvez veja citado para outro modelo em outra máquina.

Onde fica a fronteira do kernel e por que ela termina aí

O megakernel não engole o modelo inteiro. Cada lançamento abrange desde a atenção de uma camada do decodificador até a projeção qkv e a gravação de KV da camada seguinte, o que é um limite deliberado, e não uma limitação do nerve. O motivo de funcionar é que a maior parte do que a atenção de uma camada lê — o histórico de KV — foi gravada em etapas anteriores e não depende de nada em andamento, então um novo lançamento pode começar a carregá-la no instante em que estiver disponível, antes que a saída da camada anterior exista. As primeiras e as últimas camadas têm operações ligeiramente diferentes, e o post observa isso também.

Dentro de um único lançamento, a configuração é fixa: um bloco de threads de doze warps em cada um dos 144 SMs, lançado em clusters de quatro, usando memória compartilhada distribuída onde uma redução de menor latência está disponível. Os doze warps não são intercambiáveis. O warp 0 carrega KV e pesos e usa sua folga para pré-buscar os pesos da próxima operação. Os warps {{1}}1 e 2{{/1}} executam ativações e tensor cores. Os warps {{2}}3 a 11{{/2}} dividem atenção, os GEMMs densos, o GEMM de especialistas e os recebimentos de all-reduce e all-gather. Nas camadas de janela deslizante, os warps 4 a 11 carregam seu próprio KV em vez de esperar pelo 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.

Esse detalhe importa mais do que parece. As equipes de inferência geralmente aceitam que um kernel mais rápido produza números ligeiramente diferentes e ajustam discretamente seus limites de avaliação. Um kernel bit a bit igual à paralelização do lado do treinamento elimina essa discussão por completo, e é o tipo de propriedade que faz uma otimização ser adotada, e não apenas admirada.

Três contas que um grão deste tamanho tem de pagar

O megakernel não é de graça. O texto é incomumente franco quanto aos custos indiretos, e são eles a parte que vale a pena ler duas vezes, porque também são a parte que revela quanta margem ainda resta.

A primeira conta é o cache de instruções. Um SM armazena em cache cerca de 128 KB de instruções, e esses kernels têm de 215 a 314 KB de código de máquina cada, porque cada um contém todas as fases de uma camada. Um lançamento traz apenas os primeiros 32 a 64 KB para perto dos SMs; tudo além dessa janela é buscado linha a linha no primeiro uso. Os autores fizeram um experimento limpo sobre isso e relataram: dois blocos de 100 KB que, sozinhos, levavam 12,4 µs cada passaram a levar 19,5 µs quando emendados um após o outro, e reduzir um bloco para 20 KB devolveu ao outro sua velocidade isolada. As instruções passam pelo mesmo cache L2 que os dados, então um kernel que excede o cache de instruções acaba competindo com seus próprios carregamentos de pesos.

A segunda conta é o prefetch de pesos, e o post o trata como um problema de posicionamento com três respostas. A memória compartilhada recebe um anel de 5 × 32 KB por SM, cerca de 23 MB em toda a GPU. O L2 é onde o restante vai, e o L2 do GB300 comporta cerca de 135 MB, embora qualquer coisa no L2 precise de uma segunda viagem, mais curta, até a memória compartilhada, e o L2 é compartilhado com tudo o mais em execução. Os dois se sobrepõem: perto do fim de uma camada, o anel já contém os primeiros pesos qkv da próxima camada. Além disso, o warp 3 lê para quais especialistas a etapa anterior roteou e puxa para o L2 as fatias desta GPU dos 14 mais usados, cerca de 50 MB por GPU, apostando que etapas consecutivas roteiam de forma semelhante. Essa aposta vale cerca de 0,2 ms de uma etapa de 5,3 ms, e pode custar largura de banda quando está errada.

A terceira conta é a sincronização, e é aquela para a qual não há uma resposta limpa. Quando um warp espera por outra GPU através de NVLink, não há interrupção nem espera por endereço no hardware; o warp faz polling de um flag até que este mude. Os microbenchmarks do artigo situam uma sondagem de aquisição com âmbito de GPU em 141 a 145 ns, um flag unidirecional entre duas GPUs no mesmo nó em cerca de 1,9 µs, e uma sondagem em fila atrás de cópias em massa em 1,0 a 2,6 µs. As fences custam 0,4 µs quando inativas e 3,6 a 4,0 µs com carregamentos TMA em curso. Mesmo com trabalho perfeitamente distribuído, as GPUs podem ficar atrasadas umas em relação às outras até 4 µs, e cada paragem na camada espera pela mais lenta. O megakernel substitui as fences por valores sentinela escritos em slots de armazenamento comuns — um flag preenchido com dados, +0/−0 para o all-reduce, NaN com todos os bits a 1 para os gathers bf16 — e reinicia a cópia inativa de cada slot enquanto a outra está em uso.

Estes são os números que fazem o resultado parecer engenharia em vez de mágica: 4 µs de desalinhamento da GPU em relação a um passo de 5,9 ms é pouco, mas quatro paradas de sincronização por camada multiplicadas ao longo de 62 camadas são exatamente o tipo de sobrecarga que consome um ganho de 2× se você errar.

Em que os 5,9 ms ainda são gastos

O artigo inclui sua própria crítica, e o trecho mais útil nele é sobre a lacuna após o segundo GEMM. Os blocos terminam esse GEMM em algum ponto entre 55 e 64 µs dentro da camada, entre o 10º e o 90º percentil, mas a etapa de combinação espera pela saída de todas as GPUs, então não pode começar antes de 65 µs. Do fim do GEMM 2 ao fim da camada, um bloco fica sem trabalho por cerca de 12,5 µs em média — quase metade do seu tempo ocioso total. Essa cadeia, e não a sobrecarga de lançamento de kernel que o megakernel já eliminou, é o que os autores dizem estar examinando em seguida.

Eles também explicam por que a atenção da próxima camada não pode simplesmente começar mais cedo. O PDL a lança aos 79 µs, cerca de 4 µs antes de o primeiro bloco sair, mas seus blocos só chegam quando os anteriores deixam suas SMs aos 84 µs, então lançar mais cedo não faria nada começar mais cedo. A solução alternativa é pré-carregar suas linhas KV antigas na L2 aos 65 µs, onde elas esperam. Esse é o mesmo padrão do pré-carregamento de pesos: quando você não consegue remover uma espera, você move trabalho para dentro dela.

A análise de perfil é o outro problema estrutural, e é um bom problema: o kernel parece uma única barra longa por camada para qualquer profiler externo, então ele mesmo mede o tempo. A lane 0 de cada bloco grava um timestamp em um array global por meio do registro globaltimer, e os autores medem o custo dessa instrumentação em cerca de 2,5 µs por lançamento de camada, 2,8% de um passo, com saídas mantidas bit a bit idênticas e uma flag de compilação para removê-la na compilação. Um kernel que precisa carregar seu próprio profiler é um kernel que se tornou um sistema operacional, que é exatamente o enquadramento que o post usa: um processo é um kernel detendo todos os recursos da GPU, threads são warps sem preempção, esperar por E/S é espera ativa em uma flag, DMA é TMA rastreado por mbarrier, mensagens são escritas simples pela NVLink.

O modelo de rascunho também ficou mais rápido, e isso já é metade da vitória.

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.

A equipe aplicou o mesmo tratamento megakernel ao rascunho: cinco camadas de rascunho são executadas cada uma como um único lançamento, além de outras três, o que leva o rascunho de 175 kernels para 8 e de 1,16 ms por passo para 440 µs, uma melhoria de 2,6× nessa metade do passo. Vale a pena colocar os números de aceitação ao lado da comparação pública mais próxima que o post oferece.

• 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

• Por que 16 — dezesseis tokens custam mais para verificar porque ativam mais especialistas, mas a equipe achou mais fácil alcançar 6,6 com uma árvore mais ampla do que alcançar o mesmo comprimento com uma cadeia mais estreita

• Pilha de rascunho — 8000 linhas de CUDA e PTX para o modelo grande, com o código publicado em facebookresearch/sgl-muse-mega

Leia a comparação com atenção, porque não se trata de um benchmark entre dois modelos. É uma comparação de duas configurações de decodificação especulativa, uma publicada pela Xiaomi e outra pela equipe do SGLang, em hardware diferente e cargas de trabalho diferentes. A única afirmação que ela sustenta é que uma árvore de verificação mais ampla é um caminho viável para o mesmo comprimento de aceitação, o que é uma afirmação de design, não um ranking.

O que este número é, e o que ele não é

Eis a parte que um plano de capacidade precisa acertar, e é a parte sobre a qual o próprio post é honesto. O valor 1.109 é a mediana de 20 etapas de decodificação em uma única requisição GSM8K, com grau 8 de paralelismo de tensores, em oito GB300s, com batch size um e decodificação especulativa ativada. O post afirma claramente que o megakernel hoje suporta apenas batch size um e que estendê-lo para batches maiores é o próximo trabalho.

Cada uma dessas restrições elimina uma alegação diferente. Tamanho de lote igual a um significa que este não é um resultado de throughput no sentido que uma equipe de serving usa para a palavra: não há ganho de batching nele, nenhuma fila, nenhuma mistura de comprimentos de prompt. Um único conjunto de dados significa que o número está atrelado à distribuição de prompts e saídas do GSM8K e a um padrão de roteamento medido de 49 especialistas por camada que os autores então usaram como roteamento sintético em seu próprio profiling. Vinte passos de uma única requisição significa que a variância é limitada pelo número de passos, não pelo número de requisições.

E uma única requisição é o pior caso para sobreposição, como os próprios autores dizem: “com uma requisição, toda espera fica exposta; com mais, quase sempre há trabalho pronto.” Então o 2,1× do megakernel é uma medição do regime em que ocultar latência é mais difícil, e o post deixa explícito que lotes maiores deveriam ser mais fáceis, não mais difíceis. Isso é um argumento a favor do resultado, e ainda é um argumento, não uma medição.

O que é possível saber hoje, então, é limitado e concreto: o código existe em um repositório público, a publicação descreve o mecanismo com detalhe suficiente para ser verificada, a aritmética interna é consistente entre três razões independentes, e os autores documentaram os três sobrecustos que limitam ganhos adicionais. O que não é possível saber é se uma equipe independente obtém 1.109 em seus próprios oito GB300s, se a abordagem sobrevive a um lote real e se ela chega a uma versão lançada do SGLang — a publicação não afirma que isso já aconteceu, e o repositório é um artefato de pesquisa, não um pacote que se pode instalar. A comparação que mais resolveria a questão seria contra um baseline ajustado no mesmo hardware, já que a "cadeia padrão do SGLang" é um alvo móvel e a equipe que escreveu o megakernel também ajusta o baseline que ele supera.

Como isso se relaciona com os 169 tokens por segundo que você pode alugar hoje

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.

A lacuna é a coisa mais útil desta história para quem ainda não possui o hardware. Se um endpoint responde a 169 tokens por segundo e uma pilha de oito GPUs construída à mão responde a 1.109, a diferença não é motivo para construir a pilha. Oito GB300s são uma compra de capital medida em milhões, o kernel é código de pesquisa, e a restrição de lote de tamanho um significa que a pilha atualmente é boa exatamente para a carga de trabalho que um endpoint hospedado já lida bem: uma requisição por vez. Onde o megakernel realmente muda o cálculo é para uma equipe que já comprou as GPUs para treinamento ou para uma carga de trabalho com restrições de residência de dados, e agora pergunta qual utilização consegue extrair delas com lote baixo. Para essa equipe, 2,1× no mesmo silício com a mesma numérica é um número real com uma fatura real anexada, e a fatura já está paga.

Para todos os demais, a leitura honesta é que esta publicação é um sinal sobre para onde a infraestrutura de inferência está indo, e não um produto que você possa adotar esta semana. A direção é consistente com o que o mesmo laboratório vem publicando o ano todo: menos kernels e maiores, mais da decisão de agendamento deslocada para dentro do kernel, e a GPU tratada como uma máquina a ser programada, em vez de um conjunto de chamadas de biblioteca a serem sequenciadas.

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.

O que mudaria esta peça

Quatro coisas, em ordem decrescente do quanto elas mudariam o cenário.

• Um lote maior que um. Os autores dizem que lotes maiores são o próximo passo e argumentam que são mais fáceis, porque com mais tokens em processamento algumas das saídas da operação anterior já estão prontas enquanto o restante ainda chega. No momento em que existe um número em lote, a comparação deixa de ser "um fluxo versus um fluxo" e passa a ser uma alegação de throughput que uma equipe de serving pode usar.

• Uma reprodução independente. Cada número aqui vem do próprio harness dos autores, no hardware deles, com o próprio roteamento sintético deles para a passagem de profiling. Isso é normal para trabalho com kernel e ainda é uma única fonte. Uma segunda equipe que iguale 1.109, ou que não consiga, é o fato que resolveria a questão.

• Uma build lançada do SGLang que o inclui. O post publica o código e descreve o design; ele não afirma que o megakernel substituiu o caminho de decodificação padrão deste modelo em uma versão lançada. Até que isso aconteça, adotá-lo significa adotar um repositório de pesquisa.

• O fechamento da lacuna GEMM-2-to-combine. Os autores já apontaram 12,5 µs de tempo ocioso por bloco ali como quase metade do que resta. Se essa cadeia for corrigida, o mesmo hardware produz um tempo de etapa significativamente melhor, e o 2,1× passa a ser o piso de um número maior, em vez de a manchete.

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.