Cursor为Blackwell从零构建MXFP8内核，MoE层提速3.5倍，端到端训练提速1.5倍

在AI模型性能竞赛中追求更强算力的传统路径看似简单：只需不断升级硬件设备。

但Cursor团队发现NVIDIA新一代Blackwell B200 GPU的实际表现远非即插即用般理想——当他们从Hopper H100升级至新架构后遭遇”升级陷阱”：虽然硬件性能翻倍提升，但混合专家（MoE）层效率低下导致整体训练速度不升反降。

这种新型GPU架构反而放大了数据搬运和量化操作的成本消耗。

面对这一挑战，团队回归底层设计思维，在GPU内核层面重新构建了MoE训练系统——如同为赛车重新研发专用轮胎般彻底重构计算流程。

通过摆脱对现有CUDA库的依赖并自主编写核心代码： – MXFP8量化方案采用微缩放技术将张量分割为32元素块并独立计算缩放因子 – 定制化内核消除TMEM与CUDA核心间低效的数据搬运 – 采用Warp专精策略实现线程组级任务分配 – 开发专家级超分组L2缓存优化算法 最终实现： MoE层前向/反向传播速度提升3.5倍 Blackwell平台端到端训练加速1.5倍 对比初始Hopper方案整体提速达2倍 在低精度计算领域普遍采用FP8格式时遇到的关键难题——数值精度丢失被MXFP8巧妙化解：通过将张量分割成多个小块并分别为每个块计算独立缩放因子（scale factor），确保所有数据都能被有效压缩至FP8表示范围同时保留必要精度。

针对Blackwell架构特有的TMEM内存瓶颈展开深度优化： – 张量核心计算结果直接利用tcgen05.mma原生指令处理 – 采用双CTA模式使两个流式处理器协同工作降低内存流量 – 线程组专精策略实现加载/计算任务并行化 实测数据显示： 反量化耗时从Hopper上的矩阵乘法1.03倍激增至Blackwell平台1.76倍 自研量化内核带宽达6.2TB/s超越开源方案4.5TB/s水平 分组矩阵乘法性能损耗控制在4%以内 经过多维度优化后，在保持训练质量的前提下（BF16与MXFP8损失曲线完全重合），成功将反向传播中的转置-量化开销从76%占比降至可控范围，并通过专属低精度配方（FP8E4M3+32块缩放）实现了速度与精度的最佳平衡。