Muse Spark Open: 1109 tokenów na sekundę to zmiana jądra, a nie zmiana modelu

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.

Punkt odniesienia, którego nikt nie benchmarkuje: 1450 uruchomień na 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.

To jest koszt, w który uderza megakernel, a autorzy sięgają po analogię z lat 60., aby go wyjaśnić: wczesne maszyny wykonywały zadania wsadowe, przy czym jedno zadanie zajmowało całą maszynę, a procesor pozostawał bezczynny, podczas gdy taśma się obracała. Wieloprogramowość rozwiązała ten problem, utrzymując w pamięci kilka zadań i przełączając się na inne, gdy bieżące zostało zablokowane. Megakernel to ten pomysł przeniesiony na GPU — jedno jądro pozostaje rezydentne na każdym SM i uruchamia wiele zadań naraz, dzięki czemu warpy mogą rozpocząć pracę nad następną operacją, zanim zakończy się poprzednia.

Zmierzony efekt, widoczny na rysunku towarzyszącym sekcji dotyczącej wydajności, to krok dekodowania skracający się z 12,9 ms do 5,9 ms. Rysunek ten dzieli krok na część draft i część verify: 1,16 ms i 10,9 ms przy 1450 jądrach w wersji bazowej oraz 0,44 ms i 5,3 ms przy 97 jądrach w megakernelu. Przejście z 12,9 do 5,9 to około 2,2×, a 1109 przez 517 to około 2,1×, więc dwa główne stosunki są ze sobą zgodne. Jest jeszcze trzecia kontrola dostępna dla każdego, kto zechce wykonać dzielenie: krok weryfikuje 16 tokenów i akceptuje 6,6, więc 6,6 tokena w 5,9 ms daje w przybliżeniu 1120 tokenów na sekundę. Arytmetyka jest wewnętrznie spójna, co jest pierwszą rzeczą, którą należy ustalić w przypadku takiej liczby, ale nie ostatnią.

Jednej rzeczy liczby nie mówią: jaka prędkość przypada na poszczególne GPU. Wpis podaje łączną przepustowość dla całego wdrożenia z ośmioma GPU. Podzielona przez osiem daje około 139 tokenów na sekundę na GPU — to moje obliczenie na podstawie ich liczby, a nie wartość publikowana we wpisie, i nie można jej porównywać z liczbą na GPU, jaką można spotkać w odniesieniu do innego modelu na innej maszynie.

Gdzie przebiega granica jądra i dlaczego właśnie tam się kończy

Megakernel nie połyka całego modelu. Każde uruchomienie obejmuje uwagę jednej warstwy dekodera aż po projekcję qkv i zapis KV następnej warstwy, co jest celową granicą, a nie ograniczeniem nerve. Działa to dlatego, że większość tego, co czyta uwaga warstwy — historia KV — została zapisana na wcześniejszych krokach i nie zależy od niczego, co jest w toku, więc nowe uruchomienie może zacząć ją ładować, gdy tylko będzie dostępna, jeszcze zanim powstanie wynik poprzedniej warstwy. Pierwsza i ostatnia warstwa mają nieco inne operacje, co również odnotowuje post.

W ramach jednego uruchomienia struktura jest stała: jeden blok wątków złożony z dwunastu warpow na każdym ze 144 SM, uruchamiany w klastrach po cztery, z wykorzystaniem rozproszonej pamięci współdzielonej tam, gdzie dostępna jest redukcja o niższym opóźnieniu. Dwanaście warpow nie jest wzajemnie wymiennych. Warp 0 wczytuje KV i wagi, a w wolnym czasie pobiera z wyprzedzeniem wagi następnej operacji. Warpy 1 i 2 obsługują aktywacje i rdzenie tensorowe. Warpy od 3 do 11 dzielą między siebie mechanizm attention, gęste GEMM-y, GEMM ekspertów oraz odbiory all-reduce i all-gather. W warstwach z przesuwnym oknem warpy od 4 do 11 wczytują własne KV, zamiast czekać na 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.

Ten szczegół ma większe znaczenie, niż się wydaje. Zespoły zajmujące się wnioskowaniem zwykle godzą się z tym, że szybszy kernel daje nieco inne liczby, i po cichu dostrajają swoje progi ewaluacji. Kernel bitowo zgodny z paralelizacją po stronie treningu całkowicie usuwa ten argument — a to właśnie tego rodzaju właściwość sprawia, że optymalizacja zostaje przyjęta, a nie tylko podziwiana.

Trzy rachunki, które musi zapłacić ziarno tej wielkości.

