Muse Spark Open：每秒 1,109 个 Token 是内核层面的变更，而非模型变更

Muse Spark Open 于 2026 年 10 月 15 日发布，本周其权重没有任何变化。改变的是这些权重底层的软件。2026 年 11 月 11 日，LMSYS 的 SGLang 团队发布了一篇题为“Muse Spark Open Ultrafast：将每秒 token 数推至 1,000 以上”的工程文章，其中的头条数字是：在八块 GB300 上，同一模型、同一硬件的解码吞吐量从每秒 517 个 token 提升到每秒 1,109 个 token。其机制是 megakernel：每层一个常驻内核，而不是每个解码步骤约 1,450 次独立的内核启动。所以这是一个服务层面的故事，而不是模型层面的故事；有趣的问题不是 1,109 是否亮眼——它确实亮眼——而是这些数字中哪些可以带入容量规划。简短的回答是：比你希望的要少：该基准测试是单请求、单数据集、批大小为 1，而且作者之外尚无人复现。

无人用作基准的基线：每个代币 1,450 次启动

一个普通的 GPU 内核只做一件事。它把一次矩阵乘法或一次注意力计算作为由分布在各流式多处理器上的线程块组成的网格来运行，然后退出。像 SGLang 这样的服务引擎会把这些内核串成一条链，而对于 Muse Spark Open 来说，这条链很长：文中给出的数字是每个解码步骤大约 1,450 次内核启动。在两个连续内核之间，GPU 会排空并重新填充。已完成各自分块的 SM 会空闲等待下一次启动；在 SM 内部，各个 warp 也会空闲，因为一个 warp 在发出加载指令后，直到数据返回之前都无事可做。

这正是巨型内核所要解决的开销，作者借用20世纪60年代的类比来解释它：早期的机器运行批处理作业，一个作业独占整台机器，磁带旋转时处理器却处于空闲状态。多道程序设计解决了这个问题：让多个作业常驻，每当当前作业阻塞时就切换到另一个作业。巨型内核就是把这一思想搬到 GPU 上——一个内核常驻于每个 SM，同时运行许多作业，因此 warp 可以在上一个操作完成之前就开始下一个操作的工作。

根据性能章节随附的图表，测得的效应是解码步骤从 12.9 ms 降到 5.9 ms。该图表将该步骤拆成草稿和验证两半：基线中为 1.16 ms 和 10.9 ms，有 1,450 个内核；megakernel 中为 0.44 ms 和 5.3 ms，有 97 个内核。12.9 到 5.9 约为 2.2×，1,109 除以 517 约为 2.1×，因此这两个主要比率彼此吻合。还有第三项检查，任何愿意做除法的人都可以进行：该步骤验证 16 个 token 并接受 6.6 个，因此 5.9 ms 内处理 6.6 个 token，约合每秒 1,120 个 token。这些算术在内部是一致的；对于这样的数字，这是首先要确认的事，而不是最后一件事。

这些数字没有说明的一点是，速度摊到每块 GPU 上究竟是多少。文章报告的是整个八 GPU 部署的总吞吐量。除以八，大约是每块 GPU 每秒 139 个 token——这是我在他们给出的数字上算出来的，并非文章发布的数字，也不能与你可能在另一台机器上看到的另一个模型的每 GPU 数字相提并论。

内核边界位于何处，以及它为何止步于此

megakernel 不会吞掉整个模型。每次启动覆盖的是一个解码器层的注意力，直到下一层的 qkv 投影和 KV 写入；这是刻意设定的边界，而非 nerve 的限制。它之所以可行，是因为一个层的注意力所读取的大部分内容——KV 历史——是在更早的步骤中写入的，不依赖任何正在执行中的内容，因此新的启动可以在其可用的一瞬间就开始加载它，甚至在上一层输出尚不存在之前。第一层和最后一层的操作略有不同，文章也指出了这一点。

在一次启动内，形状是固定的：在 144 个 SM 的每一个上，有一个由十二个 warp 组成的线程块，以四个为一簇启动，并在可获得更低延迟归约的地方使用分布式共享内存。这十二个 warp 并不是可互换的。Warp 0 加载 KV 和权重，并利用其空闲时间预取下一操作的权重。Warp 1 和 2 运行激活和 Tensor Core。Warp 3 到 11 分担 attention、稠密 GEMM、专家 GEMM 以及 all-reduce 和 all-gather 接收。在滑动窗口层中，warp 4 到 11 加载自己的 KV，而不是等待 warp 0。

