Muse Spark Open：每秒 1,109 個 Token 是核心層面的變更，而非模型的變更

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.

這正是 megakernel 所要對付的成本，而作者們借用一九六〇年代的類比來解釋：早期機器執行批次工作，一個工作獨佔整台機器，當磁帶轉動時處理器就閒置著。多重程式設計（multiprogramming）解決了這個問題，做法是讓數個工作常駐，每當目前的工作受阻就切換到另一個。megakernel 就是把這個想法搬到 GPU 上——一個 kernel 常駐在每個 SM 上，同時執行許多工作，於是 warp 可以在前一個運算完成之前就開始下一個運算的工作。

根據效能章節所附圖表，量測到的效果是解碼步驟從 12.9 ms 縮短到 5.9 ms。該圖將這個步驟拆成草稿與驗證兩半：基線版本有 1,450 個核心時，分別是 1.16 ms 與 10.9 ms；巨型核心版本有 97 個核心時，則是 0.44 ms 與 5.3 ms。十二點九到五點九約為 2.2 倍，而 1,109 對 517 約為 2.1 倍，所以這兩個主要比率彼此吻合。任何願意做除法的人還能做第三項檢查：這個步驟驗證 16 個 token、接受 6.6 個，因此 5.9 ms 內處理 6.6 個 token，換算下來大約是每秒 1,120 個 token。算術本身內部一致，這是面對這類數字時首先要確認的事，但不是最後一件。

有件事這些數字沒有說清楚：速度分攤到每顆 GPU 上究竟是多少。該貼文報導的是整個八 GPU 部署的總吞吐量。除以八，大約是每顆 GPU 每秒 139 個 token——這是我根據他們提供的數字算出來的，不是該貼文公布的數字，也不能拿來和你在另一台機器上、另一個模型可能看到的每 GPU 數字相比。

核心邊界位於何處，以及它為何止於那裡

megakernel 並不會吞噬整個模型。每次啟動涵蓋一個解碼器層的注意力，直到下一層的 qkv 投影與 KV 寫入；這是刻意設定的邊界，而非 nerve 的限制。它能運作的原因在於，一層注意力所讀取的大部分內容——KV 歷史——是在較早步驟中寫入的，且不依賴任何仍在進行中的事物，因此新的啟動可以在其可用時立即開始載入，甚至早於前一層的輸出存在。第一層與最後一層的操作略有不同，文章也提到了這一點。

在一次啟動內，形狀是固定的：在 144 個 SM 上各有一個由 12 個 warp 組成的執行緒區塊，並以四個為一組的叢集啟動；在可取得較低延遲縮減的地方使用分散式共享記憶體。這十二個 warp 不可互換。Warp 0 載入 KV 與權重，並利用其餘裕預取下一項作業的權重。Warp 1 和 2 執行激活值與 Tensor Core 運算。Warp 3 到 11 分擔注意力、稠密 GEMM、專家 GEMM，以及 all-reduce 與 all-gather 的接收。在滑動視窗層中，warp 4 到 11 自行載入其 KV，而不是等待 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.

這個細節的重要性超乎表面。推論團隊通常接受較快的核心會產生略有不同的數值，並默默重新調整評估門檻。一個與訓練端平行化在位元層級完全相同的核心，能完全消除這個爭論，而這正是能讓最佳化被採用、而不只是被讚賞的那種特性。

這麼大的核心得付三筆帳單

巨型核心並非沒有代價。這篇文章對額外開銷異常坦率，而這些開銷正是值得讀兩遍的部分，因為它們也是告訴你還有多大餘裕的部分。

第一筆代價是指令快取。一顆 SM 大約快取 128 KB 的指令，而這些核心每個是 215 到 314 KB 的機器碼，因為每個核心都包含一層的每一個階段。一次啟動只會把前 32 到 64 KB 帶到接近 SM 的位置；超出這個窗口的一切，都會在首次使用時逐條快取線抓取。作者對此做了一個乾淨的實驗並報告了結果：兩個各 100 KB 的區塊，單獨執行時各耗時 12.4 µs，但背靠背拼接起來後耗時 19.5 µs；把其中一個區塊縮小到 20 KB，就讓另一個區塊恢復到單獨執行時的速度。指令和資料一樣會經過同一個 L2 快取，因此一個超出指令快取容量的核心，最後會與自身的權重載入爭奪資源。

