在AI模型性能提升的竞赛中，硬件升级一直是主流路径。但Cursor团队近期的实测数据表明，单纯依赖最新GPU硬件并不能直接释放Blackwell架构的全部潜能。当他们将NVIDIA Hopper H100s升级为Blackwell B200s时，意外发现性能提升背后隐藏着更深层的瓶颈。这种升级带来的硬件性能翻倍，反而在实际训练中被MoE架构的效率问题拖累。

这种现象就像给赛车换上动力更强的引擎，却发现原有轮胎无法承载新动力。Cursor团队在实验中发现，Blackwell架构的TMEM特性与现有CUDA库存在兼容性问题，导致数据搬运和量化过程成为新的性能消耗点。这直接导致训练速度未达预期，甚至出现性能倒退的情况。团队通过深度分析发现，现有MoE层的设计在Blackwell架构下存在系统性缺陷。

为解决这个问题，Cursor团队选择从底层架构入手。他们摒弃传统依赖，采用纯CUDA和PTX汇编语言，对MoE训练层进行全栈式重写。这种定制化改造让GPU硬件特性与算法流程完美契合，最终实现了性能突破。通过优化数据流管线、量化逻辑和内存带宽利用率，团队成功将Blackwell架构的性能潜力释放到极致。

为什么现有MoE内核在Blackwell上失效？

传统低精度计算方案在Blackwell架构下暴露出明显短板。以FP8量化为例，简单转换高精度数值会导致信息丢失。微缩放（MX）技术通过为每个数据块分配独立缩放因子，有效解决了这个问题。Cursor团队采用的MXFP8格式正是这种技术的典型应用，它能在保持精度的同时获得低精度计算优势。

但Blackwell架构引入的TMEM特性给量化流程带来新挑战。与Hopper架构相比，TMEM存储方式需要经历TMEM→寄存器→CUDA核心→TMEM的数据往返过程。这种异步传输模式在张量核心计算管线中产生"气泡"，显著降低执行效率。数据显示，在特定配置下，Blackwell架构的反量化耗时达到矩阵乘法本身的1.76倍，远超Hopper架构的1.03倍。

被忽视的"量化税"问题同样值得关注。以典型MoE矩阵乘法为例，计算本身仅需1.16毫秒，但量化过程就需要搬运近2.9GB数据，耗时0.44毫秒。这种数据搬运开销在反向传播中因转置-量化操作而翻倍，达到0.88毫秒，占比高达76%。这意味着，若不优化量化流程，MXFP8带来的性能提升可能被完全抵消。

Cursor如何从零重写MoE层？

面对这些挑战，Cursor团队选择放弃传统依赖。他们采用纯CUDA和PTX汇编语言，对MoE层进行全栈式重写。这种定制化改造让GPU硬件特性与算法流程完美契合，最终实现了性能突破。通过优化数据流管线、量化逻辑和内存带宽利用率，团队成功将Blackwell架构的性能潜力释放到极致。

优化策略

拥抱原生硬件指令

Cursor团队没有与TMEM架构对抗，而是围绕原生的tcgen05.mma指令构建内核。这种设计让GPU硬件自身处理MXFP8所需的缩放，完全消除TMEM和CUDA核心之间的低效数据移动。这种深度适配确保了计算流程的流畅性，避免了传统方案中的性能损耗。

设计高效的数据流水线

团队实现了一套复杂的流水线系统，采用" Warp 专精"技术将不同任务分配给线程组。例如，Warp 0负责数据加载，Warp 1处理缩放因子，Warp 2完成数据迁移，Warp 3执行矩阵乘法。这种分工模式让各个环节高度并行，显著提升整体效率。

2-CTA模式允许两个SM协同完成单个矩阵乘法，通过共享B矩阵减少内存流量，带来15-20%的性能提升。这种设计巧妙利用硬件特性，将计算资源利用率提升到新高度。

针对MoE工作负载进行优化

针对MoE训练中的分组矩阵乘法，团队应用"专家级超分组"优化算法。这种L2缓存优化策略确保内存访问模式高效，将标准矩阵乘法与分组矩阵乘法之间的性能差距控制在4%以内。这种细粒度优化让硬件资源利用达到最佳状态。

"秘密武器"：量化内核与低精度配方

团队开发的自定义MXFP8量化内核成为关键突破点。微基准测试显示，其内存带宽超过6.2TB/s，相比现有开源工具的4.5TB/s有明显提升。这种内核输出的数据内存布局与tcgen05.mma指令完全匹配，避免了其他工具所需的额外重塑步骤。

团队还确定了一种特定的低精度配方，使用FP8E4M3格式在保持训练质量的同时实现最高速度。通过对比实验发现，这种格式的训练损失收敛情况与BF16格式几乎一致，证明性能提升并未以牺牲准确性为代价。

最终测试数据显示，两种方法在10k步训练中的损失曲线几乎没有差异。这种性能突破为AI模型训练提供了新的优化思路，也为AI鞋履和服装工具的性能提升提供了技术参考。