Muse Spark Open: 1109 токенов в секунду — это изменение ядра, а не модели

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-х, чтобы объяснить её: ранние машины выполняли пакетные задания, причём одно задание владело всей машиной, а процессор простаивал, пока крутилась лента. Мультипрограммирование решило эту проблему, держа несколько заданий резидентными и переключаясь на другое всякий раз, когда текущее блокировалось. Мегаядро — это та же идея, перенесённая на GPU: одно ядро остаётся резидентным на каждом SM и выполняет множество заданий сразу, так что варпы могут начать работу следующей операции до того, как завершится предыдущая.

Измеренный эффект по рисунку к разделу о производительности — это сокращение шага декодирования с 12,9 мс до 5,9 мс. На этом рисунке шаг разбит на две половины — draft и verify: 1,16 мс и 10,9 мс при 1 450 ядрах в базовой версии, 0,44 мс и 5,3 мс при 97 ядрах в мегаядре. От 12,9 до 5,9 — это примерно 2,2×, а 1 109 / 517 — примерно 2,1×, так что два ключевых соотношения согласуются друг с другом. Есть и третья проверка, доступная любому, кто готов выполнить деление: шаг проверяет 16 токенов и принимает 6,6, то есть 6,6 токена за 5,9 мс дают примерно 1 120 токенов в секунду. Арифметика внутренне согласована — это первое, что нужно установить насчёт такого числа, и далеко не последнее.

Чего цифры не говорят, так это того, какова скорость в расчёте на один GPU. В посте приводится совокупная пропускная способность для всего развёртывания из восьми GPU. Если разделить на восемь, получится около 139 токенов в секунду на GPU — это мой расчёт на основе их цифры, а не число, которое публикует пост, и его нельзя сравнивать с показателем на один GPU, который вам могут назвать для другой модели на другой машине.

Где проходит граница ядра и почему она заканчивается именно там

Мегаядро не поглощает всю модель. Каждый запуск охватывает путь от внимания одного слоя декодера до qkv-проекции и записи KV следующего слоя — это намеренная граница, а не ограничение nerve. Работает это потому, что большая часть того, что читает внимание слоя, — история KV — была записана на более ранних шагах и не зависит ни от чего из ещё выполняющегося, поэтому новый запуск может начать загружать её сразу, как только она становится доступной, ещё до того, как появится выход предыдущего слоя. У первого и последнего слоёв операции немного отличаются, и в посте это тоже отмечается.

В рамках одного запуска конфигурация фиксирована: один блок потоков из двенадцати варпов на каждом из 144 SM, запускаемый кластерами по четыре, с использованием распределённой разделяемой памяти там, где доступна редукция с меньшей задержкой. Двенадцать варпов не взаимозаменяемы. Варп 0 загружает KV и веса и тратит своё свободное время на предварительную выборку весов следующей операции. Варпы 1 и 2 выполняют активации и работу на тензорных ядрах. Варпы с 3 по 11 делят между собой внимание, плотные GEMM, экспертные GEMM, а также приём данных all-reduce и all-gather. В слоях со скользящим окном варпы с 4 по 11 загружают собственный KV, а не ждут варп 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 КБ инструкций, а эти ядра содержат от 215 до 314 КБ машинного кода каждое, потому что каждое включает все фазы слоя. При запуске только первые 32–64 КБ оказываются близко к SM; всё за пределами этого окна подгружается построчно при первом использовании. Авторы провели на эту тему чистый эксперимент и сообщили о нём: два блока по 100 КБ, каждый из которых по отдельности выполнялся за 12,4 мкс, при последовательном соединении вместе заняли 19,5 мкс, а уменьшение одного блока до 20 КБ вернуло другой к его одиночной скорости. Инструкции проходят через тот же кэш L2, что и данные, поэтому ядро, которое перерастает кэш инструкций, в итоге конкурирует со своими собственными загрузками весов.

Второй счёт — предварительная загрузка весов, и в посте это рассматривается как задача размещения с тремя решениями. Разделяемая память получает кольцо из 5 × 32 КБ на SM, около 23 МБ на всю GPU. L2 — это куда уходит остальное, и L2 у GB300 вмещает около 135 МБ, хотя всё в L2 требует второго, более короткого прохода в разделяемую память, и L2 делится со всем остальным, что находится в работе. Эти два подхода складываются: ближе к концу слоя кольцо уже содержит первые qkv-веса следующего слоя. Вдобавок warp 3 считывает, к каким экспертам маршрутизировал предыдущий шаг, и подтягивает в L2 принадлежащие этой GPU срезы 14 самых используемых из них, около 50 МБ на GPU, делая ставку на то, что последовательные шаги маршрутизируются одинаково. Эта ставка стоит около 0,2 мс из 5,3 мс шага и может обойтись в пропускную способность, когда она неверна.

