Muse Spark Open: 1.109 tokens por segundo es un cambio de kernel, no un cambio 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.

La línea base que nadie evalúa: 1.450 lanzamientos 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.

Ese es el coste al que ataca el megakernel, y los autores recurren a una analogía de los años sesenta para explicarlo: las primeras máquinas ejecutaban trabajos por lotes, un trabajo ocupaba toda la máquina y el procesador quedaba inactivo mientras giraba una cinta. La multiprogramación lo solucionó manteniendo varios trabajos residentes y cambiando a otro siempre que el actual se bloqueaba. Un megakernel es esa idea trasladada a la GPU: un kernel permanece residente en cada SM y ejecuta muchos trabajos a la vez, de modo que los warps pueden empezar el trabajo de la siguiente operación antes de que haya terminado la anterior.

El efecto medido, según la figura que acompaña a la sección de rendimiento, es un paso de decodificación que pasa de 12,9 ms a 5,9 ms. Esa figura divide el paso en sus mitades de borrador y verificación: 1,16 ms y 10,9 ms con 1.450 kernels en la línea base, 0,44 ms y 5,3 ms con 97 kernels en el megakernel. De doce coma nueve a cinco coma nueve es aproximadamente 2,2×, y 1.109 sobre 517 es aproximadamente 2,1×, así que las dos ratios principales coinciden entre sí. Hay una tercera comprobación disponible para cualquiera que esté dispuesto a hacer la división: el paso verifica 16 tokens y acepta 6,6, así que 6,6 tokens en 5,9 ms equivalen a aproximadamente 1.120 tokens por segundo. La aritmética es internamente coherente, que es lo primero que hay que establecer sobre un número como este, y no lo último.

Una cosa que las cifras no dicen es a cuánto llega la velocidad por GPU. La entrada informa el rendimiento agregado de todo el despliegue de ocho GPU. Dividido por ocho, eso da aproximadamente 139 tokens por segundo por GPU —un cálculo mío a partir de su cifra, no una cifra que publique la entrada, y no es comparable con una cifra por GPU que puedas ver citada para otro modelo en otra máquina.

Dónde se ubica la frontera del kernel, y por qué se detiene justo ahí

El megakernel no se traga todo el modelo. Cada lanzamiento abarca desde la atención de una capa del decodificador hasta la proyección qkv y la escritura KV de la capa siguiente, lo que constituye un límite deliberado y no una limitación de nerve. La razón por la que funciona es que la mayor parte de lo que lee la atención de una capa —el historial KV— se escribió en pasos anteriores y no depende de nada en curso, así que un nuevo lanzamiento puede empezar a cargarlo en el instante en que esté disponible, antes de que exista la salida de la capa anterior. La primera y la última capa tienen operaciones ligeramente distintas, y la publicación también lo señala.

Dentro de un mismo lanzamiento, la forma es fija: un bloque de hilos de doce warps en cada una de las 144 SMs, lanzados en clústeres de cuatro, usando memoria compartida distribuida donde hay disponible una reducción de menor latencia. Los doce warps no son intercambiables. El warp 0 carga KV y pesos y dedica su margen a precargar los pesos de la siguiente operación. Los warps 1 y 2 ejecutan activaciones y tensor cores. Los warps 3 a 11 se reparten la atención, los GEMMs densos, el GEMM experto y las recepciones de all-reduce y all-gather. En las capas de ventana deslizante, los warps 4 a 11 cargan su propio KV en lugar de esperar al 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.

Ese detalle importa más de lo que parece. Los equipos de inferencia suelen aceptar que un kernel más rápido produzca números ligeramente distintos y reajustan en silencio sus umbrales de evaluación. Un kernel que es bit a bit igual a la paralelización del lado del entrenamiento elimina ese argumento por completo, que es el tipo de propiedad que hace que una optimización se adopte en lugar de simplemente admirarse.

Tres cuentas tiene que pagar un grano de maíz así de grande.

El megakernel no es gratis. La publicación es inusualmente franca sobre las sobrecargas, y son la parte que vale la pena leer dos veces, porque también son la parte que te dice cuánto margen queda.

