Muse Spark Open: 초당 1,109 토큰은 모델 변경이 아니라 커널 변경이다

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.9ms에서 5.9ms로 줄어드는 것입니다. 그 그림은 이 단계를 드래프트와 검증의 두 부분으로 나누는데, 베이스라인에서는 1,450개 커널로 1.16ms와 10.9ms이고, 메가커널에서는 97개 커널로 0.44ms와 5.3ms입니다. 12.9에서 5.9는 약 2.2배이고, 1,109를 517로 나눈 값은 약 2.1배이므로, 두 대표 비율은 서로 일치합니다. 나눗셈을 할 의향이 있는 사람이라면 누구나 이용할 수 있는 세 번째 확인이 있습니다. 이 단계는 16개 토큰을 검증하고 6.6개를 채택하므로, 5.9ms 동안 6.6개 토큰은 대략 초당 1,120개 토큰이 됩니다. 이 계산은 내부적으로 일관성이 있으며, 이는 이런 수치에서 가장 먼저 확인해야 할 것이지 마지막으로 확인할 것은 아닙니다.

숫자만으로는 알 수 없는 한 가지는 GPU당 속도가 어느 정도인지입니다. 게시글은 8개 GPU 전체 배포의 총 처리량을 보고합니다. 이를 8로 나누면 GPU당 초당 약 139토큰입니다. 이는 제가 그들의 수치로 계산한 것이지 게시글이 공개한 숫자가 아니며, 다른 모델을 다른 머신에서 인용할 때 볼 수 있는 GPU당 숫자와 비교할 수 있는 것도 아닙니다.

커널 경계가 어디에 위치하며, 왜 거기서 멈추는가

메가커널은 모델 전체를 삼키지 않는다. 각 실행은 한 디코더 레이어의 어텐션부터 다음 레이어의 qkv 프로젝션 및 KV 쓰기까지를 포괄하며, 이는 nerve의 한계라기보다 의도된 경계다. 이것이 작동하는 이유는 한 레이어의 어텐션이 읽는 것 대부분, 즉 KV 히스토리가 이전 단계에서 기록되었고 진행 중인 어떤 것에도 의존하지 않기 때문이다. 따라서 새 실행은 그것이 사용 가능해지는 즉시, 이전 레이어의 출력이 존재하기 전에 로드하기 시작할 수 있다. 첫 번째와 마지막 레이어는 연산이 약간 다르며, 글에서도 그 점을 언급한다.

한 번의 런치 안에서 구성은 고정된다. 144개 SM 각각에 12개 워프로 이루어진 스레드 블록 하나가 있고, 4개씩 클러스터로 런치되며, 더 낮은 지연의 리덕션이 가능한 분산 공유 메모리를 사용한다. 12개 워프는 서로 바꿔 쓸 수 없다. 워프 0은 KV와 가중치를 로드하고, 남는 시간을 다음 연산의 가중치 프리페치에 쓴다. 워프 1과 2는 활성화와 텐서 코어를 실행한다. 워프 3부터 11까지는 어텐션, dense GEMM, expert GEMM, 그리고 all-reduce 및 all-gather 수신을 나눠 맡는다. 슬라이딩 윈도 레이어에서는 워프 4부터 11까지가 워프 0을 기다리지 않고 자체 KV를 로드한다.

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은 약 128KB의 명령어를 캐시하는데, 이 커널들은 각각 215~314KB의 기계어 코드입니다. 각 커널이 한 레이어의 모든 단계를 담고 있기 때문입니다. 커널 실행 시 SM 가까이로 가져오는 것은 처음 32~64KB뿐이며, 그 범위를 넘어선 모든 것은 처음 사용될 때 라인 단위로 페치됩니다. 저자들은 이에 대해 깔끔한 실험을 수행해 보고했습니다. 각각 단독으로는 12.4µs가 걸리던 두 개의 100KB 블록을 연달아 이어 붙이면 19.5µs가 걸렸고, 한 블록을 20KB로 줄이자 다른 블록은 단독 속도로 돌아왔습니다. 명령어는 데이터와 동일한 L2 캐시를 통해 이동하므로, 명령어 캐시 용량을 넘어서는 커널은 결국 자체 가중치 로드와 경쟁하게 됩니다.