Третий счёт — синхронизация, и именно здесь нет однозначного ответа. Когда warp ждёт другую GPU через NVLink, в оборудовании нет ни прерывания, ни wait-on-address; warp опрашивает флаг, пока тот не изменится. Микробенчмарки из поста показывают: acquire-опрос в масштабе GPU занимает 141–145 нс, односторонний флаг между двумя GPU на одном узле — около 1,9 мкс, а опрос, стоящий в очереди за массовыми копиями, — 1,0–2,6 мкс. Fence-операции стоят 0,4 мкс в простое и 3,6–4,0 мкс при выполняющихся TMA-загрузках. Даже при идеально распределённой работе GPU могут отставать друг от друга до 4 мкс, а каждая остановка в слое ждёт самую медленную. Мегаядро заменяет fence-операции sentinel-значениями, записываемыми в обычные слоты записи — флаг, заполненный данными, +0/−0 для all-reduce, NaN из всех единиц для bf16-gather — и сбрасывает простаивающую копию каждого слота, пока используется другая.

Вот числа, которые заставляют результат ощущаться инженерией, а не магией: 4 µs рассинхронизации GPU на фоне шага в 5.9 ms — это мало, но четыре остановки синхронизации на слой, умноженные на 62 слоя, — это именно тот вид накладных расходов, который съедает двукратный выигрыш, если ошибиться.

На что всё ещё тратятся эти 5,9 мс

В посте есть и самокритика, и самая полезная фраза в нём — о разрыве после второго GEMM. Блоки завершают этот GEMM в интервале от 55 до 64 мкс с начала слоя (10–90-й процентиль), но операция combine ждёт вывода от каждой GPU, поэтому она не может начаться раньше 65 мкс. С конца GEMM 2 до конца слоя блок в среднем простаивает без работы около 12,5 мкс — почти половину всего своего времени простоя. Именно эта цепочка, а не накладные расходы на запуск ядра, которые мегаядро уже устранило, по словам авторов, находится в центре их дальнейшего внимания.

Они также объясняют, почему внимание следующего слоя не может просто начаться раньше. PDL запускает его на 79 мкс, примерно за 4 мкс до завершения первого блока, но его блоки размещаются только после того, как предыдущие покидают свои SM на 84 мкс, поэтому более ранний запуск ничего не запустил бы раньше. Обходной путь — предварительно загрузить его старые строки KV в L2 на 65 мкс, где они и ждут. Это тот же шаблон, что и при предзагрузке весов: когда нельзя устранить ожидание, вы переносите работу внутрь него.

Профилирование — другая структурная проблема, и притом приятная: для любого внешнего профилировщика ядро выглядит как одна длинная полоса на слой, поэтому оно замеряет время само. Lane 0 каждого блока записывает метку времени в глобальный массив через регистр globaltimer, и авторы оценивают стоимость такой инструментации примерно в 2,5 мкс на запуск слоя — 2,8% шага; выходные данные при этом остаются побитово идентичными, а флаг сборки позволяет исключить её при компиляции. Ядро, которому приходится нести собственный профилировщик, — это ядро, ставшее операционной системой; именно такую трактовку предлагает пост: процесс — это одно ядро, удерживающее все ресурсы GPU, потоки — это варпы без вытеснения, ожидание ввода-вывода — это активное ожидание флага, DMA — это TMA, отслеживаемый через mbarrier, сообщения — это обычные записи через 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.

Команда применила ту же обработку мегаядром к черновику: пять слоёв черновика выполняются каждый за один запуск, плюс ещё три, что сокращает черновик с 175 ядер до 8 и с 1,16 мс на шаг до 440 мкс — улучшение в 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 — проверка шестнадцати токенов обходится дороже, потому что они активируют больше экспертов, но команде показалось легче достичь 6.6 с более широким деревом, чем достичь той же длины с более узкой цепочкой

• Черновой стек — 8000 строк CUDA и PTX для большой модели, с кодом, опубликованным в facebookresearch/sgl-muse-mega

Внимательно изучите это сравнение, потому что это не бенчмарк между двумя моделями. Это сравнение двух конфигураций спекулятивного декодирования: одна опубликована Xiaomi, другая командой SGLang, на разном оборудовании и с разными рабочими нагрузками. Единственное утверждение, которое оно поддерживает, состоит в том, что более широкое дерево верификации — жизнеспособный путь к той же длине принятия, и это утверждение о дизайне, а не ранжирование.

