Muse Spark Open：毎秒1,109トークンはカーネルの変更であり、モデルの変更ではない

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.

誰もベンチマークしないベースライン：トークンあたり1,450回の起動

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.

それがメガカーネルの狙うコストであり、著者らはそれを説明するために1960年代のアナロジーを持ち出す。初期のマシンはバッチジョブを実行し、1つのジョブがマシン全体を占有し、テープが回転している間プロセッサはアイドル状態だった。マルチプログラミングは、複数のジョブを常駐させ、現在のジョブがブロックするたびに別のジョブへ切り替えることでこれを解決した。メガカーネルはそのアイデアをGPUへ移したものだ — 1つのカーネルがすべてのSM上に常駐して多数のジョブを同時に実行するため、ワープは前の演算が完了する前に次の演算の作業を開始できる。

パフォーマンスのセクションに付随する図から測定される効果は、デコードステップが12.9 msから5.9 msになることです。その図は、このステップをドラフト側と検証側の2つに分けています。ベースラインでは1,450カーネルで1.16 msと10.9 ms、メガカーネルでは97カーネルで0.44 msと5.3 msです。12.9から5.9は約2.2倍、1,109÷517は約2.1倍なので、この2つの主要な比率は互いに一致しています。割り算を進んで行う人なら誰でも使える3つ目の確認があります。このステップは16トークンを検証して6.6を受け入れるので、5.9 msで6.6トークンはおよそ毎秒1,120トークンになります。算術は内部的に整合しており、それはこの種の数値について最初に確立すべきことであって、最後ではありません。

数字が語っていないことの一つは、速度がGPUあたりでどの程度になるかだ。投稿が報告しているのは、8基のGPU配備全体の合計スループットだ。8で割ると、GPUあたり毎秒約139トークンになる。これは私が彼らの数値をもとに計算したもので、投稿が公表している数値ではなく、別のマシンで別のモデルについて引用されるGPUあたりの数値とは比較できない。

カーネル境界がどこに位置するのか、そしてなぜそこで止まるのか

メガカーネルはモデル全体を飲み込むわけではない。各起動は、あるデコーダ層のアテンションから次層のqkv射影とKV書き込みまでをカバーしており、これはnerveの制約ではなく意図的な境界である。うまく機能する理由は、ある層のアテンションが読み取るものの大半——KV履歴——が以前のステップで書き込まれ、実行中の何ものにも依存しないため、前の層の出力が存在する前に、利用可能になった瞬間から新しい起動がその読み込みを開始できるからだ。最初と最後の層は演算が少し異なり、投稿ではその点にも触れている。

1回の起動内では形状は固定される。144個のSMそれぞれで12ワープの1スレッドブロックが、4つ組のクラスタで起動され、低レイテンシのリダクションが利用できる分散共有メモリを使用する。12個のワープは互換ではない。ワープ0はKVと重みをロードし、その余裕を使って次の演算の重みをプリフェッチする。ワープ1と2はアクティベーションとテンソルコアを実行する。ワープ3から11はアテンション、密なGEMM、エキスパートGEMM、およびall-reduceとall-gatherの受信を分担する。スライディングウィンドウ層では、ワープ4から11はワープ0を待たずに自身のKVをロードする。

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.

その詳細は見た目以上に重要だ。推論チームは通常、より高速なカーネルがわずかに異なる数値を生み出すことを受け入れ、評価しきい値を静かに再調整する。訓練側の並列化とビット単位で等しいカーネルは、その議論を完全に排除する。これは、最適化が単に称賛されるのではなく採用される類の特性だ。

これほど大きなカーネルが支払わなければならない3つの請求書

メガカーネルは無料ではない。その投稿はオーバーヘッドについて異例なほど率直で、そこは二度読む価値がある部分だ。なぜなら、そこはどれだけ余地が残っているかを教えてくれる部分でもあるからだ。