El primer costo es la caché de instrucciones. Una SM almacena en caché aproximadamente 128 KB de instrucciones, y estos kernels son de 215 a 314 KB de código máquina cada uno, porque cada uno contiene todas las fases de una capa. Un lanzamiento solo acerca a las SM los primeros 32 a 64 KB; todo lo que queda más allá de esa ventana se obtiene línea por línea en el primer uso. Los autores realizaron un experimento limpio sobre esto y lo reportaron: dos bloques de 100 KB que por separado tardaban 12,4 µs cada uno tardaron 19,5 µs cuando se unieron uno tras otro, y reducir uno de los bloques a 20 KB devolvió al otro a su velocidad individual. Las instrucciones viajan por la misma caché L2 que los datos, así que un kernel que supera el tamaño de la caché de instrucciones termina compitiendo con sus propias cargas de pesos.

La segunda factura es la precarga de pesos, y el artículo la aborda como un problema de colocación con tres respuestas. La memoria compartida recibe un anillo de 5 × 32 KB por SM, unos 23 MB en toda la GPU. La L2 es donde va el resto, y la L2 del GB300 alberga unos 135 MB, aunque cualquier cosa en L2 necesita un segundo viaje, más corto, a la memoria compartida, y la L2 se comparte con todo lo demás en vuelo. Las dos se apilan: cerca del final de una capa, el anillo ya contiene los primeros pesos qkv de la siguiente capa. Además, el warp 3 lee a qué expertos enrutó el paso anterior y lleva a L2 las porciones de esta GPU de los 14 más usados, unos 50 MB por GPU, apostando a que los pasos consecutivos enrutan de forma similar. Esa apuesta vale unos 0,2 ms de un paso de 5,3 ms, y puede costar ancho de banda cuando falla.

El tercer costo es la sincronización, y es el que no tiene una respuesta clara. Cuando un warp espera a otra GPU a través de NVLink, no hay interrupción ni espera en dirección en el hardware; el warp sondea una bandera hasta que cambia. Los microbenchmarks del artículo sitúan un sondeo de adquisición con alcance de GPU en 141 a 145 ns, una bandera unidireccional entre dos GPU en el mismo nodo en unos 1.9 µs, y un sondeo en cola detrás de copias masivas en 1.0 a 2.6 µs. Las fences cuestan 0.4 µs cuando están inactivas y 3.6 a 4.0 µs con cargas TMA en vuelo. Incluso con un trabajo perfectamente distribuido, las GPU pueden retrasarse entre sí hasta 4 µs, y cada parada en la capa espera a la más lenta. El megakernel reemplaza las fences por valores centinela escritos en ranuras de almacenamiento ordinarias —una bandera llena de datos, +0/−0 para el all-reduce, NaN de todos unos para las recolecciones bf16— y restablece la copia inactiva de cada ranura mientras la otra está en uso.

Estos son los números que hacen que el resultado parezca ingeniería en vez de magia: 4 µs de desfase de GPU frente a un paso de 5,9 ms es poco, pero cuatro paradas de sincronización por capa multiplicadas por 62 capas es justo el tipo de sobrecarga que se come una ventaja de 2× si te equivocas.

En qué se siguen gastando los 5.9 ms

La publicación incluye su propia crítica, y la línea más útil de ella trata sobre el hueco después del segundo GEMM. Los bloques terminan ese GEMM en algún punto entre 55 y 64 µs dentro de la capa, entre los percentiles 10 y 90, pero la combinación espera la salida de todas las GPU, así que no puede empezar hasta los 65 µs. Desde el final de GEMM 2 hasta el final de la capa, un bloque permanece sin trabajo durante unos 12,5 µs en promedio —casi la mitad de su tiempo total de inactividad—. Esa cadena, y no la sobrecarga de lanzamiento del kernel que el megakernel ya eliminó, es lo que los autores dicen que están examinando a continuación.