Megakernel nie jest darmowy. Ten wpis jest niezwykle szczery w kwestii narzutów i to właśnie tę część warto przeczytać dwa razy, ponieważ jest ona również tą częścią, która mówi, ile jeszcze zostało miejsca.

Pierwszy rachunek dotyczy pamięci podręcznej instrukcji. SM buforuje około 128 KB instrukcji, a te jądra mają po 215–314 KB kodu maszynowego, ponieważ każde przechowuje wszystkie fazy warstwy. Uruchomienie przybliża do SM tylko pierwsze 32–64 KB; wszystko poza tym oknem jest pobierane linia po linii przy pierwszym użyciu. Autorzy przeprowadzili na ten temat czysty eksperyment i opisali go: dwa bloki po 100 KB, z których każdy osobno zajmował 12,4 µs, zajmowały 19,5 µs, gdy połączono je bezpośrednio jeden po drugim, a zmniejszenie jednego bloku do 20 KB przywróciło drugiemu prędkość osiąganą osobno. Instrukcje przechodzą przez tę samą pamięć podręczną L2 co dane, więc jądro, które przekracza rozmiar pamięci podręcznej instrukcji, zaczyna konkurować z własnym wczytywaniem wag.

Drugi koszt to prefetch wag, a wpis ujmuje to jako problem rozmieszczenia z trzema odpowiedziami. Pamięć współdzielona dostaje pierścień 5 × 32 KB na SM, około 23 MB w całym GPU. Reszta trafia do L2, a L2 w GB300 mieści około 135 MB, choć wszystko w L2 wymaga drugiej, krótszej podróży do pamięci współdzielonej, a L2 jest dzielona ze wszystkim innym, co jest w locie. Te dwa rozwiązania się nakładają: pod koniec warstwy pierścień trzyma już pierwsze wagi qkv następnej warstwy. Do tego warp 3 odczytuje, do których ekspertów rozesłano poprzedni krok, i wciąga do L2 należące do tego GPU fragmenty 14 najczęściej używanych z nich — około 50 MB na GPU — zakładając, że kolejne kroki rozsyłają podobnie. Ten zakład jest wart około 0,2 ms z 5,3 ms kroku, a gdy się nie sprawdzi, może kosztować przepustowość.

Trzecią pozycją kosztową jest synchronizacja i to ta, na którą nie ma czystej odpowiedzi. Gdy warp czeka na inny GPU przez NVLink, nie ma przerwania ani sprzętowego oczekiwania na adres; warp odpytuje flagę, aż ta się zmieni. Mikrobenchmarki z wpisu wskazują na odpytywanie acquire w zakresie GPU na poziomie 141–145 ns, jednokierunkową flagę między dwoma GPU w tym samym węźle na poziomie około 1,9 µs oraz odpytywanie zakolejkowane za kopiami zbiorczymi na poziomie 1,0–2,6 µs. Bariery kosztują 0,4 µs w stanie bezczynności i 3,6–4,0 µs, gdy w locie są ładowania TMA. Nawet przy idealnie rozdzielonej pracy GPU mogą pozostawać względem siebie opóźnione o do 4 µs, a każdy punkt zatrzymania w warstwie czeka na najwolniejszy z nich. Megakernel zastępuje bariery wartościami sentinel zapisywanymi do zwykłych slotów zapisu — flagę wypełnioną danymi, +0/−0 dla all-reduce, NaN z samych jedynek dla gatherów bf16 — i resetuje bezczynną kopię każdego slotu, gdy druga jest w użyciu.

To są liczby, które sprawiają, że wynik bardziej przypomina inżynierię niż magię: 4 µs rozbieżności GPU wobec kroku 5,9 ms to niewiele, ale cztery zatrzymania synchronizacji na warstwę pomnożone przez 62 warstwy to dokładnie ten rodzaj narzutu, który pochłania zysk 2×, jeśli zrobisz to źle.

Na co wciąż przeznaczane jest 5,9 ms

Wpis zawiera własną krytykę, a najbardziej przydatne zdanie w nim dotyczy luki po drugim GEMM. Bloki kończą ten GEMM w dowolnym momencie między 55 a 64 µs od początku warstwy, od 10. do 90. percentyla, ale operacja combine czeka na wynik każdego GPU, więc nie może się rozpocząć przed 65 µs. Od końca GEMM 2 do końca warstwy blok pozostaje bez pracy średnio przez około 12,5 µs — prawie połowę całkowitego czasu bezczynności. To właśnie ten łańcuch, a nie narzut związany z uruchamianiem jąder, który megakernel już wyeliminował, jest tym, czemu — jak twierdzą autorzy — przyglądają się teraz.