第二筆帳是權重預取，文章把它當成一個有三種答案的放置問題。共享記憶體讓每個 SM 獲得一個 5 × 32 KB 的環形緩衝區，整顆 GPU 約 23 MB。其餘的則放進 L2，而 GB300 的 L2 可容納約 135 MB，不過 L2 裡的任何資料都還需要再走一趟較短的旅程進入共享記憶體，而且 L2 會與其他所有正在執行中的工作共用。兩者會疊加使用：在一層接近尾聲時，環形緩衝區已持有下一層第一批 qkv 權重。在此之上，warp 3 會讀取前一步路由到哪些專家，並把最常用的 14 個專家中屬於這顆 GPU 的切片拉進 L2，每顆 GPU 約 50 MB，賭的是連續步驟的路由結果相似。這個賭注在 5.3 ms 的一個步驟中約值 0.2 ms，但一旦賭錯就可能耗費頻寬。

第三道帳單是同步，而這道帳單沒有乾淨的答案。當一個 warp 透過 NVLink 等待另一個 GPU 時，硬體中既沒有中斷，也沒有 wait-on-address；warp 會輪詢一個旗標，直到它改變。文章中的微基準測試顯示，GPU 範圍的 acquire 輪詢為 141 到 145 ns，同一節點上兩個 GPU 之間的單向旗標約為 1.9 µs，而排在大量複製後面的輪詢為 1.0 到 2.6 µs。Fence 在閒置時耗費 0.4 µs，在有 TMA 載入進行中時耗費 3.6 到 4.0 µs。即使工作完美分散，GPU 之間仍可能彼此落後最多 4 µs，而這一層中的每一次停止都要等待最慢的那一個。Megakernel 以寫入一般儲存槽的哨兵值取代 fence——一個填有資料的旗標、用於 all-reduce 的 +0/−0、用於 bf16 gather 的全 1 NaN——並在每個槽的另一份使用中時重置其閒置副本。

這些數字讓結果更像是工程，而不是魔法：相對於 5.9 毫秒的步進，4 微秒的 GPU 偏差算小，但每層四次同步停頓、再乘上 62 層，正是那種一旦處理失誤就會吃掉 2× 加速效益的額外開銷。

這 5.9 毫秒還花在什麼地方

這篇貼文包含了對自身的評論，其中最有用的說法是關於第二次 GEMM 之後的空檔。各區塊會在該層進行到 55 到 64 µs 之間完成該 GEMM（第 10 到第 90 百分位），但合併必須等待每個 GPU 的輸出，因此要到 65 µs 才能開始。從 GEMM 2 結束到該層結束，一個區塊平均約有 12.5 µs 閒置無工作——幾乎占其總閒置時間的一半。作者表示，他們接下來要研究的是這條鏈，而不是 megakernel 已經消除的核心啟動開銷。

他們也解釋了為什麼下一層的注意力不能單純提早開始。PDL 在 79 µs 時啟動它，比第一個區塊退出早約 4 µs，但它的區塊要等到前一組區塊在 84 µs 離開 SM 後才會進駐，因此更早啟動也不會讓任何工作更早開始。變通做法是在 65 µs 時將其舊的 KV 列預取到 L2，讓它們在那裡等待。這與權重預取是相同的模式：當你無法消除等待時，就把工作移進等待裡。

效能剖析是另一個結構性問題，而且還是個好問題：對任何外部剖析器來說，kernel 在每一層看起來都像一條長條，所以它乾脆自己計時。每個 block 的 lane 0 會透過 globaltimer 暫存器把時間戳寫到全域陣列，而作者測得這套 instrumentation 的成本約為每次 layer 啟動 2.5 µs、占一個 step 的 2.8%，輸出仍保持位元完全相同，並有建置旗標可將其編譯移除。一個必須自帶剖析器的 kernel，就是一個已經變成作業系統的 kernel，而這正是文章採用的說法：行程就是一個持有所有 GPU 資源的 kernel，執行緒就是不可搶占的 warp，等待 I/O 就是對旗標忙等，DMA 是由 mbarrier 追蹤的 TMA，訊息就是跨 NVLink 的一般 store。

草稿模型也變快了，而這已經是成功的一半。

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.

團隊對草稿套用了同樣的 megakernel 處理：五個草稿層各自以一次啟動執行，再加上另外三個，讓草稿部分從 175 個 kernel 降到 8 個，每步從 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——十六個 token 的驗證成本更高，因為它們會啟用更多專家，但團隊發現，用更寬的樹達到 6.6 比用更窄的鏈達到相同長度更容易

• 草稿堆疊 — 為大型模型撰寫的 8000 行 CUDA 與 PTX，程式碼發佈於 facebookresearch/sgl-muse-mega