También explican por qué la atención de la siguiente capa no puede simplemente empezar antes. PDL la lanza a los 79 µs, unos 4 µs antes de que salga el primer bloque, pero sus bloques solo llegan una vez que los anteriores abandonan sus SMs a los 84 µs, así que lanzarla antes no haría que nada empezara antes. La solución alternativa consiste en precargar sus antiguas filas KV en L2 a los 65 µs, donde esperan. Es el mismo patrón que la precarga de pesos: cuando no puedes eliminar una espera, trasladas trabajo a ella.

El profiling es el otro problema estructural, y es un buen problema: para cualquier profiler externo, el kernel parece una sola barra larga por capa, así que se cronometra a sí mismo. La lane 0 de cada bloque escribe una marca de tiempo en un array global mediante el registro globaltimer, y los autores miden el coste de esa instrumentación en unos 2,5 µs por lanzamiento de capa, un 2,8 % de un paso, con las salidas intactas bit a bit y un flag de compilación para eliminarla al compilar. Un kernel que tiene que llevar su propio profiler es un kernel que se ha convertido en un sistema operativo, que es exactamente el planteamiento que usa el post: un proceso es un kernel que retiene todos los recursos de la GPU, los hilos son warps sin preempción, esperar E/S es espera activa sobre un flag, DMA es TMA rastreado por mbarrier, los mensajes son simples stores a través de NVLink.

El modelo de borrador también se volvió más rápido, y eso es la mitad de la victoria.

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.

El equipo aplicó el mismo tratamiento de megakernel al borrador: cinco capas de borrador se ejecutan cada una como un solo lanzamiento, más otras tres, lo que lleva el borrador de 175 kernels a 8 y de 1,16 ms por paso a 440 µs, una mejora de 2,6× en esa mitad del paso. Vale la pena poner las cifras de aceptación junto a la comparación pública más cercana que ofrece el artículo.

• 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 qué 16: verificar dieciséis tokens cuesta más porque activan más expertos, pero al equipo le resultó más fácil alcanzar 6,6 con un árbol más ancho que alcanzar la misma longitud con una cadena más estrecha

• Pila de borrador — 8000 líneas de CUDA y PTX para el modelo grande, con el código publicado en facebookresearch/sgl-muse-mega

Lee la comparación con atención, porque no es una comparación entre dos modelos. Es una comparación entre dos configuraciones de decodificación especulativa, una publicada por Xiaomi y otra por el equipo de SGLang, sobre hardware diferente y cargas de trabajo diferentes. La única afirmación que respalda es que un árbol de verificación más amplio es una vía viable para alcanzar la misma longitud de aceptación, lo cual es una afirmación de diseño, no una clasificación.

Qué es este número y qué no es

Esta es la parte que un plan de capacidad tiene que acertar, y es la parte sobre la que el propio artículo es honesto. La cifra de {{1}}1,109{{/1}} es la mediana de {{2}}20 pasos de decodificación{{/2}} en una sola solicitud {{3}}GSM8K{{/3}}, con {{4}}grado de paralelismo tensorial 8{{/4}}, en ocho {{5}}GB300s{{/5}}, con {{6}}tamaño de lote uno{{/6}} y {{7}}decodificación especulativa activada{{/7}}. El artículo afirma claramente que el megakernel hoy solo admite {{8}}tamaño de lote uno{{/8}}, y que extenderlo a lotes más grandes es el siguiente trabajo.

Cada una de esas restricciones elimina una afirmación distinta. Un tamaño de lote de uno significa que este no es un resultado de throughput en el sentido en que un equipo de serving usa la palabra: no hay ganancia de batching en ello, ni cola, ni mezcla de longitudes de prompt. Un solo dataset significa que la cifra está atada a la distribución de prompts y salidas de GSM8K, y a un patrón de enrutamiento medido de 49 expertos por capa que los autores luego usaron como enrutamiento sintético en su propio profiling. Veinte pasos de una sola solicitud significa que la varianza está acotada por el número de pasos, no por el número de solicitudes.

