Muse Spark Open: 1.109 token mỗi giây là thay đổi ở kernel, không phải thay đổi ở mô hình

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.

Mức cơ sở không ai đánh giá chuẩn: 1,450 lần khởi chạy mỗi 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.

Đó là cái giá mà megakernel nhắm vào, và các tác giả mượn một phép so sánh từ thập niên 1960 để giải thích điều đó: những cỗ máy thời đầu chạy các công việc theo lô, một công việc chiếm toàn bộ máy, và bộ xử lý ngồi không trong khi băng từ quay. Đa chương trình đã khắc phục điều này bằng cách giữ nhiều công việc thường trú và chuyển sang công việc khác mỗi khi công việc hiện tại bị chặn. Megakernel là ý tưởng đó được chuyển lên GPU — một kernel thường trú trên mọi SM và chạy nhiều công việc cùng lúc, nhờ đó các warp có thể bắt đầu công việc của thao tác tiếp theo trước khi thao tác trước đó hoàn tất.

Hiệu ứng đo được, từ hình đi kèm phần hiệu năng, là một bước giải mã giảm từ 12,9 ms xuống 5,9 ms. Hình đó chia bước này thành hai nửa nháp và xác minh: 1,16 ms và 10,9 ms với 1.450 kernel ở đường cơ sở, 0,44 ms và 5,3 ms với 97 kernel trong megakernel. Từ mười hai phẩy chín xuống năm phẩy chín là khoảng 2,2×, và 1.109 trên 517 là khoảng 2,1×, nên hai tỷ lệ chính khớp với nhau. Có một phép kiểm tra thứ ba mà bất kỳ ai sẵn lòng thực hiện phép chia đều có thể dùng: bước này xác minh 16 token và chấp nhận 6,6, vậy 6,6 token trong 5,9 ms tương đương khoảng 1.120 token mỗi giây. Phép tính nhất quán nội tại, đây là điều đầu tiên cần xác lập về một con số như thế này và không phải điều cuối cùng.

Một điều mà những con số không cho biết là tốc độ đạt được trên mỗi GPU là bao nhiêu. Bài đăng báo cáo thông lượng tổng hợp cho toàn bộ triển khai tám GPU. Chia cho tám, tức khoảng 139 token mỗi giây trên mỗi GPU — phép tính của tôi dựa trên con số của họ, không phải con số mà bài đăng công bố, và không thể so sánh với con số trên mỗi GPU mà bạn có thể thấy được trích dẫn cho một mô hình khác trên một cỗ máy khác.

Ranh giới kernel nằm ở đâu, và tại sao nó lại dừng ở đó

Megakernel không nuốt trọn toàn bộ mô hình. Mỗi lần khởi chạy bao phủ từ phần attention của một lớp decoder đến phép chiếu qkv và ghi KV của lớp kế tiếp, đây là ranh giới có chủ đích chứ không phải hạn chế của nerve. Lý do nó hiệu quả là phần lớn những gì attention của một lớp đọc — lịch sử KV — đã được ghi ở các bước trước đó và không phụ thuộc vào bất kỳ thứ gì còn đang xử lý dở, nên một lần khởi chạy mới có thể bắt đầu nạp nó ngay khi nó sẵn sàng, trước khi đầu ra của lớp trước tồn tại. Lớp đầu tiên và lớp cuối cùng có các phép toán hơi khác nhau, và bài viết cũng ghi chú điều đó.

Trong một lần khởi chạy, cấu trúc là cố định: một khối luồng gồm mười hai warp trên mỗi SM trong số 144 SM, được khởi chạy theo các cụm bốn, sử dụng bộ nhớ chia sẻ phân tán khi có thể thực hiện phép giảm với độ trễ thấp hơn. Mười hai warp không thể hoán đổi cho nhau. Warp 0 tải KV và trọng số, đồng thời dùng thời gian rảnh để prefetch trọng số của phép toán tiếp theo. Warp 1 và 2 chạy các activation và tensor core. Warp 3 đến 11 chia nhau xử lý attention, các dense GEMM, GEMM chuyên gia và các phép nhận all-reduce và all-gather. Trong các lớp cửa sổ trượt, warp 4 đến 11 tự tải KV của riêng mình thay vì chờ 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.

