Cursor自曝黑科技重写MoE生成机制让Blackwell推理性能翻倍

时间：2026-08-02 19:50

Cursor团队在BlackwellGPU上重构MoE生成机制，采用输出中心的WarpDecode，绕过传统专家中心路径的内存低效，实现高效推理。推理速度提升至1 84倍，吞吐量达到118-121tokens s，输出精度接近FP32标准，提升至1 4倍。该成果显著优化了模型部署效率与性能。

近期，Cursor 研发团队针对 Blackwell GPU 架构，对混合专家模型的生成机制展开了全面重构。实测结果显示，推理速度提升至原来的 1.84 倍，吞吐量稳定达到 118–121 tokens/s，同时输出精度与 FP32 的接近度提升了 1.4 倍。

Cursor 在 X 平台发文表示，他们重新设计了 MoE 模型在 Blackwell GPU 上生成 token 的底层逻辑，让推理速度实现了 1.84 倍的显著提升。

这组数据确实引人注目：

吞吐量从 64–66 tokens/s 直接跃升至 118–121 tokens/s，涨幅达 1.84 倍；
输出质量更贴近全精度 FP32 标准，精度接近度提升了 1.4 倍。

该方案的核心在于“以输出为中心”的 Warp Decode 技术，它直接针对传统 MoE 模型“以专家为中心”的生成方式中存在的内存效率低下与精度不足问题，进行了系统性优化，实现了全面突破。

今天，我们将深入解析：Cursor 的 Warp Decode 究竟做了哪些创新？为何能同时实现速度与精度的双重飞跃？

传统MoE：专家调度模式，效率瓶颈明显

当前主流的顶级大语言模型，大多采用混合专家架构——一个模型内部容纳数十甚至上百个“专家子网络”，在实际推理时仅激活其中一部分。例如，在某一层从 128 个专家中挑选 8 个参与计算，这样既能维持庞大的参数量，又能有效控制实际计算开销。

传统 MoE 的计算流程大致如下：

首先通过路由机制确定每个 token 应分配给哪个专家；
随后将分配给同一专家的 token 进行集中收集；
专家完成计算后，再将结果重新整合输出。

这种传统路径在大批量处理场景下表现良好，因为每个专家上的共享计算量足以分摊数据整理带来的额外开销。然而，在自回归解码阶段——也就是我们生成代码时——由于每次仅生成少量 token，无法提供足够的共享计算来支撑。传统路径中的八个阶段里，有五个阶段纯粹用于“数据管理”，本身并不执行任何实质性的计算任务。

落实到实际应用中，结果是：理论上的 MoE 高效性，在现实中却因大量时间耗费在数据搬运上，导致 GPU 带宽利用率偏低，运行速度缓慢。

Warp Decode：以“输出”为核心，精简中间环节

面对数据搬运效率低下的问题，Cursor 团队另辟蹊径。

我们首先来理解 warp decode 的具体含义。根据官方描述：在 Blackwell GPU 上进行小批量解码时，围绕输出而非专家来组织 kernel 能够获得更佳效果，Cursor 将这种新方法命名为“warp decode”。

现代 GPU 以由 32 条并行处理通道组成的组（即一个 warp）来执行指令。在 warp decode 中，每个 warp 只专注于计算一个输出值。warp 会直接从内存中流式读取所需的权重数据，将所有 8 个路由专家的结果汇总到一个持续累加的总值中，最终输出一个结果。

那么 warp decode 的具体运行机制是怎样的？简单来说，核心思路是从围绕“专家”转向围绕“输出”，尽可能精简中间环节。

warp decode 主要通过两种机制实现性能提升：一是移除传统路径所需的多余阶段和缓冲区，二是实现 warp 的完全独立性，从而带来更优的调度效率与更出色的延迟隐藏能力。

具体实现方式

每个 GPU warp 只负责一个输出标量，并且该 warp 在整个计算过程中“专一”地执行这一任务。
warps 之间完全独立，不存在跨 warp 的同步操作或共享可变状态。
整个 MoE 层被极致压缩为仅两个融合内核：
moe_gate_up_3d_batched 负责处理 gate 和 up 投影，warp 独立完成点积运算、SiLU 激活等操作，中间值直接在寄存器中计算，无需写入共享内存；
moe_down_3d_batched 负责处理 down 投影，每个 warp 循环遍历 top-k 专家，累加结果，最后通过 warp 级别的蝶形归约将部分和合并为最终输出。整个过程几乎全部在寄存器层面完成，避免了大量中间缓冲区和内存数据的往返传输。
蝶形归约技术将 warp 内 32 个 lane 的局部部分和快速合并为最终的输出标量。当 moe_down_3d_batched 内核处理完一个 token 对应的所有 top-k 专家后，每个 warp 已将来自不同专家的贡献累加到自身私有的 FP32 寄存器累加器中。此时，一条 __shfl_xor_sync 指令即可执行 warp 级别的蝶形归约，该指令在底层编译为 PTX 指令 shfl.sync.bfly。

显著效果：速度与精度同步提升

完全绕过共享内存，无需将中间结果写入 shared memory 再读回；无 L1 缓存往返、无 bank 冲突，所有操作在寄存器层面完成，延迟极低；无需显式屏障，同步逻辑已内置在指令的 lane mask 中，直接保证正确性。

根据官方测试，效果确实令人印象深刻：

在 Cursor 内部推理系统上，针对运行于 NVIDIA B200 的 Qwen-3 风格模型进行测试：
速度方面，端到端解码吞吐量提升 1.84 倍，在不同上下文长度下表现稳定（纯生成阶段优化）。
精度方面，输出与完整 FP32 参考值相比，接近程度提升 1.4 倍。
硬件效率方面，B200 在连续内存读取上的实测峰值为 6.8 TB/s (通过 copy kernel 测得)。在 B=32 时，warp decode 可稳定达到 3.95 TB/s，相当于该峰值的 58%。