Muse Spark Open : 1 109 jetons par seconde, c’est un changement de noyau, pas un changement de modèle

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 référence que personne n’évalue : 1 450 lancements par 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.

C’est le coût que le megakernel attaque, et les auteurs recourent à une analogie des années 1960 pour l’expliquer : les premières machines exécutaient des tâches par lots, une tâche monopolisant toute la machine, et le processeur restait inactif pendant qu’une bande tournait. La multiprogrammation a résolu cela en gardant plusieurs tâches résidentes et en passant à une autre chaque fois que celle en cours se bloquait. Un megakernel, c’est cette idée transposée au GPU — un noyau reste résident sur chaque SM et exécute de nombreuses tâches à la fois, de sorte que les warps peuvent commencer le travail de l’opération suivante avant que l’opération précédente soit terminée.

L’effet mesuré, d’après la figure qui accompagne la section sur les performances, est une étape de décodage qui passe de 12,9 ms à 5,9 ms. Cette figure décompose l’étape en ses moitiés de brouillon et de vérification : 1,16 ms et 10,9 ms avec 1 450 noyaux dans la version de référence, 0,44 ms et 5,3 ms avec 97 noyaux dans le megakernel. De douze virgule neuf à cinq virgule neuf, cela fait environ 2,2×, et 1 109 sur 517 environ 2,1×, donc les deux ratios phares concordent. Il existe une troisième vérification à la portée de quiconque accepte de faire la division : l’étape vérifie 16 jetons et en accepte 6,6, donc 6,6 jetons en 5,9 ms équivalent à environ 1 120 jetons par seconde. L’arithmétique est cohérente en interne, ce qui est la première chose à établir au sujet d’un nombre comme celui-ci, et non la dernière.

Une chose que les chiffres ne disent pas, c’est où se situe la vitesse par GPU. L’article rapporte un débit agrégé pour l’ensemble du déploiement à huit GPU. Divisé par huit, cela fait environ 139 tokens par seconde par GPU — c’est mon calcul à partir de leur chiffre, pas un nombre que l’article publie, et ce n’est pas comparable à un nombre par GPU que vous pourriez voir cité pour un autre modèle sur une autre machine.

Où se situe la frontière du noyau, et pourquoi elle s'arrête là

Le mégakernel n'englobe pas tout le modèle. Chaque lancement couvre l'attention d'une couche de décodeur jusqu'à la projection qkv et à l'écriture KV de la couche suivante, ce qui constitue une frontière délibérée plutôt qu'une limitation de nerve. La raison pour laquelle cela fonctionne, c'est que l'essentiel de ce que lit l'attention d'une couche — l'historique KV — a été écrit lors d'étapes antérieures et ne dépend de rien en cours d'exécution, de sorte qu'un nouveau lancement peut commencer à le charger dès qu'il est disponible, avant même que la sortie de la couche précédente n'existe. Les première et dernière couches ont des opérations légèrement différentes, et le post le mentionne aussi.

Au sein d’un seul lancement, la configuration est figée : un bloc de threads de douze warps sur chacun des 144 SM, lancé en clusters de quatre, en utilisant la mémoire partagée distribuée lorsqu’une réduction à plus faible latence est disponible. Les douze warps ne sont pas interchangeables. Le warp 0 charge les KV et les poids et consacre son temps mort à précharger les poids de l’opération suivante. Les warps 1 et 2 exécutent les activations et les cœurs tensoriels. Les warps 3 à 11 se répartissent l’attention, les GEMM denses, le GEMM des experts et les réceptions all-reduce et all-gather. Dans les couches à fenêtre glissante, les warps 4 à 11 chargent leurs propres KV au lieu d’attendre le 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.

Ce détail compte plus qu’il n’y paraît. Les équipes d’inférence acceptent généralement qu’un noyau plus rapide produise des nombres légèrement différents et réajustent discrètement leurs seuils d’évaluation. Un noyau bit à bit identique à la parallélisation côté entraînement élimine complètement cet argument, ce qui est le genre de propriété qui fait qu’une optimisation est adoptée plutôt que simplement admirée.

Trois factures qu’un noyau aussi gros doit payer

Le megakernel n’est pas gratuit. L’article est d’une franchise inhabituelle au sujet des surcoûts, et c’est la partie qui mérite d’être lue deux fois, car c’est aussi elle qui vous dit combien de marge il reste.