Chi tiết đó quan trọng hơn vẻ ngoài của nó. {{1}}Các nhóm suy luận thường chấp nhận rằng một kernel nhanh hơn tạo ra các con số hơi khác đi và lặng lẽ tinh chỉnh lại ngưỡng đánh giá của họ.{{/1}} {{2}}Một kernel bằng nhau từng bit với cách song song hóa phía huấn luyện loại bỏ hoàn toàn lập luận đó, và đó là kiểu thuộc tính khiến một tối ưu được áp dụng chứ không chỉ được ngưỡng mộ.{{/2}}

Ba hóa đơn mà một nhân cỡ này phải trả

Megakernel không miễn phí. Bài viết thẳng thắn một cách bất thường về các chi phí phụ trội, và đó chính là phần đáng đọc lại hai lần, vì chúng cũng là phần cho bạn biết còn bao nhiêu dư địa.

Chi phí đầu tiên nằm ở bộ nhớ đệm lệnh. Một SM đệm khoảng 128 KB lệnh, và mỗi kernel này có 215 đến 314 KB mã máy, vì mỗi kernel chứa mọi giai đoạn của một lớp. Một lần khởi chạy chỉ đưa 32 đến 64 KB đầu tiên đến gần các SM; mọi thứ vượt quá cửa sổ đó được tìm nạp theo từng dòng khi dùng lần đầu. Các tác giả đã chạy một thí nghiệm sạch về điều này và báo cáo: hai khối 100 KB, mỗi khối mất 12.4 µs khi chạy riêng, mất 19.5 µs khi được ghép nối tiếp nhau, và việc thu nhỏ một khối xuống 20 KB đưa khối còn lại trở về tốc độ khi chạy một mình. Lệnh đi qua cùng bộ nhớ đệm L2 như dữ liệu, vì vậy một kernel vượt quá bộ nhớ đệm lệnh sẽ phải cạnh tranh với chính các lần tải trọng số của nó.

Khoản chi phí thứ hai là nạp trước trọng số, và bài viết coi đó là một bài toán bố trí với ba lời giải. Bộ nhớ chia sẻ có một vòng đệm 5 × 32 KB mỗi SM, khoảng 23 MB trên toàn GPU. Phần còn lại đi vào L2, và L2 của GB300 chứa khoảng 135 MB, dù bất cứ thứ gì trong L2 đều cần một chuyến thứ hai, ngắn hơn, vào bộ nhớ chia sẻ và L2 được chia sẻ với mọi thứ khác đang xử lý. Hai cách này xếp chồng lên nhau: gần cuối một lớp, vòng đệm đã chứa các trọng số qkv đầu tiên của lớp kế tiếp. Trên hết, warp 3 đọc xem bước trước đã định tuyến đến những expert nào và kéo các lát của GPU này thuộc 14 expert được dùng nhiều nhất vào L2, khoảng 50 MB mỗi GPU, đặt cược rằng các bước liên tiếp định tuyến giống nhau. Canh bạc đó đáng giá khoảng 0,2 ms trong một bước 5,3 ms, và có thể tốn băng thông khi nó sai.

Hóa đơn thứ ba là đồng bộ hóa, và đây là hạng mục không có câu trả lời gọn gàng. Khi một warp chờ một GPU khác qua NVLink, không có ngắt và không có wait-on-address trong phần cứng; warp thăm dò một cờ cho đến khi nó thay đổi. Các microbenchmark trong bài đặt một poll acquire phạm vi GPU ở 141 đến 145 ns, một cờ một chiều giữa hai GPU trên cùng một node ở khoảng 1,9 µs, và một poll xếp hàng sau các bản sao chép khối lượng lớn ở 1,0 đến 2,6 µs. Fence tốn 0,4 µs khi rảnh và 3,6 đến 4,0 µs khi có các load TMA đang chạy dở. Ngay cả khi công việc được phân bố hoàn hảo, các GPU vẫn có thể lệch nhau tới 4 µs, và mọi điểm dừng trong lớp đều chờ GPU chậm nhất. Megakernel thay các fence bằng các giá trị sentinel được ghi vào các slot store thông thường — một cờ chứa dữ liệu, +0/−0 cho all-reduce, NaN toàn số một cho các gather bf16 — và đặt lại bản sao rảnh của mỗi slot trong khi bản còn lại đang được dùng.

