Cursor为Blackwell从零构建MXFP8内核，MoE层提速3.5倍，端到端训练提速1.5倍
文章摘要
【关 键 词】 AI模型、硬件升级、MoE层优化、Cursor团队、量化内核
在构建更强大 AI 模型的竞赛中，传统做法是升级硬件，但 Cursor 发现释放下一代 GPU 潜力并非易事。从 NVIDIA 的 Hopper H100s 升级到 Blackwell B200s 后，团队遭遇「升级陷阱」，硬件性能翻倍，实际训练速度却因 MoE 层效率低下而变慢，新架构放大了数据搬运和量化开销。
为解决这一问题，Cursor 回归基础，从零开始重写整个混合专家（MoE）训练层。他们抛弃对现有 CUDA 库的依赖，直接针对 TMEM 新特性设计数据流管线，避免寄存器搬运开销；将量化与反量化逻辑融入内核计算流程，压缩内存带宽占用；优化 MXFP8 的 microscaling 实现，在保证训练收敛质量的同时提升性能。最终，MoE 层前向和反向传播提速 3.5 倍，端到端训练速度在 Blackwell 上快了 1.5 倍，相比 Hopper GPU 方案实现 2 倍加速。
现有 MoE 内核在 Blackwell 上失效的原因主要有两点。一是张量内存（TMEM）瓶颈，Blackwell 引入新的 TMEM 存储累加结果，导致数据往返低效，产生「气泡」，降低执行效率，且反量化速度滞后于计算速度。二是存在被忽视的「量化税」，数据量化过程本身成为性能杀手，现有开源量化内核带宽利用率低，生成的缩放因子布局与硬件指令不兼容。
Cursor 团队放弃高层依赖，用纯 CUDA 和 PTX 汇编语言编写 MoE 层的 GPU 代码。其优化策略包括：围绕原生的 `tcgen05.mma` 指令构建内核，消除低效数据移动；设计高效的数据流水线，采用「Warp 专精」和 2 – CTA 模式等技术，实现高度并行和性能提升；针对 MoE 工作负载应用「专家级超分组」的 L2 缓存优化启发式算法，限制性能下降。此外，团队开发了自定义的 MXFP8 量化内核，其内核输出的数据内存布局与指令要求一致，避免额外「重塑」步骤。同时，确定了特定的低精度「配方」，在不影响训练质量的情况下提供最高速度。
原文和模型
【原文链接】 阅读原文 [ 1888字 | 8分钟 ]
【原文作者】 机器之心
【摘要模型】 doubao-1-5-pro-32k-250115
【摘要评分】 ★★★☆☆
相关文章