La première facture est le cache d’instructions. Un SM met en cache environ 128 Ko d’instructions, et ces kernels font chacun de 215 à 314 Ko de code machine, car chacun contient toutes les phases d’une couche. Un lancement n’amène que les premiers 32 à 64 Ko à proximité des SM ; tout ce qui dépasse cette fenêtre est récupéré ligne par ligne lors de la première utilisation. Les auteurs ont mené une expérience propre sur ce point et l’ont rapportée : deux blocs de 100 Ko qui prenaient chacun 12,4 µs seuls prenaient 19,5 µs lorsqu’ils étaient enchaînés bout à bout, et réduire un bloc à 20 Ko a ramené l’autre à sa vitesse en solo. Les instructions passent par le même cache L2 que les données, donc un kernel qui dépasse le cache d’instructions finit par entrer en concurrence avec ses propres chargements de poids.

Le deuxième poste de coût est le préchargement des poids, et l’article le traite comme un problème de placement à trois solutions. La mémoire partagée reçoit un anneau de 5 × 32 Ko par SM, soit environ 23 Mo sur l’ensemble du GPU. Le L2 accueille le reste, et le L2 du GB300 contient environ 135 Mo, mais tout ce qui se trouve dans le L2 doit effectuer un second trajet, plus court, vers la mémoire partagée, et le L2 est partagé avec tout le reste en cours d’exécution. Les deux se superposent : vers la fin d’une couche, l’anneau contient déjà les premiers poids qkv de la couche suivante. En plus de cela, le warp 3 lit quels experts l’étape précédente a routés et amène dans le L2 les tranches de ce GPU des 14 experts les plus utilisés, environ 50 Mo par GPU, en pariant que des étapes consécutives routent de façon similaire. Ce pari vaut environ 0,2 ms sur une étape de 5,3 ms, et peut coûter de la bande passante lorsqu’il est faux.

Le troisième coût est la synchronisation, et c’est celui pour lequel il n’y a pas de réponse nette. Quand un warp attend un autre GPU via NVLink, il n’y a ni interruption ni attente sur adresse dans le matériel ; le warp sonde un drapeau jusqu’à ce qu’il change. Les microbenchmarks de l’article situent un sondage d’acquisition à portée GPU à 141 à 145 ns, un drapeau unidirectionnel entre deux GPU d’un même nœud à environ 1,9 µs, et un sondage mis en file derrière des copies en masse à 1,0 à 2,6 µs. Les barrières coûtent 0,4 µs au repos et 3,6 à 4,0 µs avec des chargements TMA en cours. Même avec un travail parfaitement réparti, les GPU peuvent avoir jusqu’à 4 µs de retard les uns par rapport aux autres, et chaque point d’arrêt de la couche attend le plus lent. Le mégakernel remplace les barrières par des valeurs sentinelles écrites dans des emplacements de stockage ordinaires — un drapeau rempli de données, +0/−0 pour l’all-reduce, NaN tout-à-un pour les gathers bf16 — et réinitialise la copie inactive de chaque emplacement pendant que l’autre est utilisée.

Voici les chiffres qui donnent au résultat une allure d’ingénierie plutôt que de magie : 4 µs de décalage GPU contre un pas de 5,9 ms, c’est peu, mais quatre arrêts de synchronisation par couche, répétés sur 62 couches, représentent exactement le type de surcoût qui engloutit un gain de 2× si l’on se trompe.

À quoi les 5,9 ms sont encore consacrées

Le billet contient sa propre critique, et la phrase la plus utile y concerne l’écart après le deuxième GEMM. Les blocs terminent ce GEMM entre 55 et 64 µs après le début de la couche, du 10e au 90e percentile, mais le combine attend la sortie de chaque GPU, il ne peut donc pas démarrer avant 65 µs. Entre la fin du GEMM 2 et la fin de la couche, un bloc reste sans travail pendant environ 12,5 µs en moyenne — presque la moitié de son temps d’inactivité total. C’est cette chaîne, et non la surcharge de lancement du noyau que le mégakernel a déjà éliminée, que les auteurs disent examiner ensuite.