Đây là những con số khiến kết quả giống với kỹ thuật hơn là phép màu: độ lệch GPU 4 µs so với một bước 5,9 ms là nhỏ, nhưng bốn điểm dừng đồng bộ hóa trên mỗi lớp nhân lên trên 62 lớp đúng là kiểu chi phí phụ trội có thể nuốt mất lợi thế gấp 2× nếu bạn làm sai.

5.9 ms vẫn còn được dành cho việc gì

Bài viết có cả phần phê bình của chính nó, và câu hữu ích nhất trong đó nói về khoảng trống sau GEMM thứ hai. Các block hoàn thành GEMM đó ở đâu đó trong khoảng 55 đến 64 µs tính từ đầu layer, theo phân vị thứ 10 đến thứ 90, nhưng bước combine phải chờ đầu ra của mọi GPU, nên nó không thể bắt đầu trước 65 µs. Từ lúc kết thúc GEMM 2 đến cuối layer, một block không có việc gì để làm trong trung bình khoảng 12,5 µs — gần một nửa tổng thời gian nhàn rỗi của nó. Chính chuỗi đó, chứ không phải chi phí khởi chạy kernel mà megakernel đã loại bỏ, là thứ mà các tác giả nói họ đang xem xét tiếp theo.

Họ cũng giải thích vì sao attention của lớp tiếp theo không thể đơn giản bắt đầu sớm hơn. PDL khởi chạy nó ở 79 µs, khoảng 4 µs trước khi block đầu tiên thoát ra, nhưng các block của nó chỉ được đưa vào chạy khi các block trước rời khỏi SM của chúng ở 84 µs, nên khởi chạy sớm hơn cũng không làm bất cứ thứ gì bắt đầu sớm hơn. Cách khắc phục là nạp trước các hàng KV cũ của nó vào L2 ở 65 µs, để chúng chờ ở đó. Đó là cùng một kiểu mẫu như nạp trước weight: khi bạn không thể loại bỏ một khoảng chờ, bạn chuyển công việc vào trong đó.

Profiling là vấn đề cấu trúc còn lại, và đó là một vấn đề hay: kernel trông giống như một thanh dài cho mỗi layer đối với bất kỳ profiler bên ngoài nào, nên nó tự đo thời gian chính mình. Lane 0 của mỗi block ghi một dấu thời gian vào một mảng toàn cục thông qua thanh ghi globaltimer, và các tác giả đo chi phí của việc đánh dấu này vào khoảng 2,5 µs mỗi lần launch layer, 2,8% một step, với đầu ra vẫn giống hệt từng bit và có một cờ build để loại bỏ nó khi biên dịch. Một kernel phải mang theo profiler của riêng nó là một kernel đã trở thành một hệ điều hành, đúng như cách diễn đạt mà bài viết sử dụng: một process là một kernel nắm giữ toàn bộ tài nguyên GPU, thread là các warp không có preemption, chờ I/O là busy-wait trên một cờ, DMA là TMA được theo dõi bởi mbarrier, message chỉ là các thao tác ghi thông thường qua NVLink.

Mô hình nháp cũng nhanh hơn, và đó là một nửa thành công

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.

Nhóm đã áp dụng cùng cách xử lý megakernel cho phần draft: năm lớp draft, mỗi lớp chạy như một lần khởi chạy, cộng thêm ba lớp khác, giúp đưa phần draft từ 175 kernel xuống còn 8 và từ 1.16 ms mỗi bước xuống còn 440 µs, tức cải thiện 2.6× ở nửa đó của bước. Các con số chấp nhận đáng được đặt cạnh phép so sánh công khai gần nhất mà bài viết đưa ra.

• 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