두 번째 항목은 가중치 프리페치이며, 이 글은 이를 세 가지 해법이 있는 배치 문제로 다룹니다. 공유 메모리는 SM당 5 × 32KB의 링을 가지며, GPU 전체로는 약 23MB입니다. 나머지는 L2로 가는데, GB300의 L2는 약 135MB를 담을 수 있습니다. 다만 L2에 있는 것은 무엇이든 공유 메모리로 들어가는 두 번째의 더 짧은 이동이 필요하고, L2는 실행 중인 다른 모든 것과 공유됩니다. 두 방식은 겹쳐서 동작합니다. 한 레이어가 끝날 무렵이면 링에는 이미 다음 레이어의 첫 qkv 가중치가 들어 있습니다. 그 위에 워프 3은 이전 단계에서 어떤 전문가로 라우팅되었는지 읽고, 가장 많이 쓰인 14개 중 이 GPU에 해당하는 조각을 L2로 끌어옵니다. GPU당 약 50MB이며, 연속된 단계가 비슷하게 라우팅된다는 데 베팅하는 것입니다. 이 베팅은 5.3ms 단계 중 약 0.2ms의 가치가 있지만, 틀릴 경우 대역폭을 잡아먹을 수 있습니다.

세 번째 부담은 동기화이며, 여기에는 깔끔한 답이 없다. 한 warp가 NVLink를 통해 다른 GPU를 기다릴 때 하드웨어에는 인터럽트도, 주소 대기 기능도 없다. warp는 플래그가 바뀔 때까지 폴링한다. 글의 마이크로벤치마크에 따르면 GPU 범위 acquire 폴링은 141~145ns, 동일 노드의 두 GPU 사이 단방향 플래그는 약 1.9µs, 대량 복사 뒤에 큐잉된 폴링은 1.0~2.6µs다. 펜스는 유휴 시 0.4µs, TMA 로드가 진행 중일 때 3.6~4.0µs가 든다. 작업이 완벽히 분산되어 있어도 GPU들은 서로 최대 4µs까지 뒤처질 수 있고, 레이어의 모든 정지 지점은 가장 느린 GPU를 기다린다. 메가커널은 펜스를 일반 저장 슬롯에 기록한 센티널 값으로 대체한다. 즉 데이터로 채운 플래그, all-reduce용 +0/−0, bf16 gather용 all-ones NaN을 사용하고, 한 슬롯을 사용하는 동안 다른 유휴 복사본을 재설정한다.

이 숫자들은 결과가 마법이 아니라 엔지니어링처럼 느껴지게 만드는 수치입니다: 5.9 ms 단계 대비 4 µs의 GPU 스큐는 작지만, 레이어당 네 번의 동기화 중단이 62개 레이어에 걸쳐 곱해지면 잘못했을 때 2× 이득을 잡아먹는 바로 그런 종류의 오버헤드입니다.

5.9ms가 여전히 무엇에 쓰이고 있는지

이 게시물은 자체 비판을 담고 있는데, 그중 가장 유용한 대목은 두 번째 GEMM 이후의 공백에 관한 것이다. 블록들은 레이어 시작 후 55~64µs 사이(10~90백분위)에 해당 GEMM을 완료하지만, 결합(combine)은 모든 GPU의 출력을 기다리므로 65µs가 되어야 시작할 수 있다. GEMM 2가 끝난 뒤 레이어가 끝날 때까지 블록은 평균 약 12.5µs 동안 아무 작업 없이 머문다. 이는 전체 유휴 시간의 거의 절반에 해당한다. 저자들이 다음으로 살펴보고 있다고 말하는 것은 메가커널이 이미 제거한 커널 실행 오버헤드가 아니라 바로 이 사슬이다.