請仔細閱讀這項比較，因為這並不是兩個模型之間的基準測試。它比較的是兩種推測解碼配置，一種由小米發表，另一種由 SGLang 團隊提出，兩者運行於不同的硬體與不同的工作負載。它唯一能支持的論點是：更寬的驗證樹是達到相同接受長度的可行途徑，而這是一項設計主張，並非排名。

這個數字是什麼，以及它不是什麼

這正是容量規劃必須弄對的部分，也是文章本身坦承的部分。1,109 這個數字，是在八張 GB300 上、以張量平行度 8、批次大小為一並開啟推測解碼時，單一 GSM8K 請求的 20 個解碼步驟的中位數。文章明白指出，這個 megakernel 目前只支援批次大小為一，而將其擴展到更大的批次是下一步工作。

這些限制各自排除了不同的主張。批次大小為一意味著，這不是推論服務團隊使用該詞意義下的吞吐量結果：其中沒有批次處理增益、沒有佇列，也沒有混合不同提示長度。單一資料集意味著這個數字綁定於 GSM8K 的提示與輸出分布，以及綁定於每層 49 個專家的實測路由模式，而作者隨後在自己的效能剖析中將其用作合成路由。單一請求的二十個步驟意味著變異量受步驟數限制，而不是受請求數限制。

而單一請求是重疊的最壞情況，作者自己也這麼說：「只有一個請求時，每一次等待都會顯露出來；請求更多時，幾乎總是有已就緒的工作。」因此，megakernel 的 2.1× 是對延遲隱藏最困難之情境的量測，而文章明確表示，較大的批次應該更容易，而非更困難。這是支持該結果的論點，而且仍然只是論點，而非量測。

那麼，今天可知的事情範圍有限但真實：程式碼存在於公開儲存庫中，文章對機制的描述詳細到足以查核，內部算術在三個獨立比率之間彼此一致，而作者也記錄了三項限制進一步增益的額外開銷。無法得知的是，獨立團隊能否在自己的八張 GB300s 上達到 1,109，該方法能否經得起真實批次，以及它能否進入已發布的 SGLang 版本——文章並未聲稱已經如此，而該儲存庫是研究產物，不是可直接安裝的套件。最能釐清問題的比較，是在相同硬體上對照調校過的基準，因為「SGLang 的預設鏈」是不斷變動的目標，而且撰寫巨型核心的團隊也調校了它所擊敗的基準。

它與你今日可租用的每秒 169 個 token 有何關係

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 個 token，而手工打造、由八張 GPU 組成的堆疊每秒回應 1,109 個，這個差異並不是打造該堆疊的理由。八張 GB300 是一筆以百萬計的資本支出，這個 kernel 是研究用程式碼，而批次大小為 1 的限制意味著，這套堆疊目前恰好擅長託管端點早已處理得很好的工作負載：一次一個請求。megakernel 真正改變盤算的地方，是對於已經為了訓練買下這些 GPU、或是有資料駐留限制的工作負載的團隊，而他們現在問的是：在低批次下能從這些 GPU 榨出多少使用率。對那個團隊來說，在同一批晶片上、以相同的數值運算達到 2.1×，是一個實實在在的數字，而且背後有一張真實的帳單，而這張帳單已經付清了。

對其他人來說，誠實的解讀是：這篇文章是一個關於推論基礎設施走向的訊號，而不是一個你這週就能採用的產品。這個方向與同一個實驗室整年來持續發表的內容一致：更少、更大的 kernel，更多排程決策被移進 kernel 內部，以及把 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.

什麼會改變這件作品？

四件事，依其能多大程度改變局面，由高到低排列。

• 批次大於一。作者表示接下來會採用更大的批次，並主張這樣更容易，因為有更多 token 同時在處理中時，前一個運算的部分輸出已經就緒，其餘的仍在送達。一旦存在批次化的數字，比較就不再是「單一串流對單一串流」，而開始成為服務團隊可用的輸送量主張。

• 獨立重現。這裡的每一個數字都來自作者自己的測試框架、他們自己的硬體，以及他們為效能剖析階段自行設計的合成路由。這在 kernel 工作中很正常，而且它仍然只是單一來源。第二個團隊若能重現 1,109，或無法重現，才是能一槌定音的事實。

• 一個已發行且包含此功能的 SGLang 建置版本。該文章公開了程式碼並說明了設計；它並未表示 megakernel 已在某個已發行版本中取代此模型的預設解碼路徑。在它這麼做之前，採用它就意味著採用一個研究用儲存庫。

• GEMM-2-to-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.