Wyjaśniają też, dlaczego attention kolejnej warstwy nie może po prostu rozpocząć się wcześniej. PDL uruchamia je w 79 µs, około 4 µs przed zakończeniem pierwszego bloku, ale jego bloki trafiają na miejsce dopiero, gdy poprzednie opuszczą swoje SM w 84 µs, więc wcześniejsze uruchomienie nie rozpoczęłoby niczego wcześniej. Obejściem jest wstępne pobranie jego starych wierszy KV do L2 w 65 µs, gdzie czekają. To ten sam wzorzec co wstępne pobieranie wag: gdy nie można usunąć oczekiwania, przenosi się do niego pracę.

Profilowanie to drugi problem strukturalny — i to dobry problem: dla każdego zewnętrznego profilera jądro wygląda jak jeden długi pasek na warstwę, więc mierzy się samo. Lane 0 każdego bloku zapisuje znacznik czasu do globalnej tablicy przez rejestr globaltimer, a autorzy szacują koszt tej instrumentacji na około 2,5 µs na uruchomienie warstwy, czyli 2,8% kroku, przy czym wyniki pozostają bitowo identyczne, a flaga kompilacji pozwala ją wyłączyć. Jądro, które musi nosić własny profiler, to jądro, które stało się systemem operacyjnym — dokładnie tak ujmuje to wpis: proces to jedno jądro obejmujące wszystkie zasoby GPU, wątki to warpy bez wywłaszczania, oczekiwanie na I/O to aktywne oczekiwanie na flagę, DMA to TMA śledzone przez mbarrier, a komunikaty to zwykłe zapisy przez NVLink.

Model roboczy również przyspieszył, a to już połowa sukcesu.

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.

Zespół zastosował ten sam zabieg z megakernel do draftu: pięć warstw draftu działa jako jedno uruchomienie, plus trzy inne, co zmniejsza liczbę jąder draftu ze 175 do 8, a czas na krok z 1,16 ms do 440 µs — 2,6× poprawy w tej połowie kroku. Liczby dotyczące akceptacji warto zestawić z najbliższym publicznym porównaniem, jakie oferuje post.

• 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

• Dlaczego 16 — szesnaście tokenów kosztuje więcej przy weryfikacji, ponieważ aktywują więcej ekspertów, ale zespół uznał, że łatwiej jest osiągnąć 6,6 za pomocą szerszego drzewa niż osiągnąć tę samą długość za pomocą węższego łańcucha

• Stos wersji roboczej — 8000 linii CUDA i PTX dla dużego modelu, z kodem opublikowanym pod adresem facebookresearch/sgl-muse-mega

Przeczytaj to porównanie uważnie, ponieważ nie jest to benchmark między dwoma modelami. To porównanie dwóch konfiguracji dekodowania spekulatywnego — jednej opublikowanej przez Xiaomi, a drugiej przez zespół SGLang — na różnym sprzęcie i różnych obciążeniach. Jedyna teza, którą ono wspiera, to że szersze drzewo weryfikacji jest realną drogą do uzyskania tej samej długości akceptacji, co jest tezą projektową, a nie rankingiem.

Czym jest ta liczba, a czym nie jest

Oto element, który plan przepustowości musi właściwie ująć — i co do którego sam post jest uczciwy. Liczba 1,109 to mediana z 20 kroków dekodowania pojedynczego żądania GSM8K, przy stopniu równoległości tensorowej 8, na ośmiu GB300, z rozmiarem batcha równym jeden i włączonym dekodowaniem spekulacyjnym. Post wprost stwierdza, że megajądro obsługuje dziś wyłącznie batch size równy jeden, a rozszerzenie go na większe batche to następny etap prac.

Każde z tych ograniczeń eliminuje inne twierdzenie. Rozmiar partii równy jeden oznacza, że nie jest to wynik przepustowości w rozumieniu, jakiego używa zespół obsługujący modele: nie ma w nim zysku z batchowania, nie ma kolejki, nie ma mieszanki długości promptów. Jeden zbiór danych oznacza, że liczba jest powiązana z rozkładem promptów i wyników GSM8K oraz ze zmierzonym wzorcem routingu 49 ekspertów na warstwę, którego autorzy następnie użyli jako syntetycznego routingu we własnym profilowaniu. Dwadzieścia kroków jednego żądania oznacza, że wariancja jest ograniczona przez liczbę kroków, a nie przez liczbę żądań.

A pojedyncze żądanie to najgorszy przypadek dla nakładania się, co sami autorzy mówią: „przy jednym żądaniu każde oczekiwanie jest odsłonięte; przy większej liczbie prawie zawsze jest gotowa praca”. Zatem 2,1× megajądra to pomiar reżimu, w którym ukrywanie opóźnień jest najtrudniejsze, a wpis wyraźnie stwierdza, że większe partie powinny być łatwiejsze, a nie trudniejsze. To argument na korzyść wyniku, ale wciąż argument, a nie pomiar.