또한 다음 레이어의 어텐션이 단순히 더 일찍 시작할 수 없는 이유도 설명한다. PDL은 첫 블록이 종료되기 약 4µs 전인 79µs에 이를 시작하지만, 해당 블록들은 이전 블록들이 84µs에 SM에서 빠져나온 뒤에야 배치되므로 더 일찍 시작해도 더 빨리 시작되는 것은 없다. 해결책은 기존 KV 행을 65µs에 L2로 프리페치해 그곳에서 대기시키는 것이다. 이는 가중치 프리페치와 같은 패턴이다. 대기를 제거할 수 없을 때는 그 대기 속으로 작업을 옮기는 것이다.

프로파일링은 또 다른 구조적 문제이며, 이는 좋은 문제입니다: 커널은 모든 외부 프로파일러에게 레이어당 하나의 긴 막대처럼 보이므로 스스로 시간을 측정합니다. 각 블록의 레인 0은 globaltimer 레지스터를 통해 전역 배열에 타임스탬프를 기록하며, 저자들은 해당 계측의 비용을 레이어 실행당 약 2.5µs, 스텝의 2.8%로 측정했고, 출력은 비트 단위로 동일하게 유지되며 빌드 플래그로 컴파일에서 제외할 수 있습니다. 자체 프로파일러를 포함해야 하는 커널은 운영 체제가 된 커널이며, 이는 바로 해당 게시물이 사용하는 프레이밍입니다: 프로세스는 모든 GPU 리소스를 보유한 하나의 커널이고, 스레드는 선점 없는 워프이며, I/O 대기는 플래그에 대한 바쁜 대기이고, DMA는 mbarrier로 추적되는 TMA이며, 메시지는 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 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인가 — 16개 토큰은 더 많은 전문가를 활성화하므로 검증 비용이 더 들지만, 팀은 더 좁은 체인으로 같은 길이에 도달하는 것보다 더 넓은 트리로 6.6에 도달하는 것이 더 쉽다고 판단했습니다.

• 드래프트 스택 — 대형 모델을 위한 8,000줄의 CUDA 및 PTX, 코드는 facebookresearch/sgl-muse-mega에 공개

비교를 주의 깊게 읽어 보십시오. 이는 두 모델 간의 벤치마크가 아닙니다. 서로 다른 하드웨어와 서로 다른 워크로드에서 Xiaomi가 발표한 구성과 SGLang 팀이 발표한 구성, 즉 두 가지 추측 디코딩 구성을 비교한 것입니다. 이 비교가 뒷받침하는 주장은 더 넓은 검증 트리가 동일한 수용 길이에 도달하는 실행 가능한 경로라는 것뿐이며, 이는 순위가 아니라 설계에 관한 주장입니다.

이 숫자가 무엇인지, 그리고 무엇이 아닌지

용량 계획이 반드시 제대로 짚어야 하는 부분이 바로 여기이며, 이는 글 자체가 솔직하게 인정하는 부분이기도 합니다. 1,109라는 수치는 단일 GSM8K 요청에서 20개의 디코드 단계를 측정한 중앙값으로, 텐서 병렬 차수 8, 8대의 GB300, 배치 크기 1, 추측 디코딩이 켜진 상태에서 나온 값입니다. 이 글은 메가커널이 현재 배치 크기 1만 지원하며, 이를 더 큰 배치로 확장하는 것이 다음 작업이라고 분명히 밝힙니다.

이들 제약은 각각 서로 다른 주장을 제거한다. 배치 크기가 1이라는 것은 서빙 팀이 사용하는 의미의 처리량 결과가 아니라는 뜻이다. 여기에는 배치 이득도, 큐도, 프롬프트 길이의 혼합도 없다. 데이터셋이 하나라는 것은 그 수치가 GSM8K의 프롬프트 및 출력 분포와, 저자들이 이후 자체 프로파일링에서 합성 라우팅으로 사용한 층당 49개 전문가의 측정된 라우팅 패턴에 묶여 있음을 의미한다. 하나의 요청에 대한 20단계라는 것은 분산이 요청 수가 아니라 단계 수에 의해 제한됨을 의미한다.

