Muse Spark Open: 1,109 Token per Detik Adalah Perubahan Kernel, Bukan Perubahan Model

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.

Baseline yang tidak dijadikan tolok ukur oleh siapa pun: 1,450 peluncuran per 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.

Itulah biaya yang diserang oleh megakernel, dan para penulis menggunakan analogi tahun 1960-an untuk menjelaskannya: mesin-mesin awal menjalankan job batch, dengan satu job memiliki seluruh mesin, dan prosesor menganggur sementara pita berputar. Multiprogramming memperbaikinya dengan menjaga beberapa job tetap residen dan beralih ke job lain setiap kali job saat ini terhambat. Megakernel adalah ide itu yang dipindahkan ke GPU — satu kernel tetap residen di setiap SM dan menjalankan banyak job sekaligus, sehingga warp dapat memulai pekerjaan operasi berikutnya sebelum operasi sebelumnya selesai.

Efek terukur, dari gambar yang menyertai bagian performa, adalah langkah decode yang turun dari 12,9 ms menjadi 5,9 ms. Gambar tersebut memecah langkah itu menjadi bagian draft dan verifikasi: 1,16 ms dan 10,9 ms dengan 1.450 kernel pada baseline, 0,44 ms dan 5,3 ms dengan 97 kernel pada megakernel. Dari 12,9 ke 5,9 sekitar 2,2×, dan 1.109 dibagi 517 sekitar 2,1×, sehingga kedua rasio utama itu saling sesuai. Ada pemeriksaan ketiga yang tersedia bagi siapa pun yang mau melakukan pembagian: langkah tersebut memverifikasi 16 token dan menerima 6,6, jadi 6,6 token dalam 5,9 ms menghasilkan sekitar 1.120 token per detik. Aritmetika ini konsisten secara internal, yang merupakan hal pertama yang perlu ditetapkan tentang angka seperti ini dan bukan yang terakhir.

Satu hal yang tidak disebutkan oleh angka-angka itu adalah berapa kecepatan yang dihasilkan per GPU. Tulisan tersebut melaporkan throughput agregat untuk keseluruhan deployment delapan GPU. Dibagi delapan, hasilnya sekitar 139 token per detik per GPU — hitungan saya sendiri atas angka mereka, bukan angka yang diterbitkan dalam tulisan tersebut, dan tidak dapat dibandingkan dengan angka per GPU yang mungkin Anda lihat dikutip untuk model lain di mesin lain.

Di mana batas kernel berada, dan mengapa ia berhenti di sana

Megakernel tidak menelan seluruh model. Setiap peluncuran mencakup attention satu layer decoder hingga proyeksi qkv dan penulisan KV layer berikutnya, yang merupakan batas yang disengaja, bukan keterbatasan nerve. Alasan ini berhasil adalah bahwa sebagian besar yang dibaca oleh attention suatu layer — riwayat KV — ditulis pada langkah-langkah sebelumnya dan tidak bergantung pada apa pun yang sedang berjalan, sehingga peluncuran baru dapat mulai memuatnya begitu tersedia, sebelum output layer sebelumnya ada. Layer pertama dan terakhir memiliki operasi yang sedikit berbeda, dan tulisan itu juga mencatatnya.

Dalam satu peluncuran, bentuknya tetap: satu blok thread berisi dua belas warp pada masing-masing 144 SM, diluncurkan dalam klaster berisi empat, menggunakan memori bersama terdistribusi ketika reduksi dengan latensi lebih rendah tersedia. Kedua belas warp tersebut tidak dapat dipertukarkan. Warp 0 memuat KV dan bobot serta menggunakan waktu senggangnya untuk melakukan prefetch bobot operasi berikutnya. Warp 1 dan 2 menjalankan aktivasi dan tensor core. Warp 3 hingga 11 membagi attention, GEMM padat, GEMM expert, serta penerimaan all-reduce dan all-gather. Pada layer sliding-window, warp 4 hingga 11 memuat KV mereka sendiri alih-alih menunggu 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.

Detail itu lebih penting daripada kelihatannya. Tim inferensi biasanya menerima bahwa kernel yang lebih cepat menghasilkan angka yang sedikit berbeda dan diam-diam menyetel ulang ambang evaluasi mereka. Kernel yang sama secara bitwise dengan paralelisasi di sisi pelatihan menghilangkan argumen itu sepenuhnya, yaitu sifat yang membuat sebuah optimisasi diadopsi, bukan sekadar dikagumi.

Tiga tagihan yang harus dibayar oleh kernel sebesar ini

Megakernel tidak gratis. Tulisan itu luar biasa terbuka soal overhead-nya, dan bagian itulah yang layak dibaca dua kali, karena bagian itu juga yang memberi tahu seberapa besar ruang yang tersisa.