Ils expliquent aussi pourquoi l’attention de la couche suivante ne peut pas simplement commencer plus tôt. PDL la lance à 79 µs, environ 4 µs avant la sortie du premier bloc, mais ses blocs n’arrivent qu’une fois que les précédents quittent leurs SMs à 84 µs ; lancer plus tôt ne ferait donc rien démarrer plus tôt. La solution de contournement consiste à précharger ses anciennes lignes KV dans L2 à 65 µs, où elles attendent. C’est le même schéma que le préchargement des poids : lorsqu’on ne peut pas supprimer une attente, on y déplace du travail.

Le profilage est l’autre problème structurel, et c’est un bon problème : pour tout profileur externe, le kernel ressemble à une seule longue barre par couche, il se chronomètre donc lui-même. La lane 0 de chaque bloc écrit un horodatage dans un tableau global via le registre globaltimer, et les auteurs mesurent le coût de cette instrumentation à environ 2,5 µs par lancement de couche, soit 2,8 % d’un pas, les sorties restant identiques bit pour bit et une option de compilation permettant de l’exclure. Un kernel qui doit embarquer son propre profileur est un kernel devenu un système d’exploitation, ce qui est exactement le cadrage employé par le billet : un processus est un kernel détenant toutes les ressources GPU, les threads sont des warps sans préemption, attendre des E/S revient à faire de l’attente active sur un indicateur, le DMA est le TMA suivi par mbarrier, les messages sont de simples écritures via NVLink.

Le modèle draft est aussi devenu plus rapide, et c’est la moitié de la victoire.

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.

Le même traitement « megakernel » a été appliqué par l’équipe au draft : cinq couches du draft s’exécutent chacune en un seul lancement, plus trois autres, ce qui fait passer le draft de 175 kernels à 8 et de 1,16 ms par étape à 440 µs, soit une amélioration de 2,6× sur cette moitié de l’étape. Les chiffres d’acceptation méritent d’être mis en regard de la comparaison publique la plus proche proposée dans l’article.

• 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

• Pourquoi 16 — seize jetons coûtent plus cher à vérifier parce qu’ils activent davantage d’experts, mais l’équipe a trouvé plus facile d’atteindre 6,6 avec un arbre plus large que d’atteindre la même longueur avec une chaîne plus étroite

• Pile de brouillon — 8 000 lignes de CUDA et PTX pour le grand modèle, avec le code publié sur facebookresearch/sgl-muse-mega

Lisez la comparaison attentivement, car il ne s'agit pas d'un benchmark entre deux modèles. Il s'agit d'une comparaison de deux configurations de décodage spéculatif, l'une publiée par Xiaomi et l'autre par l'équipe SGLang, sur des matériels et des charges de travail différents. La seule affirmation qu'elle étaye est qu'un arbre de vérification plus large est une voie viable pour atteindre la même longueur d'acceptation, ce qui est une affirmation de conception, non un classement.

Ce qu’est ce numéro, et ce qu’il n’est pas

Voici la partie qu'un plan de capacité doit absolument bien traiter, et c'est celle sur laquelle l'article lui-même est honnête. Le chiffre de 1 109 correspond à la médiane de 20 étapes de décodage sur une seule requête GSM8K, avec un parallélisme tensoriel de degré 8, sur huit GB300, avec une taille de lot égale à un et le décodage spéculatif activé. L'article indique clairement que le megakernel ne prend aujourd'hui en charge qu'une taille de lot égale à un, et que son extension à des lots plus grands constitue le prochain travail à réaliser.

Chacune de ces contraintes retire une affirmation différente. Une taille de lot de un signifie qu’il ne s’agit pas d’un résultat de débit au sens où une équipe de serving emploie ce terme : il n’y a là aucun gain de batching, aucune file d’attente, aucun mélange de longueurs de prompt. Un seul jeu de données signifie que le chiffre est lié à la distribution des prompts et des sorties de GSM8K, ainsi qu’à un motif de routage mesuré de 49 experts par couche que les auteurs ont ensuite utilisé comme routage synthétique dans leur propre profilage. Vingt étapes pour une seule requête signifient que la variance est bornée par le nombre d’étapes, et non par le nombre de requêtes.

Et une seule requête est le pire cas pour le chevauchement, ce que les auteurs disent eux-mêmes : « à une requête, chaque attente est exposée ; avec davantage, il y a presque toujours du travail prêt. » Ainsi, le 2,1× du megakernel est une mesure du régime où le masquage de latence est le plus difficile, et l’article est explicite : des lots plus grands devraient être plus faciles plutôt que plus difficiles. C’est un argument en faveur du résultat, et cela reste un argument plutôt qu’une mesure.