그리고 단일 요청은 오버랩에 있어 최악의 경우인데, 저자들도 스스로 말합니다: "요청이 하나일 때는 모든 대기가 드러나고, 더 많을 때는 거의 항상 준비된 작업이 있습니다." 따라서 메가커널의 2.1×는 지연 숨기기가 가장 어려운 영역을 측정한 것이며, 이 글은 더 큰 배치가 더 어려워지기보다는 더 쉬워져야 한다고 분명히 밝히고 있습니다. 이는 결과에 유리한 논거이며, 여전히 측정이 아니라 논거입니다.

그렇다면 오늘날 알 수 있는 것은 좁지만 실질적이다: 코드는 공개 저장소에 존재하고, 게시물은 검증 가능할 만큼 충분히 자세히 메커니즘을 설명하며, 내부 산술은 세 가지 독립적인 비율에 걸쳐 일치하고, 저자들은 추가 이득을 제한하는 세 가지 오버헤드를 문서화했다. 알 수 없는 것은 독립적인 팀이 자체 여덟 대의 GB300에서 1,109를 얻는지, 이 접근 방식이 실제 배치에서 살아남는지, 그리고 출시된 SGLang 빌드에 포함되는지다 — 게시물은 그렇다고 주장하지 않으며, 저장소는 설치할 수 있는 패키지라기보다 연구 산출물이다. 가장 많은 것을 판가름할 비교는 동일한 하드웨어에서 튜닝된 베이스라인과의 비교인데, "SGLang의 기본 체인"은 움직이는 표적이고 메가커널을 작성한 팀이 자신들이 앞서는 베이스라인도 튜닝하기 때문이다.

오늘날 대여할 수 있는 초당 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토큰으로 응답하고 직접 구축한 8-GPU 스택이 1,109로 응답한다면, 그 차이는 스택을 구축해야 할 이유가 아니다. GB300 8대는 수백만 달러 단위로 측정되는 자본 구매이며, 커널은 연구용 코드이고, 배치 크기 1이라는 제약은 현재 이 스택이 호스팅 엔드포인트가 이미 잘 처리하는 바로 그 워크로드, 즉 한 번에 하나의 요청에만 잘 맞는다는 뜻이다. 메가커널이 계산을 진짜로 바꾸는 경우는 이미 학습용으로 GPU를 구매했거나 데이터 상주 제약이 있는 워크로드를 위해 GPU를 구매한 팀이, 이제 낮은 배치에서 이들로부터 어떤 활용도를 얻을 수 있는지 묻는 상황이다. 그런 팀에게 동일한 실리콘, 동일한 수치 연산에서의 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보다 큰 배치. 저자들은 더 큰 배치가 다음 단계라고 말하며, 더 많은 토큰이 처리 중일 때 이전 연산의 출력 일부는 이미 준비되어 있고 나머지는 아직 도착하므로 더 큰 배치가 더 쉽다고 주장한다. 배치된 수치가 존재하는 순간, 비교는 "하나의 스트림 대 하나의 스트림"이기를 멈추고 서빙 팀이 활용할 수 있는 처리량 주장이 되기 시작한다.

• 독립적 재현. 여기 나오는 모든 수치는 저자 자신들의 하네스에서, 자신들의 하드웨어에서, 프로파일링 패스용으로 자신들이 만든 합성 라우팅을 사용해 나온 것입니다. 이는 커널 작업에서는 정상적인 일이며, 그래도 여전히 단일 출처이기는 합니다. 두 번째 팀이 1,109를 똑같이 재현하든, 재현하지 못하든, 그것이야말로 이 문제를 판가름할 사실입니다.

• 이를 포함하는 릴리스된 SGLang 빌드. 해당 글은 코드를 공개하고 설계를 설명하지만, 출시된 버전에서 메가커널이 이 모델의 기본 디코드 경로를 대체했다고는 말하지 않는다. 그렇게 되기 전까지는 이를 채택하는 것이 곧 연구용 리포지토리를 채택하는 것을 의미한다.

• 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.