Biaya pertama adalah cache instruksi. Sebuah SM menyimpan sekitar 128 KB instruksi dalam cache, dan kernel-kernel ini masing-masing berukuran 215 hingga 314 KB kode mesin, karena setiap kernel memuat setiap fase sebuah lapisan. Satu peluncuran hanya membawa 32 hingga 64 KB pertama mendekati SM; semua yang di luar jendela itu diambil baris demi baris saat pertama kali digunakan. Para penulis menjalankan eksperimen bersih tentang ini dan melaporkannya: dua blok 100 KB yang masing-masing membutuhkan 12,4 µs saat sendirian membutuhkan 19,5 µs ketika disambung berurutan, dan mengecilkan satu blok menjadi 20 KB mengembalikan blok lainnya ke kecepatan sendiriannya. Instruksi melewati cache L2 yang sama dengan data, sehingga kernel yang melampaui kapasitas cache instruksi akhirnya bersaing dengan pemuatan bobotnya sendiri.

Butir kedua adalah prefetch bobot, dan tulisan ini memperlakukannya sebagai masalah penempatan dengan tiga jawaban. Shared memory mendapat cincin 5 × 32 KB per SM, sekitar 23 MB di seluruh GPU. L2 adalah tempat sisanya, dan L2 GB300 menampung sekitar 135 MB, meskipun apa pun di L2 memerlukan perjalanan kedua yang lebih singkat ke shared memory dan L2 digunakan bersama dengan segala hal lain yang sedang berjalan. Keduanya bertumpuk: menjelang akhir sebuah layer, cincin tersebut sudah menyimpan bobot qkv pertama dari layer berikutnya. Di atas itu, warp 3 membaca expert mana yang dirutekan oleh langkah sebelumnya dan menarik irisan GPU ini dari 14 expert yang paling sering digunakan ke dalam L2, sekitar 50 MB per GPU, dengan bertaruh bahwa langkah-langkah berurutan merutekan secara serupa. Taruhan itu bernilai sekitar 0,2 ms dari langkah 5,3 ms, dan dapat memakan bandwidth saat tebakan itu salah.

Biaya ketiga adalah sinkronisasi, dan itulah yang tidak memiliki jawaban bersih. Ketika sebuah warp menunggu GPU lain melalui NVLink, tidak ada interupsi dan tidak ada wait-on-address di perangkat keras; warp melakukan polling pada sebuah flag hingga flag tersebut berubah. Microbenchmark dalam postingan menempatkan polling acquire dengan cakupan GPU pada 141 hingga 145 ns, flag satu arah antara dua GPU pada node yang sama sekitar 1,9 µs, dan polling yang mengantre di belakang penyalinan massal pada 1,0 hingga 2,6 µs. Fence memakan biaya 0,4 µs saat idle dan 3,6 hingga 4,0 µs dengan pemuatan TMA yang sedang berlangsung. Bahkan dengan pekerjaan yang terdistribusi sempurna, GPU dapat tertinggal satu sama lain hingga 4 µs, dan setiap titik henti dalam lapisan menunggu yang paling lambat. Megakernel menggantikan fence dengan nilai sentinel yang ditulis ke slot penyimpanan biasa — flag yang diisi data, +0/−0 untuk all-reduce, NaN semua-satu untuk pengumpulan bf16 — dan mereset salinan idle dari setiap slot saat salinan lainnya sedang digunakan.

Ini angka-angka yang membuat hasilnya terasa seperti rekayasa, bukan sihir: pergeseran GPU 4 µs terhadap langkah 5,9 ms memang kecil, tetapi empat titik henti sinkronisasi per lapisan yang dikalikan di 62 lapisan adalah tepat jenis overhead yang memakan keuntungan 2× jika Anda salah menanganinya.

Untuk apa 5,9 ms itu masih dihabiskan

Pos ini menyertakan kritiknya sendiri, dan kalimat paling berguna di dalamnya adalah tentang celah setelah GEMM kedua. Blok menyelesaikan GEMM itu di suatu titik antara 55 dan 64 µs sejak awal layer, persentil ke-10 hingga ke-90, tetapi combine menunggu keluaran setiap GPU, sehingga tidak dapat dimulai sampai 65 µs. Dari akhir GEMM 2 hingga akhir layer, sebuah blok tidak mengerjakan apa pun selama sekitar 12,5 µs secara rata-rata — hampir setengah dari total waktu menganggurnya. Rantai itulah, bukan overhead peluncuran kernel yang sudah dihilangkan oleh megakernel, yang menurut penulis sedang mereka periksa berikutnya.