To, co jest dziś poznawalne, jest zatem wąskie i realne: kod istnieje w publicznym repozytorium, post opisuje mechanizm wystarczająco szczegółowo, by można go było zweryfikować, wewnętrzna arytmetyka jest spójna w trzech niezależnych stosunkach, a autorzy udokumentowali trzy narzuty, które ograniczają dalsze zyski. Nie można natomiast wiedzieć, czy niezależny zespół uzyska 1,109 na własnych ośmiu GB300s, czy podejście przetrwa rzeczywistą partię wsadową i czy trafi do wydanego buildu SGLang — post nie twierdzi, że tak się stało, a repozytorium jest artefaktem badawczym, a nie pakietem, który można zainstalować. Porównaniem, które rozstrzygnęłoby najwięcej, jest zestawienie z dostrojoną linią bazową na tym samym sprzęcie, ponieważ „domyślny łańcuch SGLang” jest ruchomym celem, a zespół, który napisał megakernel, dostraja również linię bazową, którą przewyższa.

Jak to się ma do 169 tokenów na sekundę, które można dziś wynająć

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.

Ta luka jest najbardziej użyteczną rzeczą w tej historii dla każdego, kto jeszcze nie ma tego sprzętu. Jeśli endpoint odpowiada z szybkością 169 tokenów na sekundę, a ręcznie zbudowany stos ośmiu GPU odpowiada z szybkością 1 109, ta różnica nie jest powodem, by budować ten stos. Osiem GB300 to zakup kapitałowy liczony w milionach, jądro to kod badawczy, a ograniczenie rozmiaru partii do jednego oznacza, że stos jest obecnie dobry dokładnie w obciążeniu, które hostowany endpoint już dobrze obsługuje: jedno żądanie naraz. Megakernel naprawdę zmienia kalkulację dla zespołu, który już kupił GPU do trenowania albo dla obciążenia z ograniczeniami dotyczącymi rezydencji danych, i który teraz pyta, jakie wykorzystanie może z nich uzyskać przy niskim batchu. Dla tego zespołu 2,1× na tym samym krzemie z tą samą numeryką to realna liczba z realną fakturą, a faktura została już opłacona.

Dla wszystkich innych uczciwa interpretacja jest taka, że ten wpis jest sygnałem, dokąd zmierza infrastruktura wnioskowania, a nie produktem, który można wdrożyć w tym tygodniu. Kierunek jest spójny z tym, co to samo laboratorium publikuje przez cały rok: mniej, ale większych jąder, więcej decyzji dotyczących szeregowania przenoszonych do wnętrza jądra oraz GPU traktowany jako maszyna, którą się programuje, a nie zestaw wywołań bibliotecznych, które trzeba szeregować.

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.

Co zmieniłoby ten element?

Cztery rzeczy, w malejącym porządku według tego, jak bardzo zmieniłyby obraz sytuacji.

• Partia większa niż jeden. Autorzy twierdzą, że następne będą większe partie, i argumentują, że są łatwiejsze, ponieważ gdy w locie jest więcej tokenów, część wyników poprzedniej operacji jest już gotowa, podczas gdy pozostałe wciąż napływają. Gdy tylko pojawia się liczba dla partii, porównanie przestaje być zestawieniem „jednego strumienia z jednym strumieniem”, a staje się twierdzeniem o przepustowości, z którego zespół serwujący może skorzystać.

• Niezależna reprodukcja. Każda liczba pochodzi z własnego stanowiska testowego autorów, na ich własnym sprzęcie, z ich własnym syntetycznym routingu na potrzeby przebiegu profilowania. To normalne w pracy nad kernelami i wciąż pozostaje jednym źródłem. Drugi zespół, który uzyskałby te same 1 109 — albo któremu by się to nie udało — to fakt, który by to rozstrzygnął.

• Wydana kompilacja SGLang, która to zawiera. Wpis publikuje kod i opisuje projekt; nie stwierdza, że megakernel zastąpił domyślną ścieżkę dekodowania dla tego modelu w wydanej wersji. Dopóki tak się nie stanie, przyjęcie go oznacza przyjęcie repozytorium badawczego.

• Zamykanie luki GEMM-2-to-combine. Autorzy wskazali już 12,5 µs czasu bezczynności na blok jako prawie połowę tego, co pozostaje. Jeśli ten łańcuch zostanie naprawiony, ten sam sprzęt zapewni istotnie lepszy czas kroku, a 2,1× stanie się dolną granicą większej wartości, a nie głównym wynikiem.

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.