• Tại sao 16 — mười sáu token tốn nhiều chi phí xác minh hơn vì chúng kích hoạt nhiều chuyên gia hơn, nhưng nhóm nhận thấy dễ đạt 6.6 với một cây rộng hơn so với việc đạt cùng độ dài bằng một chuỗi hẹp hơn

• Ngăn xếp bản nháp — 8000 dòng CUDA và PTX cho mô hình lớn, với mã được công bố tại facebookresearch/sgl-muse-mega

Hãy đọc kỹ phần so sánh này, vì đây không phải là một phép đánh giá chuẩn giữa hai mô hình. Đây là so sánh giữa hai cấu hình giải mã suy đoán, một do Xiaomi công bố và một do nhóm SGLang công bố, trên phần cứng khác nhau và khối lượng công việc khác nhau. Khẳng định duy nhất mà nó hỗ trợ là một cây xác minh rộng hơn là một hướng khả thi để đạt cùng độ dài chấp nhận, đây là một khẳng định về thiết kế chứ không phải là một bảng xếp hạng.

Con số này là gì và nó không phải là gì

Đây là phần mà một kế hoạch dung lượng phải làm đúng, và cũng là phần mà chính bài viết thừa nhận một cách thẳng thắn. Con số 1.109 là trung vị của 20 bước giải mã trên một yêu cầu GSM8K duy nhất, ở độ song song tensor là 8, trên tám GB300, với kích thước lô là một và bật giải mã suy đoán. Bài viết nói rõ rằng megakernel hiện nay chỉ hỗ trợ kích thước lô là một, và việc mở rộng nó cho các lô lớn hơn là hạng mục công việc tiếp theo.

Mỗi ràng buộc đó loại bỏ một khẳng định khác nhau. Kích thước batch bằng một nghĩa là đây không phải là kết quả về thông lượng theo nghĩa mà một đội ngũ serving dùng từ này: không có lợi ích gộp batch trong đó, không có hàng đợi, không có sự pha trộn độ dài prompt. Một bộ dữ liệu nghĩa là con số này gắn với phân phối prompt và đầu ra của GSM8K, và gắn với một mẫu định tuyến đo được của 49 chuyên gia mỗi lớp mà sau đó các tác giả đã dùng làm định tuyến tổng hợp trong phần profiling của riêng họ. Hai mươi bước của một yêu cầu nghĩa là phương sai bị giới hạn bởi số bước, không phải bởi số yêu cầu.

Và một yêu cầu đơn lẻ là trường hợp tệ nhất đối với chồng lấp, điều mà chính các tác giả nói: "với một yêu cầu, mọi khoảng chờ đều bị phơi ra; với nhiều yêu cầu hơn, gần như luôn có công việc sẵn sàng." Vì vậy, con số 2,1× của megakernel là một phép đo chế độ mà việc ẩn độ trễ khó nhất, và bài viết nói rõ rằng các lô lớn hơn phải dễ hơn chứ không phải khó hơn. Đó là một lập luận ủng hộ kết quả, và nó vẫn là một lập luận chứ không phải một phép đo.

Vậy thì, những gì có thể biết được hôm nay là hẹp và có thật: mã nguồn tồn tại trong một kho lưu trữ công khai, bài đăng mô tả cơ chế đủ chi tiết để có thể kiểm chứng, phần tính toán nội bộ khớp nhau qua ba tỷ lệ độc lập, và các tác giả đã ghi lại ba chi phí phụ trội giới hạn mức lợi ích tiếp theo. Những gì không thể biết là liệu một nhóm độc lập có đạt 1.109 trên tám GB300 của riêng họ hay không, liệu cách tiếp cận có trụ được qua một lô thực tế hay không, và liệu nó có được đưa vào một bản dựng SGLang đã phát hành hay không — bài đăng không tuyên bố là đã có, và kho lưu trữ là một sản phẩm nghiên cứu chứ không phải một gói bạn có thể cài đặt. Phép so sánh có thể giải quyết được nhiều nhất là so với một baseline đã được tinh chỉnh trên cùng phần cứng, vì "chuỗi mặc định của SGLang" là một mục tiêu luôn thay đổi và nhóm đã viết megakernel cũng chính là nhóm tinh chỉnh baseline mà nó đánh bại.