Mereka juga menjelaskan mengapa atensi lapisan berikutnya tidak bisa begitu saja dimulai lebih awal. PDL meluncurkannya pada 79 µs, sekitar 4 µs sebelum blok pertama keluar, tetapi blok-bloknya baru tiba setelah blok-blok sebelumnya meninggalkan SM mereka pada 84 µs, jadi meluncurkan lebih awal tidak akan memulai apa pun lebih awal. Solusinya adalah melakukan prefetch baris KV lamanya ke L2 pada 65 µs, tempat baris-baris itu menunggu. Itu pola yang sama dengan prefetch bobot: ketika Anda tidak dapat menghilangkan suatu penantian, Anda memindahkan pekerjaan ke dalamnya.

Profiling adalah masalah struktural lainnya, dan ini masalah yang menarik: kernel tampak seperti satu batang panjang per layer bagi profiler eksternal mana pun, jadi kernel mengukur waktunya sendiri. Lane 0 dari setiap blok menulis stempel waktu ke array global melalui register globaltimer, dan para penulis mengukur biaya instrumentasi itu sekitar 2,5 µs per peluncuran layer, 2,8% dari satu langkah, dengan output dibiarkan identik bit dan sebuah flag build untuk menghilangkannya saat kompilasi. Kernel yang harus membawa profiler-nya sendiri adalah kernel yang telah menjadi sistem operasi, yang persis merupakan framing yang digunakan postingan itu: sebuah proses adalah satu kernel yang memegang semua sumber daya GPU, thread adalah warp tanpa preemption, menunggu I/O adalah busy-wait pada sebuah flag, DMA adalah TMA yang dilacak oleh mbarrier, pesan adalah penyimpanan biasa lintas NVLink.

Model drafnya juga jadi lebih cepat, dan itu sudah setengah kemenangan.

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.

Tim menerapkan perlakuan megakernel yang sama pada draf: lima lapisan draf masing-masing dijalankan sebagai satu peluncuran, ditambah tiga lainnya, yang membawa draf dari 175 kernel menjadi 8 dan dari 1,16 ms per langkah menjadi 440 µs, peningkatan 2,6× pada separuh langkah tersebut. Angka penerimaan layak disandingkan dengan perbandingan publik terdekat yang ditawarkan artikel itu.

• 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

• Mengapa 16 — enam belas token lebih mahal untuk diverifikasi karena mengaktifkan lebih banyak pakar, tetapi tim menemukan bahwa lebih mudah mencapai 6.6 dengan pohon yang lebih lebar daripada mencapai panjang yang sama dengan rantai yang lebih sempit

• Stack draf — 8000 baris CUDA dan PTX untuk model besar, dengan kode yang dipublikasikan di facebookresearch/sgl-muse-mega

Baca perbandingan ini dengan cermat, karena ini bukan tolok ukur antara dua model. Ini adalah perbandingan dua konfigurasi speculative decoding, satu dipublikasikan oleh Xiaomi dan satu oleh tim SGLang, pada perangkat keras dan beban kerja yang berbeda. Satu-satunya klaim yang didukungnya adalah bahwa pohon verifikasi yang lebih lebar merupakan jalur yang layak untuk mencapai panjang penerimaan yang sama, yang merupakan klaim desain, bukan peringkat.

Apa arti angka ini, dan apa yang bukan artinya.

Inilah bagian yang harus dilakukan dengan benar oleh sebuah rencana kapasitas, dan inilah bagian yang diakui secara jujur oleh tulisan itu sendiri. Angka 1,109 adalah median dari 20 langkah decode pada satu permintaan GSM8K, pada tensor parallel degree 8, di delapan GB300, dengan batch size satu dan speculative decoding aktif. Tulisan tersebut menyatakan dengan gamblang bahwa megakernel saat ini hanya mendukung batch size satu, dan bahwa memperluasnya ke batch yang lebih besar adalah pekerjaan berikutnya.

Batasan-batasan itu masing-masing menghilangkan klaim yang berbeda. Ukuran batch satu berarti ini bukan hasil throughput dalam arti yang digunakan tim serving: tidak ada keuntungan batching di dalamnya, tidak ada antrean, tidak ada campuran panjang prompt. Satu dataset berarti angka tersebut terikat pada distribusi prompt dan output GSM8K, serta pada pola routing terukur dari 49 expert per lapisan yang kemudian digunakan penulis sebagai routing sintetis dalam profiling mereka sendiri. Dua puluh langkah dari satu permintaan berarti variansnya dibatasi oleh jumlah langkah, bukan oleh jumlah permintaan.

Dan satu permintaan adalah kasus terburuk untuk overlap, yang dikatakan sendiri oleh para penulis: "pada satu permintaan setiap penantian terekspos; dengan lebih banyak permintaan, hampir selalu ada pekerjaan yang siap." Jadi 2.1× dari megakernel adalah pengukuran rezim ketika penyembunyian latensi paling sulit, dan tulisan itu secara eksplisit menyatakan bahwa batch yang lebih besar seharusnya lebih mudah, bukan lebih sulit. Itu adalah argumen yang mendukung hasil tersebut, dan tetap merupakan argumen, bukan pengukuran.

