Cursor 凌晨自曝黑科技：重写 MoE 生成机制，Blackwell 推理性能直接翻倍！网友：直接榨干B200最后一滴带宽！

今天早上，Cursor 在X上发布一条推文：“我们重建了 MoE 模型在 Blackwell GPU 上生成 Tokens 的方式，导致推理速度快了 1.84 倍。”

今天早上，Cursor 在X上发布一条推文：“我们重建了 MoE 模型在 Blackwell GPU 上生成 Tokens 的方式，导致推理速度快了 1.84 倍。”

数据表现超级亮眼：

Cursor使用的“输出中心”的 Warp Decode直接把传统 MoE 模型的“专家中心”生成方式中内存使用效率低、准确性低等问题，一次性全部解决！

今天，我们来拆解一下：Cursor 的 Warp Decode 到底做了什么？为什么能同时实现速度和精度的双提升？

现在的顶级大模型，大多采用Mixture of Experts架构—— 一个模型里面有几十个甚至上百个“专家”子网络，输入的时候只激活其中一部分专家（例如在某一层从 128 个专家中选择 8 个来干活），这样既能保持超大参数量，又能控制实际计算量。

传统MoE的计算路径大概是这样的：

这种传统的MoE 路径在大批量场景下效果很好，因为每个专家上的共享工作足以摊薄整理数据的额外开销。

但在自回归解码阶段——也就是我们用AI生成代码时，由于一次只生成几个 token，没有足够的共享工作来支撑。传统路径中的八个阶段里，有五个阶段纯粹是“数据管理”，本身并不进行任何实际计算。

到了我们的实际的应用里，结果就是：在理论上MoE很高效，但实际上将太多时间用来运送数据，跑起来GPU带宽利用率低，速度慢。

既然搬运数据太慢，Cursor 直接换一条路走。

我们先来了解一下warp decode 具体是什么？根据官方描述：

在 Blackwell GPU 上进行小批量解码时，围绕输出而非专家来组织 kernel 效果更好。Cursor 将这种方法称为“warp decode”。

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

而 warp decode 是如何运行的呢？

简单来说，从围绕“专家”到围绕“输出”，中间的环节能砍则砍。

warp decode 主要通过两种机制提升性能：一是去掉传统路径所需的阶段和缓冲区，二是实现 warp 的独立性，从而带来更优的调度效果和更好的延迟隐藏能力。

具体做法：

每个GPU warp只负责一个输出标量（scalar），并且这个warp在整个计算过程中“终身”只干这一件事。
warps之间完全独立，没有跨warp的同步或共享可变状态。
整个MoE层被极致压缩成仅两个融合内核（fused kernels）：moe_gate_up_3d_batched：处理gate和up投影，warp独立完成dot product、SiLU激活等，中间值直接在寄存器里算，不写共享内存。moe_down_3d_batched：处理down投影，每个warp循环遍历top-k专家，累加结果，最后用warp级别的butterfly reduction（__shfl_xor_sync指令）把部分和合并成最终输出。整个过程几乎全在寄存器里完成，避免了大量中间缓冲区和内存往返。
蝶形归约（Butterfly Reduction）：将warp内32个lane的局部部分和（lane-local partial sums）快速合并成最终的输出标量。当moe_down_3d_batched内核处理完一个token对应的所有top-k专家后，每个warp已经把来自不同专家的贡献累加到了自己私有的FP32寄存器累加器。这时，一条指令__shfl_xor_sync来做warp级别的蝶形归约，并且直接编译成底层的PTX指令shfl.sync.bfly