最初の代償は命令キャッシュだ。SMは約128KBの命令をキャッシュするが、これらのカーネルはそれぞれ215〜314KBのマシンコードであり、各カーネルがレイヤの全フェーズを保持しているためだ。起動時にSMの近くに置かれるのは最初の32〜64KBのみで、そのウィンドウを超えた部分は初回使用時にライン単位でフェッチされる。著者らはこれについてクリーンな実験を行い、報告している。単独ではそれぞれ12.4µsかかった2つの100KBブロックを連結すると19.5µsかかり、一方のブロックを20KBに縮小すると、もう一方は単独時の速度に戻った。命令はデータと同じL2キャッシュを通るため、命令キャッシュに収まらなくなったカーネルは、自身の重みのロードと競合することになる。

第2の案は重みプリフェッチであり、投稿ではこれを3つの答えを持つ配置問題として扱っている。共有メモリにはSMあたり5 × 32 KBのリングが確保され、GPU全体では約23 MBになる。残りはL2に置かれ、GB300のL2は約135 MBを保持するが、L2にあるものは共有メモリへより短い2度目の移動が必要で、L2は実行中の他のあらゆるものと共有される。この2つは積み重なる。層の終わり近くでは、リングにはすでに次層の最初のqkv重みが入っている。さらに、warp 3は前のステップがどのエキスパートにルーティングしたかを読み、最も使われる14個のうちこのGPUのスライスをL2に引き込む。GPUあたり約50 MBで、連続するステップは同じようにルーティングされると賭けている。その賭けは5.3 msのステップのうち約0.2 msの価値があり、外れたときは帯域を消費し得る。

3つ目の代償は同期であり、これには明快な答えがない。ワープがNVLink経由で別のGPUを待つとき、ハードウェアには割り込みもアドレス待機もない。ワープはフラグが変わるまでポーリングする。記事のマイクロベンチマークでは、GPUスコープのacquireポーリングは141〜145 ns、同一ノード上の2つのGPU間の片方向フラグは約1.9 µs、バルクコピーの後ろに並んだポーリングは1.0〜2.6 µsだった。フェンスはアイドル時で0.4 µs、TMAロードが進行中だと3.6〜4.0 µsかかる。作業が完全に分散されていても、GPU間で最大4 µsの遅れが生じうるし、レイヤー内のすべての停止は最も遅いものを待つ。メガカーネルはフェンスを、通常のストアスロットに書き込まれるセンチネル値に置き換える。すなわち、データで埋めたフラグ、all-reduce用の+0/−0、bf16ギャザー用の全ビット1のNaNであり、一方を使用中にもう一方のアイドルコピーをリセットする。

結果を魔法ではなくエンジニアリングだと感じさせるのは、こうした数字だ。5.9 ms のステップに対して 4 µs の GPU スキューは小さいが、レイヤーごとに4回の同期停止が62レイヤーにわたって積み重なると、間違えた場合に2×の高速化を食い尽くす、まさにその種のオーバーヘッドになる。

5.9 ms はまだ何に費やされているのか

その投稿にはそれ自体への批判も含まれており、中でも最も有用な一節は2回目のGEMM後のギャップについて述べている。ブロックはそのGEMMをレイヤー開始から55〜64µsの間（10〜90パーセンタイル）に終えるが、結合処理はすべてのGPUの出力を待つため、65µsになるまで開始できない。GEMM 2の終了からレイヤーの終了までの間、ブロックは平均で約12.5µs、何の作業もなく待機する——これはブロックの総アイドル時間のほぼ半分に当たる。著者らが次に注目しているのは、メガカーネルがすでに排除したカーネル起動オーバーヘッドではなく、この連鎖だという。

彼らは、次の層のアテンションが単純に早く開始できない理由も説明している。PDL はそれを 79 µs に起動する。最初のブロックが抜ける約 4 µs 前だが、そのブロックは前のブロックが 84 µs に SM を離れて初めて着地するため、より早く起動しても何も早く開始しない。回避策は、その古い KV 行を 65 µs に L2 へプリフェッチし、そこで待たせることだ。これは重みのプリフェッチと同じパターンである。待ちをなくせないなら、その待ち時間に仕事を移すのだ。