要做到这一点，需要改变模型的拆分方式。Muse Spark Open 是一个专家混合模型，将每个 token 路由到 256 个专家中的 8 个，专家以 4 位 NVFP4 存储，部分层则保留为 bf16。它有 62 个解码器层，其中 46 个对最近 2,048 个 token 使用滑动窗口注意力，16 个使用全局注意力；此外还有一个 pooler 投影，在进入专家之前压缩每个 token 向量，以及一个 unpooler，在之后将其恢复。该 megakernel 在 N 维度上对两个 MoE GEMM 进行分片，这迫使每次 GEMM 后都要进行一次 all-gather——第一个之后是 fp4 all-gather，第二个之后是 bf16 all-gather。这比标准张量并行多了一个通信步骤；标准张量并行在中间维度上分片，只需要一次 all-reduce，但作者认为它总体上传输的数据更少、能更好地扩展到他们接下来想要的批大小，并且可以与专家并行做到按位一致，从而匹配训练数值。

这个细节比看起来更重要。推理团队通常默认：更快的 kernel 会产生略有不同的数值，然后悄悄重新调整其评估阈值。一个与训练侧并行化按位相等的 kernel 则彻底消除了这一争议，而正是这类特性会让一项优化被真正采用，而不只是被赞赏。

这么大的一个内核要付三笔账单

megakernel 并非没有代价。这篇文章对各项开销异常坦诚，而这部分值得反复阅读，因为它们同样能告诉你还剩下多少空间。

第一项代价是指令缓存。一个 SM 大约缓存 128 KB 的指令，而这些内核每个都有 215 到 314 KB 的机器码，因为每个内核都包含某一层的所有阶段。一次启动只会把最初的 32 到 64 KB 带到 SM 附近；超出这个窗口的所有内容都会在首次使用时逐行取入。作者对此做了一个干净的实验并报告了结果：两个 100 KB 的块单独执行时各需 12.4 µs，而背靠背拼接后需要 19.5 µs；把其中一个块缩小到 20 KB，另一个块就恢复到了单独执行时的速度。指令与数据经过同一个 L2 缓存，因此一个超出指令缓存容量的内核最终会与自己的权重加载争抢缓存。

第二项成本是权重预取，文章将其视为一个放置问题，并给出三种答案。共享内存为每个 SM 提供一个 5 × 32 KB 的环形缓冲区，整个 GPU 上约 23 MB。其余部分放入 L2，GB300 的 L2 容量约为 135 MB，但 L2 中的任何内容都需要再经历一次更短的传输进入共享内存，而且 L2 还要与所有其他正在执行的任务共享。两者叠加：在一层接近尾声时，环形缓冲区已经保存了下一层的第一批 qkv 权重。在此基础上，warp 3 会读取上一步路由到了哪些专家，并把最常用的 14 个专家在本 GPU 上的切片拉入 L2，每个 GPU 约 50 MB，赌的是连续步骤的路由相似。这个赌注在 5.3 ms 的一步中约值 0.2 ms，而一旦赌错就可能消耗带宽。

第三项代价是同步，而它正是没有干净答案的那一项。当一个 warp 通过 NVLink 等待另一个 GPU 时，硬件中既没有中断，也没有 wait-on-address；warp 会轮询一个标志，直到它发生变化。文章中的微基准测试显示：GPU 作用域的 acquire 轮询为 141 到 145 ns；同一节点上两个 GPU 之间的单向标志约为 1.9 µs；排在批量拷贝之后的轮询为 1.0 到 2.6 µs。栅栏在空闲时开销为 0.4 µs，在有 TMA 加载正在进行时为 3.6 到 4.0 µs。即使工作分配得完美，GPU 之间也可能彼此落后最多 4 µs，而该层中的每一次停止都要等待最慢的那个。这个 megakernel 用写入普通存储槽的哨兵值取代了栅栏——用数据填充的标志、全归约使用 +0/−0、bf16 gather 使用全 1 的 NaN——并在其中一个副本被使用时重置每个槽位的空闲副本。

这些数字让结果更像是工程而不是魔法：相对于 5.9 ms 的步长，4 µs 的 GPU 偏斜很小，但每层四次同步停顿再乘以 62 层，正是那种一旦处理不当就会吞掉 2× 收益的开销。

这 5.9 毫秒仍然花在了什么上

这篇文章包含了它自己的评析，其中最有价值的一句话是关于第二个 GEMM 之后的空档。各个块会在进入该层后的 55 到 64 µs 之间完成那个 GEMM（第 10 到第 90 百分位），但 combine 阶段要等待每个 GPU 的输出，因此直到 65 µs 才能开始。从 GEMM 2 结束到该层结束，一个块平均约有 12.5 µs 处于无工作可做的状态——这几乎占其总空闲时间的一半。作者表示，他们接下来要关注的是这条链，而不是 megakernel 已经消除的内核启动开销。