Y una sola solicitud es el peor caso para el solapamiento, como dicen los propios autores: «con una solicitud, cada espera queda expuesta; con más, casi siempre hay trabajo disponible». Así que el 2,1× del megakernel es una medición del régimen en el que ocultar la latencia resulta más difícil, y la publicación deja claro que los lotes más grandes deberían ser más fáciles, no más difíciles. Eso es un argumento a favor del resultado, y sigue siendo un argumento en lugar de una medición.

Lo que hoy puede saberse, por tanto, es limitado y real: el código existe en un repositorio público, la publicación describe el mecanismo con suficiente detalle para poder verificarlo, la aritmética interna es coherente en tres ratios independientes, y los autores han documentado las tres sobrecargas que acotan las ganancias adicionales. Lo que no puede saberse es si un equipo independiente obtiene 1.109 en sus propias ocho GB300, si el enfoque sobrevive a un lote real y si llega a una compilación publicada de SGLang; la publicación no afirma que lo haya logrado, y el repositorio es un artefacto de investigación más que un paquete que se pueda instalar. La comparación que más esclarecería las cosas es la que se haga contra una línea base optimizada en el mismo hardware, ya que «la cadena predeterminada de SGLang» es un blanco móvil y el equipo que escribió el megakernel también optimiza la línea base a la que supera.

Cómo se relaciona con los 169 tokens por segundo que puedes alquilar hoy

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.

La diferencia es lo más útil de esta historia para cualquiera que aún no tenga el hardware. Si un endpoint responde a 169 tokens por segundo y una pila de ocho GPU montada a mano responde a 1.109, la diferencia no es una razón para construir la pila. Ocho GB300 son una compra de capital de millones, el kernel es código de investigación y la restricción de tamaño de lote uno significa que la pila actualmente es buena exactamente para la carga de trabajo que un endpoint alojado ya maneja bien: una solicitud a la vez. Donde el megakernel realmente cambia el cálculo es para un equipo que ya compró las GPU para entrenamiento o para una carga de trabajo con restricciones de residencia de datos, y ahora se pregunta qué utilización puede obtener de ellas con lote bajo. Para ese equipo, 2,1× en el mismo silicio con los mismos cálculos numéricos es un número real con una factura real asociada, y la factura ya está pagada.

Para todos los demás, la lectura honesta es que esta publicación es una señal de hacia dónde va la infraestructura de inferencia, no un producto que puedas adoptar esta semana. La dirección es coherente con lo que el mismo laboratorio ha estado publicando todo el año: menos kernels y de mayor tamaño, una mayor parte de la decisión de planificación trasladada dentro del kernel, y la GPU tratada como una máquina que se programa en lugar de un conjunto de llamadas de biblioteca que se secuencian.

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.

¿Qué cambiaría esta pieza?

Cuatro cosas, en orden descendente de cuánto cambiarían el panorama.

• Un lote mayor que uno. Los autores señalan que lo siguiente son lotes más grandes y sostienen que son más fáciles, porque con más tokens en vuelo algunas de las salidas de la operación anterior ya están listas mientras el resto sigue llegando. En el momento en que existe una cifra por lotes, la comparación deja de ser "un flujo frente a un flujo" y pasa a ser una afirmación sobre el throughput que un equipo de serving puede aprovechar.

• Una reproducción independiente. Cada cifra que aparece aquí procede del propio banco de pruebas de los autores, en su propio hardware, con su propio enrutamiento sintético para la pasada de perfilado. Eso es normal en el trabajo con kernels y sigue siendo una única fuente. Que un segundo equipo iguale 1.109, o no lo consiga, es el dato que lo zanjaría.

• Una versión publicada de SGLang que lo incluya. La publicación difunde el código y describe el diseño; no afirma que el megakernel haya sustituido la ruta de decodificación predeterminada para este modelo en una versión distribuida. Hasta que lo haga, adoptarlo implica adoptar un repositorio de investigación.

• El cierre de la brecha GEMM-2-to-combine. Los autores ya han señalado 12,5 µs de tiempo de inactividad por bloque ahí como casi la mitad de lo que queda. Si se corrige esa cadena, el mismo hardware produce un tiempo de paso significativamente mejor, y el 2,1× se convierte en la base de un número mayor en lugar de ser el titular.

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.