Nó liên quan thế nào đến 169 token mỗi giây mà bạn có thể thuê ngay hôm nay

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.

Khoảng cách là thứ hữu ích nhất trong câu chuyện này đối với bất kỳ ai chưa sở hữu phần cứng đó. Nếu một endpoint trả lời ở 169 token mỗi giây và một ngăn xếp tám GPU tự dựng trả lời ở 1.109, thì khác biệt đó không phải lý do để dựng ngăn xếp. Tám GB300 là một khoản mua vốn tính bằng hàng triệu, kernel là mã nghiên cứu, và ràng buộc kích thước batch bằng một nghĩa là ngăn xếp hiện tại chỉ giỏi đúng loại khối lượng công việc mà một endpoint được host đã xử lý tốt: mỗi lần một yêu cầu. Nơi megakernel thực sự thay đổi bài toán là với một nhóm đã mua GPU để huấn luyện hoặc cho một khối lượng công việc có ràng buộc về lưu trú dữ liệu, và giờ đang hỏi có thể tận dụng được bao nhiêu từ chúng ở batch thấp. Với nhóm đó, 2,1× trên cùng phần cứng silicon với cùng phép toán số học là một con số thật gắn với một hóa đơn thật, và hóa đơn đã được thanh toán.

Đối với mọi người khác, cách hiểu trung thực là bài đăng này là một tín hiệu về hướng đi của hạ tầng suy luận, chứ không phải một sản phẩm mà bạn có thể áp dụng ngay trong tuần này. Hướng đi này nhất quán với những gì cùng phòng thí nghiệm đó đã công bố suốt cả năm nay: ít kernel hơn nhưng lớn hơn, nhiều quyết định lập lịch hơn được chuyển vào bên trong kernel, và GPU được coi như một cỗ máy cần được lập trình thay vì một tập hợp các lệnh gọi thư viện cần được sắp xếp tuần tự.

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.

Điều gì sẽ thay đổi phần này

Bốn điều, xếp theo thứ tự giảm dần về mức độ chúng có thể làm thay đổi cục diện.

• Một lô lớn hơn một. Các tác giả nói rằng các lô lớn hơn là bước tiếp theo và lập luận rằng chúng dễ hơn, vì khi có nhiều token đang được xử lý đồng thời, một số đầu ra của phép toán trước đó đã sẵn sàng trong khi phần còn lại vẫn đang đến. Ngay khi có một con số được gộp lô, phép so sánh không còn là "một luồng so với một luồng" nữa mà bắt đầu trở thành một tuyên bố về thông lượng mà một đội vận hành phục vụ có thể dùng.

• {{1}}Một bản tái tạo độc lập.{{/1}} Mọi số liệu ở đây đều đến từ khung kiểm thử của chính các tác giả, trên phần cứng của chính họ, với định tuyến tổng hợp của riêng họ cho lượt phân tích hiệu năng. Điều đó là bình thường đối với công việc kernel và nó vẫn chỉ là một nguồn duy nhất. Việc một nhóm thứ hai tái lập được 1,109, hoặc không làm được, mới là dữ kiện có thể giải quyết dứt điểm.

• Một bản dựng SGLang đã phát hành có chứa nó. Bài đăng công bố mã và mô tả thiết kế; bài viết không nói rằng megakernel đã thay thế đường giải mã mặc định cho mô hình này trong một phiên bản đã phát hành. Cho đến khi điều đó xảy ra, việc áp dụng nó đồng nghĩa với việc áp dụng một kho nghiên cứu.

• Việc thu hẹp khoảng trống GEMM-2-to-combine. Các tác giả đã xác định 12.5 µs thời gian nhàn rỗi trên mỗi khối ở đó là gần một nửa phần còn lại. Nếu chuỗi đó được khắc phục, cùng phần cứng đó sẽ tạo ra thời gian bước tốt hơn đáng kể, và mức 2.1× sẽ trở thành mức sàn của một con số lớn hơn thay vì là điểm nhấn chính.

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.