他们还解释了为什么下一层注意力的启动不能简单地提前。PDL 在 79 µs 时启动它，比第一个块退出早约 4 µs，但它的块只有在先前的块于 84 µs 离开 SM 后才会落地，因此更早启动并不会让任何工作更早开始。解决办法是在 65 µs 时将其旧的 KV 行预取到 L2 中，让它们在那里等待。这与权重预取是相同的模式：当你无法消除等待时，就把工作移入等待之中。

性能分析是另一个结构性问题，而且是个好问题：在任何外部性能分析器看来，这个 kernel 每一层都像一条长长的横条，所以它自己给自己计时。每个 block 的 lane 0 通过 globaltimer 寄存器把时间戳写入一个全局数组，作者测得这种插桩的成本约为每次层启动 2.5 µs，占一个 step 的 2.8%，同时输出保持逐位相同，并提供一个构建标志把它编译剔除。一个必须自带性能分析器的 kernel，已经变成了一个操作系统，而这正是文章采用的框架：一个进程就是一个持有全部 GPU 资源的 kernel，线程就是不可抢占的 warp，等待 I/O 就是在标志位上忙等，DMA 就是由 mbarrier 跟踪的 TMA，消息就是跨 NVLink 的普通存储。

草稿模型也变快了，而这正是成功的一半。

Muse Spark Open 的解码路径是推测式的：草稿模型提出 token，大模型在单遍中验证它们。这里的草稿模型是 DFlash V2，它在单遍中猜测接下来的 15 个 token，而大模型则在一次覆盖 16 个 token 的遍历中检查所有这些 token。验证使用由草稿生成的候选 token 树，而不是线性链；文章称这略好一些。

团队对草稿模型采用了同样的 megakernel 处理方式：五个草稿层各自作为一次启动运行，另外还有三个其他层；这使得草稿模型从 175 个 kernel 降到 8 个，每步从 1.16 ms 降到 440 µs，在步骤的这一半上实现了 2.6× 的提升。接受率数据值得与文中给出的最接近的公开对比放在一起看。

• 接受长度——Muse Spark Open 每步验证 16 个 token，平均接受 6.6 个；小米的 MiMo-V2.5-Pro 此前已突破每秒 1,000 个 token，每步验证 8 个，平均接受 6.3 个

• 为什么是 16——16 个 token 的验证成本更高，因为它们会激活更多专家，但团队发现，用更宽的树达到 6.6，比用更窄的链达到相同长度更容易

• 草稿栈——为大模型编写的 8000 行 CUDA 和 PTX，代码已发布在 facebookresearch/sgl-muse-mega

请仔细阅读这项对比，因为它不是两个模型之间的基准测试。它比较的是两种投机解码配置：一种由 Xiaomi 发布，另一种由 SGLang 团队发布，二者运行在不同的硬件和不同的工作负载上。它唯一能支持的论断是：更宽的验证树是实现相同接受长度的一条可行路径，这是一种设计主张，而不是排名。

这个数字是什么，以及它不是什么

这是容量规划必须做对的部分，也是文章本身坦诚承认的部分。1,109 这个数字是在单个 GSM8K 请求上、张量并行度为 8、使用八块 GB300、批大小为 1 且开启推测解码时，20 个解码步的中位数。文章明确表示，如今的 megakernel 仅支持批大小为 1，将其扩展到更大的批是下一步工作。

{{1}}这些限制各自排除的是一种不同的说法。{{/1}} {{2}}批大小为一意味着，在服务团队使用这个词的意义上，这并不是一个吞吐量结果：其中没有批处理收益，没有队列，也没有不同提示长度的混合。{{/2}} {{3}}单一数据集意味着，这个数字与 GSM8K 的提示和输出分布绑定，也与 {{4}}每层 49 个专家{{/4}} 的实测路由模式绑定，而作者随后在自己的性能剖析中将其用作合成路由。{{/3}} {{5}}一个请求的二十个步骤意味着，方差受步骤数约束，而不是受请求数约束。{{/5}}

而单个请求是重叠的最坏情况，作者自己也这么说：“只有一个请求时，每一次等待都会暴露出来；请求更多时，几乎总有就绪的工作可做。”因此，megakernel 的 2.1× 是对延迟隐藏最难情形的一次测量，而且文章明确表示，更大的批次应当更容易，而不是更难。这是有利于该结果的一个论据，而且它仍然只是论据，而不是测量。