Ce que l’on peut savoir aujourd’hui, donc, est limité et bien réel : le code existe dans un dépôt public, l’article décrit le mécanisme avec suffisamment de détails pour être vérifié, l’arithmétique interne concorde sur trois ratios indépendants, et les auteurs ont documenté les trois surcoûts qui limitent les gains supplémentaires. Ce que l’on ne peut pas savoir, c’est si une équipe indépendante obtient 1 109 sur ses propres huit GB300s, si l’approche survit à un vrai lot, et si elle est intégrée à une version publiée de SGLang — l’article ne prétend pas que c’est le cas, et le dépôt est un artefact de recherche plutôt qu’un paquet installable. La comparaison qui trancherait le plus est celle avec une référence optimisée sur le même matériel, car « la chaîne par défaut de SGLang » est une cible mouvante et l’équipe qui a écrit le megakernel optimise aussi la référence qu’il bat.

Comment cela se rapporte aux 169 tokens par seconde que vous pouvez louer aujourd’hui

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.

L’écart est l’élément le plus utile de cette histoire pour quiconque ne possède pas déjà le matériel. Si un point de terminaison répond à 169 tokens par seconde et qu’une pile de huit GPU assemblée à la main répond à 1 109, la différence n’est pas une raison de construire cette pile. Huit GB300 représentent un investissement en capital de l’ordre de plusieurs millions, le noyau est du code de recherche, et la contrainte de taille de lot de un signifie que la pile n’est actuellement bonne que pour exactement la charge de travail qu’un point de terminaison hébergé gère déjà bien : une requête à la fois. Là où le mégakernel change véritablement la donne, c’est pour une équipe qui a déjà acheté les GPU pour l’entraînement ou pour une charge de travail soumise à des contraintes de résidence des données, et qui se demande maintenant quel taux d’utilisation elle peut en tirer à faible taille de lot. Pour cette équipe, un facteur 2,1× sur le même silicium avec les mêmes calculs numériques est un chiffre réel auquel correspond une facture réelle, et la facture est déjà payée.

Pour tous les autres, la lecture honnête est que cette publication est un signal sur la direction que prend l’infrastructure d’inférence, et non un produit que vous pouvez adopter cette semaine. La direction est cohérente avec ce que le même laboratoire publie depuis le début de l’année : des noyaux moins nombreux et plus gros, une part accrue des décisions d’ordonnancement déplacée à l’intérieur du noyau, et le GPU considéré comme une machine à programmer plutôt que comme un ensemble d’appels de bibliothèque à enchaîner.

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'est-ce qui changerait cette pièce ?

Quatre choses, par ordre décroissant de l’impact qu’elles auraient sur la situation.

• Un lot supérieur à un. Les auteurs annoncent que des lots plus grands arrivent ensuite et soutiennent qu'ils sont plus faciles, car avec davantage de jetons en cours de traitement, certaines sorties de l'opération précédente sont prêtes pendant que les autres arrivent encore. Dès qu'un chiffre issu d'un traitement par lots existe, la comparaison cesse d'être « un flux contre un flux » et devient une affirmation de débit qu'une équipe chargée du service d'inférence peut utiliser.

• Une reproduction indépendante. Chaque chiffre ici provient du banc d’essai des auteurs eux-mêmes, sur leur propre matériel, avec leur propre routage synthétique pour la passe de profilage. C’est normal pour du travail sur noyau, et cela reste une source unique. Qu’une seconde équipe retrouve 1,109, ou n’y parvienne pas, voilà le fait qui trancherait.

• Une version publiée de SGLang qui l’intègre. L’article publie du code et décrit la conception ; il ne dit pas que le mégakernel a remplacé le chemin de décodage par défaut pour ce modèle dans une version livrée. Tant que ce n’est pas le cas, l’adopter revient à adopter un dépôt de recherche.

• La réduction de l’écart entre GEMM-2 et combine. Les auteurs ont déjà identifié 12,5 µs de temps d’inactivité par bloc à cet endroit comme représentant près de la moitié de ce qui reste. Si cette chaîne est corrigée, le même matériel produit un temps par étape nettement meilleur, et le 2,1× devient le plancher d’un chiffre plus élevé plutôt que le résultat phare.

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.