Что такое это число и чем оно не является

Вот та часть, которую план по мощностям должен учесть правильно, и именно в ней сам пост честен. Цифра 1,109 — это медиана по 20 шагам декодирования на одном запросе GSM8K, при степени тензорного параллелизма 8, на восьми GB300, с размером батча один и включённым спекулятивным декодированием. В посте прямо сказано, что сегодня мегаядро поддерживает только размер батча один, и что расширение его на большие батчи — следующая задача.

Каждое из этих ограничений снимает своё отдельное утверждение. Размер батча, равный единице, означает, что это не результат по пропускной способности в том смысле, в каком это слово использует команда, занимающаяся обслуживанием: в нём нет выигрыша от батчинга, нет очереди, нет смешения длин промптов. Один набор данных означает, что число привязано к распределению промптов и выходных данных GSM8K, а также к измеренному шаблону маршрутизации из 49 экспертов на слой, который авторы затем использовали как синтетическую маршрутизацию в своём собственном профилировании. Двадцать шагов одного запроса означают, что разброс ограничен количеством шагов, а не количеством запросов.

И одиночный запрос — худший случай для перекрытия, о чём сами авторы говорят: «при одном запросе каждое ожидание остаётся не скрытым; при большем числе почти всегда есть готовая работа». Так что 2,1× мегаядра — это измерение режима, в котором скрывать задержки труднее всего, и в посте прямо сказано, что с более крупными пакетами должно быть легче, а не труднее. Это аргумент в пользу результата, и всё же это аргумент, а не измерение.

То, что сегодня можно знать, — узко и реально: код существует в публичном репозитории, в посте механизм описан достаточно подробно для проверки, внутренняя арифметика согласуется по трём независимым соотношениям, а авторы задокументировали три накладных расхода, которые ограничивают дальнейший выигрыш. Чего знать нельзя — это получит ли независимая команда 1,109 на своих восьми GB300s, выживет ли подход при реальном батче и попадёт ли он в выпущенную сборку SGLang; в посте не утверждается, что это уже произошло, а репозиторий — это исследовательский артефакт, а не пакет, который можно установить. Сравнение, которое прояснило бы больше всего, — это сравнение с настроенным baseline на том же оборудовании, поскольку «стандартная цепочка SGLang» — движущаяся мишень, а команда, написавшая мегаядро, также настраивает baseline, который оно превосходит.

Как это соотносится со 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 токенов в секунду, а собранный вручную стек из восьми GPU — 1 109, эта разница не является поводом собирать такой стек. Восемь GB300 — это капитальная закупка, измеряемая миллионами, ядро — это исследовательский код, а ограничение batch-size-one означает, что сейчас стек хорошо подходит ровно для той нагрузки, с которой уже хорошо справляется размещённый эндпоинт: один запрос за раз. Мегаядро по-настоящему меняет расклад для команды, которая уже купила GPU для обучения или для нагрузки с ограничениями на хранение данных, и теперь спрашивает, какую утилизацию она может из них получить при малом batch. Для такой команды 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.

Что изменило бы эту часть?

Четыре вещи, в порядке убывания того, насколько сильно они могли бы изменить картину.

• Батч больше единицы. Авторы говорят, что дальше будут более крупные батчи, и утверждают, что так проще, потому что при большем числе токенов в обработке часть выходов предыдущей операции уже готова, пока остальные ещё поступают. Как только появляется число, полученное при батчинге, сравнение перестаёт быть сравнением «одного потока с одним потоком» и становится заявлением о пропускной способности, которое может использовать команда, отвечающая за обслуживание.

• Независимое воспроизведение. Все приведённые здесь цифры получены на собственном стенде авторов, на их собственном оборудовании, с их собственной синтетической маршрутизацией для прохода профилирования. Для работы с ядрами это нормально, но всё равно остаётся единственным источником. Вторая команда, которая воспроизведёт 1,109 — или не сможет этого сделать, — вот факт, который бы всё решил.

• Выпущенная сборка SGLang, которая его содержит. В публикации выложен код и описана архитектура; в ней не утверждается, что мегаядро заменило стандартный путь декодирования для этой модели в поставляемой версии. Пока это не так, его внедрение означает внедрение исследовательского репозитория.

• Закрытие разрыва между GEMM-2 и combine. Авторы уже назвали тамошние 12,5 мкс простоя на блок почти половиной того, что остаётся. Если эту цепочку исправить, то же оборудование даст заметно лучшее время шага, а 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.