Jadi, yang dapat diketahui hari ini sempit dan nyata: kodenya ada di repositori publik, postingan tersebut menjelaskan mekanismenya dengan cukup rinci untuk diperiksa, aritmetika internalnya konsisten di tiga rasio independen, dan para penulis telah mendokumentasikan tiga overhead yang membatasi keuntungan lebih lanjut. Yang tidak dapat diketahui adalah apakah tim independen mendapatkan 1.109 pada delapan GB300 miliknya sendiri, apakah pendekatan tersebut bertahan pada batch nyata, dan apakah pendekatan itu masuk ke build SGLang yang dirilis — postingan tersebut tidak mengklaim telah melakukannya, dan repositori itu adalah artefak riset, bukan paket yang dapat Anda pasang. Perbandingan yang paling menentukan adalah dengan baseline yang telah disetel pada perangkat keras yang sama, karena "rantai default SGLang" adalah target yang bergerak dan tim yang menulis megakernel juga menyetel baseline yang dikalahkannya.

Bagaimana kaitannya dengan 169 token per detik yang dapat Anda sewa hari ini

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.

Selisih itu adalah hal paling berguna dalam kisah ini bagi siapa pun yang belum memiliki perangkat kerasnya. Jika sebuah endpoint menjawab pada 169 token per detik dan tumpukan delapan GPU rakitan sendiri menjawab pada 1.109, perbedaannya bukan alasan untuk membangun tumpukan itu. Delapan GB300 adalah pembelian modal yang diukur dalam jutaan, kernelnya adalah kode riset, dan batasan ukuran batch satu berarti tumpukan itu saat ini bagus tepat untuk beban kerja yang sudah ditangani dengan baik oleh endpoint terkelola: satu permintaan pada satu waktu. Tempat megakernel benar-benar mengubah perhitungan adalah bagi tim yang sudah membeli GPU untuk pelatihan atau untuk beban kerja dengan batasan residensi data, dan sekarang bertanya berapa utilisasi yang bisa mereka dapatkan darinya pada batch rendah. Bagi tim itu, 2,1× pada silikon yang sama dengan numerik yang sama adalah angka nyata dengan tagihan nyata yang menyertainya, dan tagihannya sudah dibayar.

Bagi semua orang lainnya, tafsiran jujurnya adalah bahwa tulisan ini merupakan sinyal tentang ke mana infrastruktur inferensi akan bergerak, bukan produk yang bisa Anda adopsi minggu ini. Arahnya konsisten dengan apa yang telah diterbitkan oleh lab yang sama sepanjang tahun: kernel yang lebih sedikit tetapi lebih besar, lebih banyak keputusan penjadwalan yang dipindahkan ke dalam kernel, dan GPU diperlakukan sebagai mesin yang harus diprogram, bukan sebagai sekumpulan pemanggilan pustaka yang harus diurutkan.

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.

Apa yang akan mengubah bagian ini?

Empat hal, dalam urutan menurun berdasarkan seberapa besar pengaruhnya terhadap gambaran.

• Batch yang lebih besar dari satu. Para penulis mengatakan batch yang lebih besar akan menjadi langkah berikutnya dan berargumen bahwa itu lebih mudah, karena dengan lebih banyak token yang sedang diproses, sebagian keluaran dari operasi sebelumnya sudah siap sementara sisanya masih tiba. Begitu angka yang di-batch ada, perbandingannya tidak lagi menjadi "satu aliran versus satu aliran" dan mulai menjadi klaim throughput yang dapat digunakan oleh tim serving.

• Reproduksi independen. Setiap angka di sini berasal dari harness milik para penulis sendiri, pada perangkat keras mereka sendiri, dengan routing sintetis mereka sendiri untuk proses profiling. Itu normal untuk pekerjaan kernel dan tetap merupakan satu sumber. Tim kedua yang menyamai 1,109, atau gagal menyamainya, adalah fakta yang akan memastikannya.

• Sebuah build SGLang yang telah dirilis yang memuatnya. Postingan tersebut menerbitkan kode dan menjelaskan desainnya; tidak disebutkan bahwa megakernel telah menggantikan jalur decode default untuk model ini dalam versi yang sudah dikirim. Sampai hal itu terjadi, mengadopsinya berarti mengadopsi repositori riset.

• Penutupan celah GEMM-2-ke-combine. Para penulis telah menyebut 12,5 µs waktu menganggur per blok di sana sebagai hampir setengah dari sisa yang ada. Jika rantai itu diperbaiki, perangkat keras yang sama menghasilkan waktu langkah yang secara signifikan lebih baik, dan 2,1× menjadi batas bawah dari angka yang lebih besar, bukan angka utamanya.

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.