プロファイリングはもう一つの構造上の問題であり、これは良い問題だ。カーネルは外部のプロファイラから見ると各レイヤーごとに一本の長いバーのように見えるため、自身で時間を計測する。各ブロックのレーン0がglobaltimerレジスタを介してグローバル配列にタイムスタンプを書き込み、著者らはその計測のコストをレイヤー起動あたり約2.5µs、ステップの2.8%と見積もっており、出力はビット単位で同一のまま、ビルドフラグでコンパイル時に除外できる。自身のプロファイラを搭載しなければならないカーネルは、オペレーティングシステムになったカーネルであり、それはまさにこの記事が用いている枠組みである。プロセスとはすべてのGPUリソースを保持する一つのカーネルであり、スレッドとはプリエンプションのないワープであり、I/O待ちとはフラグに対するビジーウェイトであり、DMAとはmbarrierで追跡されるTMAであり、メッセージとはNVLinkを介した単純なストアである。

ドラフトモデルも速くなりました。そして、それは利点の半分にすぎません。

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.

チームはドラフトにも同じメガカーネル処理を適用した。5つのドラフト層がそれぞれ1回の起動として実行され、さらに他3つも同様で、ドラフトは175カーネルから8カーネルへ、1ステップあたり1.16 msから440 µsへと、ステップのその半分で2.6×の改善となる。受理数は、この投稿が示す最も近い公開比較と並べて見る価値がある。

• 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

• なぜ16なのか — 16個のトークンはより多くのエキスパートを起動するため検証コストが高くなるが、チームは、より狭いチェーンで同じ長さに到達するよりも、より広いツリーで6.6に到達する方が容易だとわかった

• ドラフトスタック — 大型モデル向けの8,000行のCUDAとPTX。コードはfacebookresearch/sgl-muse-megaで公開

比較は注意深く読む必要がある。これは2つのモデル間のベンチマークではない。Xiaomiが公開したものとSGLangチームによるものという、2つの投機的デコーディング構成を、異なるハードウェアと異なるワークロード上で比較したものだ。この比較が支持する唯一の主張は、より広い検証ツリーが同じアクセプタンス長へ到達するための実行可能な道筋であるということであり、これは設計上の主張であってランキングではない。

この数字が何であるか、そして何でないか

ここがキャパシティプランが正しく押さえるべき点であり、投稿自体が正直に認めている点でもある。1,109という数値は、8台のGB300上でテンソル並列度8、バッチサイズ1、投機的デコーディングを有効にした状態での、単一のGSM8Kリクエストにおける20デコードステップの中央値である。投稿は、メガカーネルが現時点ではバッチサイズ1しかサポートしておらず、より大きなバッチへ拡張することが次の作業であると明言している。

これらの制約はそれぞれ別の主張を打ち消している。バッチサイズが1ということは、サービングチームがこの言葉を使う意味でのスループット結果ではない：そこにはバッチ処理による利得も、待ち行列も、プロンプト長の混在もない。単一データセットということは、その数値がGSM8Kのプロンプトと出力の分布に結びついており、さらに著者らが自らのプロファイリングで合成ルーティングとして用いた、層ごとに49エキスパートという測定済みのルーティングパターンにも結びついている。1つのリクエストの20ステップということは、分散がリクエスト数ではなくステップ数によって制限されることを意味する。

そして、単一リクエストはオーバーラップの最悪ケースであり、著者自身もこう述べている。「リクエストが1件では、あらゆる待ちが露わになる。より多ければ、ほぼ常に実行可能な作業がある」。したがって、メガカーネルの2.1×は、レイテンシ隠蔽が最も難しい領域における測定値であり、投稿は、より大きなバッチは難しくなるどころか容易になるはずだと明言している。これは結果に有利な論拠ではあるが、依然として測定ではなく論拠である。