那么，今天可知的东西是狭窄但真实的：代码存在于一个公共仓库中，文章对机制的描述详细到足以被核查，内部算术在三个独立比率之间相互一致，而且作者记录了限制进一步提升的三项开销。不可知的是，一个独立团队能否在自己的八块 GB300 上得到 1,109，这一方法能否经受真实批次的考验，以及它是否能进入已发布的 SGLang 构建——文章并未声称已经做到，而且该仓库是研究产物，而不是你可以安装的软件包。最能定论的比较，是在相同硬件上与一个经过调优的基线进行对比，因为“SGLang 的默认链路”是一个移动靶标，而编写这个 megakernel 的团队也调优了它所击败的基线。

它与你现在可以租用的每秒 169 个 token 有什么关系

Artificial Analysis 对 M​use Spark Open 记录了一个截然不同的数字：输出速度为每秒 169 个 token，这是端到端测量的，包含首个 token 的时间，针对托管端点，首个回答 token 时间为 13.81 秒。这并不矛盾，也不是反驳。这是对不同部署的不同测量：一条流经他人的托管 API，而一个请求则经你自己的八 GPU 节点并搭配研究内核。

这个差距是这则故事中对尚未拥有这套硬件的人最有用的部分。如果一个端点每秒响应 169 个 token，而手工搭建的八 GPU 堆栈每秒响应 1,109 个，那么这一差异并不是搭建该堆栈的理由。八块 GB300 是数以百万美元计的资金采购，内核是研究代码，而批大小为 1 的限制意味着该堆栈目前恰好擅长托管端点已经处理得很好的工作负载：一次一个请求。megakernel 真正改变权衡之处，在于一个已经为训练购买了 GPU 的团队，或者一个受数据驻留约束的工作负载，现在在问：在低批大小下能从这些 GPU 中获得多少利用率。对那个团队来说，在相同芯片、相同数值计算下实现 2.1× 是一个真实数字，背后附着一张真实账单，而账单已经付清。

对其他所有人来说，最坦率的解读是：这篇文章是一个信号，表明推理基础设施正朝哪个方向发展，而不是你本周就能采用的产品。这个方向与同一个实验室全年一直在发布的内容一致：更少、更大的内核，更多调度决策被移入内核内部，以及把 GPU 当作一台可编程的机器，而不是一组需要按顺序调用的库函数。

这也是理解我们自己立场的正确框架，这一点值得直说。M​use Spark Open 并不在 OrcaRouter 上；它是一个下载项、一个受限访问的仓库和一份许可证，而不是我们提供的端点，我们也不会把它描述成端点。Muse Spark 1.2 和 Muse Spark 1.1 是 OrcaRouter 实际路由的两个 Muse 模型，它们按 Meta 自己的标价提供，因为供应商的标价是原样传递的，而不是被加价。如果你要问的是本季度该把哪个模型放在产品背后，而不是编译哪个内核，那么这条路更短：一个兼容 OpenAI 的端点，覆盖 200+ 个模型，供应商标价以 0% 加价原样传递，外加自动故障转移，以及一个路由 DSL，它能通过配置而不是重新部署代码来在模型之间转移流量。

什么会改变这件作品？

四件事，按它们对局面影响的大小从大到小排列。

• 大于 1 的批次。作者表示下一步是更大的批次，并认为这会更容易，因为当更多 token 同时在处理中时，上一个算子的部分输出已经就绪，而其余输出仍在到达。一旦存在批处理后的数字，比较就不再是“一个流对另一个流”，而开始成为服务团队可以采用的吞吐量说法。

• 独立复现。这里的每个数字都来自作者自己的测试框架、自己的硬件，以及他们为性能分析阶段自建的合成路由。这对内核工作来说很正常，但它仍然是单一来源。第二个团队能否复现出 1,109，才是能最终定论的事实。

• 一个包含该功能的已发布 SGLang 构建版本。该文章发布了代码并描述了设计；但并未说明该 megakernel 已在某个正式发布版本中取代了此模型的默认解码路径。在这一点实现之前，采用它就意味着采用一个研究仓库。

• GEMM-2-to-combine 之间的间隙闭合。作者已经指出，那里每块 12.5 µs 的空闲时间几乎占剩余时间的一半。如果这条链路得到修复，同一硬件会产生明显更好的步进时间，而 2.1× 将成为更大数字的下限，而不是头条亮点。

那份清单上的任何内容都不能成为贬低这一结果的理由。一个团队写了 8,000 行 CUDA 和 PTX，把一步从 12.9 ms 推进到 5.9 ms，公开了代码，随后又公开了剩余时间究竟花在哪里——这正是这个领域声称重视却很少做到的事。正确的总结并不是 Muse Spark Open 变快了。而是：某个特定的服务栈，在特定硬件上、特定批大小下，变快了 2.1×，并且实现下一个 2× 的配方现在已公开。