今日知り得ることは、狭いが現実的だ。コードは公開リポジトリに存在し、投稿は検証できる十分な詳細で仕組みを説明し、内部計算は3つの独立した比率で一致し、著者らはさらなる改善を制約する3つのオーバーヘッドを文書化している。知り得ないのは、独立したチームが自前の8台のGB300で1,109を出せるか、この手法が実際のバッチで耐えうるか、そしてリリース済みのSGLangビルドに入るかどうかだ。投稿はそう主張していないし、リポジトリはインストールできるパッケージではなく研究用の成果物だ。最も決着をつける比較は、同じハードウェア上のチューニング済みベースラインとの比較だ。なぜなら「SGLangのデフォルトチェーン」は動く標的であり、メガカーネルを書いたチームが、それが打ち負かすベースラインもチューニングしているからだ。

今日利用できる毎秒169トークンと、それがどう関係するのか

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.

この差は、すでにハードウェアを所有していない人にとって、この話の中で最も有用なものだ。エンドポイントが毎秒169トークンで応答し、手作りした8-GPUスタックが1,109で応答するなら、その差はスタックを構築する理由にはならない。GB300を8台というのは数百万単位の設備投資であり、カーネルは研究用コードであり、バッチサイズ1という制約は、このスタックが現在、ホスト型エンドポイントがすでにうまく処理しているワークロード、つまり一度に1リクエストというワークロードにまさに適していることを意味する。メガカーネルが本当に計算を変えるのは、訓練用にGPUをすでに購入したチーム、あるいはデータレジデンシー制約のあるワークロードを抱え、低バッチでどれだけの利用率を得られるかを今問うているチームにとってだ。そのチームにとって、同じシリコン上で同じ数値演算による2.1倍は、現実の請求書が付いた現実の数字であり、その請求書はすでに支払われている。

それ以外の人々にとっての率直な解釈は、この投稿が今週導入できる製品ではなく、推論インフラがどこへ向かっているかを示すシグナルだということだ。その方向性は、同じラボがこの一年を通じて公表してきた内容と一致している。より少なく、より大きなカーネル、スケジューリングの判断の多くをカーネル内部へ移すこと、そしてGPUを順番に実行する一連のライブラリ呼び出しではなく、プログラムされるべき機械として扱うことだ。

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.

何がこのピースを変えるでしょうか？

4つの事柄を、状況への影響が大きい順に挙げます。

• 1より大きいバッチ。著者らは次はより大きなバッチだと述べており、そのほうが簡単だと主張している。なぜなら、より多くのトークンが処理中にあると、前の演算の出力の一部はすでに準備できている一方で、残りはまだ到着しているからだ。バッチ化された数値が存在する瞬間、比較は「1つのストリーム対1つのストリーム」ではなくなり、サービングチームが使えるスループットの主張になる。

• 独立した再現。ここに示す数値はすべて、著者自身のハーネス、著者自身のハードウェア、そしてプロファイリング実行用の著者自身の合成ルーティングによるものである。これはカーネル開発ではよくあることであり、依然として単一の情報源にすぎない。別のチームが1,109を一致させるか、あるいは一致させられないか——それがこれを決着させる事実である。

• それを搭載したリリース済みのSGLangビルド。投稿はコードを公開し、設計について説明しているが、このモデルにおいてメガカーネルが出荷版でデフォルトのデコード経路を置き換えたとは述べていない。そうなるまでは、それを採用することは研究用リポジトリを採用することを意味する。

• GEMM-2 から combine までのギャップ解消。著者らはすでに、そこでのブロックあたり 12.5 µs のアイドル時間を、残りのほぼ半分と特定している。その連鎖が修正されれば、同じハードウェアでもステップ時間は有意に改善し、2.1× は見出しではなく、より大きな数値の下限となる。

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.