第一次在H100上碰FP8 GEMM的时候，我差点被一整套奇怪的数值结果劝退：矩阵乘出来不是全0就是NaN，偶尔正常一次结果又对不上参考实现。后来把关注点从“FP8就是少了一半精度”转到“更底层的数据管理和流水线调度”上，才慢慢理解像DeepGEMM这类库为什么值得拆开看。DeepGEMM是一个面向FP8精度的轻量级GEMM库，专门处理稠密矩阵乘和MoE分段矩阵乘，代码量不大，却把从数值缩放到Hopper架构底层调度的关键点全部串了起来。如果你在做推理加速、训练框架底层开发，或者只是想搞懂FP8内核背后的原理，这篇东西应该能给你省不少时间。

1. 从FP16到FP8，它卡住了不少团队

1.1 硬件算力翻倍，但软件没跟上

GPU的浮点演进其实和写代码的懒惰程度是矛盾的。老框架里到处是FP16的假设：算子注册表、自动混合精度策略、自定义CUDA kernel的模板参数，全都围着FP16转。换FP8不是把某个dtype字符串改掉就完事，几乎每个环节都要重新验证。

但硬件层面的诱惑实在太大。Hopper架构里FP8的峰值算力直接是FP16的两倍，内存带宽消耗也少一半。对大模型推理来说，这等同于在相同显存和功耗下塞进更多请求；对训练来说，则代表更大的batch和更短的迭代时间。于是问题变成：既然硬件给得这么爽快，为什么很多团队迟迟没切换？

答案很简单：软件没有跟上。FP8 GEMM不是“把float换成8位”这么朴素。乘积结果的范围、累加过程中的精度、缩放因子的计算位置，每一样都会改变内核的数据流。很多人上手之后发现，要么精度崩掉，要么性能还不如FP16，最后只能退回老方案。DeepGEMM这类库的价值，恰恰在于把这些问题在库内部一次性解决好，让上层使用者不必重造轮子。

1.2 大模型GEMM的两个“反标准”形状

传统GEMM库的设计目标是规整的大矩阵乘法：M、N、K都很大，用分块策略把SM喂得饱饱的，reduction维度尽可能长。但大模型场景里的GEMM形状往往并不规整。

第一个典型场景是推理生成阶段。M经常只有1到16，K和N可以到几千甚至上万。这种细长形状会让标准tile划分产生大量空转：每个block算不了多少活就结束了，SM的利用率被严重拉低。哪怕你有很高的FP8峰值算力，实际吞吐也会非常惨。

第二个典型场景是MoE。多个专家权重共享一份激活输入，每个专家分到的token数量不固定，如果像常规做法那样逐个专家单独调GEMM，kernel launch次数会爆炸，单个小GEMM又凑不满一个完整wave，SM不断闲置。DeepGEMM把这一类问题抽象成“分组/分段GEMM”，用单次内核调用处理多个专家的计算，目的是让本来细碎的工作变得足够稠密。这个方向，恰恰是很多通用库没有为FP8做精细优化的死角。

1.3 “小而精”的设计哲学很关键

我最初看到DeepGEMM源码时，第一反应是：核心逻辑怎么这么短？它并没有把cuBLAS几千行内部逻辑重写一遍，而是围绕FP8 GEMM的关键路径做了精准裁剪。对于库开发者来说，代码短不代表功能弱，反而说明每一行都有明确目的，便于阅读、审计和集成。

这也是我认为它值得学习的原因之一。很多高性能库动辄上万行，内部优化堆叠得密密麻麻，你根本分不清哪些是架构约束、哪些是历史包袱。DeepGEMM把优化集中在“FP8数据布局”“缩放因子处理”“Hopper异步流水线”“MoE分段调度”这几个真正影响性能的点上，剩下的保持朴素。这种极致的取舍，对想自己动手改内核的人来说是很好的教材。

2. FP8数值细节，决定整个库的设计方向

2.1 e4m3、e5m2怎么选

FP8并不是一种格式，而是两种：e4m3和e5m2。它们的差异集中在指数位和尾数位的分配上。我整理了一张常用的对照表：

e4m3因为尾数只有3位，精度很粗糙，但动态范围在4位指数下也能覆盖常见的激活分布；e5m2的指数多一位，能表达更大的数值，但尾数更少，精度进一步下降。工程里最常见的组合是：激活和权重用e4m3，梯度、误差或者需要较大动态范围的数据用e5m2。

这里有个很容易踩的坑：不要想当然地觉得所有张量都适合e4m3。如果数值分布本身很宽，比如某些层的中间激活偶尔会冒出一个明显偏大的channel，e4m3直接就会溢出到NaN。真正可靠的方案都是配合缩放因子做范围压缩，而不是裸着硬上。

2.2 scaling不是乘个系数那么简单

FP8最容易被低估的部分是scaling。刚接触时我也以为就是给每个张量乘一个全局系数，把最大绝对值限制在448以内。实际做进去才发现，per-tensor（整张量一个scale）在真实模型里经常不够用，因为激活分布并不均匀，一个离群特征会拉低全图的有效精度。

工程上更常见的是per-token结合per-channel的组合策略：对激活A按token维度分组，对权重B按channel维度分组，甚至把K维度切块后做更细粒度的分组量化。这样虽然会多出一些scale计算的额外开销，但能把量化误差控制在可接受范围。

用个生活化的类比：per-tensor像用一把长尺量全班身高，班里有个两米多的体育生，你为了不爆尺只能把刻度拉得很宽，结果普通同学的身高误差全变大；per-token / per-channel相当于分组测量，名正言顺地按人群特征用不同的尺，整体误差自然更低。

DeepGEMM在设计上支持这类复杂的scale组合，把缩放因子的传递和矩阵乘本身的流水线融合在一起。也就是说，scale不是前处理阶段单独算好再传进来，而是内核内部就处理掉了一部分缩放逻辑，避免额外读写和kernel launch带来的开销。

2.3 累加精度与延迟缩放：FP32寄存器里的故事

FP8能承受的精度有限，但GEMM有一个天然优势：乘法可以用8位，累加却不需要停在8位。标准的做法是内部累加器用FP32，整个K维的乘积一路累加完，最后写回时才落到目标输出精度。这样即使输入是8位，累加误差也不会在K很大时失控。

另一个关键策略是延迟缩放。你可以把它理解成“用上一次统计好的scale，先跑今天的计算”。训练场景里激活分布短时间不会突变，没必要每一步都重新算一遍scale，适当沿用上一轮的统计值就能稳住精度，同时省掉一次额外的数值统计kernel。在线缩放则适合动态范围变化很剧烈的场景，两者各有取舍。

把这些细节串在一起，你就会发现：FP8 GEMM的性能高低不只是看乘法和加法本身，还得看scale如何参与流水线、累加器放在哪一级、写回时有没有多余的格式转换。DeepGEMM在数值层面的每一个设计，最终都是为了让这些步骤尽量不打断内存与计算的重叠。

3. 双路径实现：Triton把功能跑通，CUDA把性能吃透

3.1 Triton路径先把功能跑通

Triton是一种面向GPU的Python式DSL，写起来简洁，社区普及率高，很多项目用它做快速原型。DeepGEMM保留了Triton实现，意味着普通工程师也能在不太懂CUDA细节的情况下，把核心流程读懂、改改参数、跑通自己的用例。

Triton自动化处理了很多琐事：tile划分、共享内存分配、基本的数据搬运。对多数常规形状来说，Triton路径已经能拿到可观的性能，因为编译器会帮你做不少调度决策。它的主要限制是不能像手写CUDA那样，精细控制异步流水线的每个环节，也很难针对某个具体GPU架构做非常极致的寄存器和共享内存布局优化。

所以可以把Triton路径看作“大众版”：覆盖广、上手快、维护成本低。如果你只是想把一个模型从FP16迁移到FP8，先跑Triton路径验证数值，是完全可行的一条路线。

3.2 CUDA路径把Hopper的潜力吃干净

真正要榨干Hopper架构的算力，还是得回到CUDA。Triton路径解决的是“能跑”，CUDA路径解决的是“跑满”。这部分的优化点我列一下：

• TMA批量拷贝：Hopper上的Tensor Memory Accelerator可以异步把全局内存里的矩阵块搬到共享内存，搬运过程不占用计算单元，和计算流水线重叠。
• Warp Speciaization：把warp分成producer和consumer两类，一部分只负责搬运数据，另一部分只负责矩阵乘，让smem始终有数据可用，消除流水线空泡。
• wgmma指令：warp group级别的矩阵乘指令，一个指令能完成更大块的FP8乘法，比早期mma指令更省指令数和寄存器压力。
• 动态wave调度：根据矩阵形状和SM数量计算需要多少个block，尽量让每个SM都分配到合适的任务，避免最后一个wave出现大量空闲。

这些点单独任何一个都不算新鲜，但组合起来以后，效果是明显的：GeMM的算术强度被拉高，内存延迟被隐藏，SM的利用率稳定在高位。DeepGEMM的CUDA路径就是把这些东西按FP8的特点重新组织了一遍，并且把代码控制在紧凑范围内，这是它最难得的地方。

3.3 为什么保留两条路径

我最初觉得保留两套实现是重复劳动，后来在实际工程里才理解这样做的好处。Triton路径是维护和教学的最优解，CUDA路径是极限性能的兜底选项。运行时可以根据矩阵形状、是否有Hopper特性可用、输入是否规整等条件自动选择走哪条路。

这种“双轨”思路对很多团队都有参考价值：不要一上来就追求手写CUDA，先用高抽象路径验证功能，再对热点场景用低层优化做加固。两条路径可以长期共存，而不是二选一。DeepGEMM的调度机制，本质上也是这种渐进式优化思路的一次实践。

4. 复现与实测：我看懂了为什么它敢叫高效

4.1 环境准备和三个最容易失败的地方

想复现DeepGEMM，硬件门槛比较清楚：Hopper架构的GPU，也就是通常说的H100、H200这类SM90设备。软件方面需要较新的CUDA环境，Triton也要跟得上版本，PyTorch作为上层依赖基本跑不掉。

我实际遇到过的三个失败场景，分享在这里：

1. 编译时CUDA arch信息没配对，导致报错或者生成出来的kernel根本起不了作用。解决办法是先确认torch.cuda.get_device_capability()输出，再设置对应架构，不要偷懒用默认值。
2. Triton版本太旧，JIT编译内核时直接炸。这个最隐蔽，因为安装依赖时不一定报错，但运行到特定路径就会挂。建议按仓库README的版本要求再往上对齐一点，别用过于保守的版本。
3. 环境混装导致第三方库冲突。最好单独建一个虚拟环境，避免系统里多个项目的CUDA依赖互相污染。

跑通环境后的第一步不要直接上benchmark，先跑内置的smoke test。它能一次性验证数值正确性、基本性能和当前架构兼容性，是最省时间的路径。

4.2 实测：大shape足够稳，小shape才是亮点

我复现时的观察分为两类形状。大shape，比如M=4096、N=4096、K=4096这种规整稠密GEMM，整体吞吐已经能逼近硬件理论峰值，稳定性也高，反复跑几轮波动很小。这类场景本来就是通用库的强项，DeepGEMM做得好的地方在于没有明显拖后腿。

真正让我眼前一亮的是小shape。M只有16、甚至更小时，传统库的利用率会掉得很厉害，但DeepGEMM还能保持相对高的有效吞吐，原因就在于它对wave调度和tile粒度做了针对性处理，不会让SM闲着。对推理场景来说，这就是实打实的收益，因为实际线上请求很少能凑出那种厂商标杆测试里的大batch。

4.3 性能对比的“公平陷阱”

跑benchmark时最容易犯的错误是拿FP8和FP16直接比FLOPS。FP8的峰值算力本来就翻倍了，比赢了也不代表你的模型效果就一样好。更合理的对比是看“最终推理吞吐”或者“训练迭代时间”，同时记录数值误差是否在可接受范围内。

具体的建议是：固定GPU频率，多轮warmup后取稳定值；对比时用相同的内存布局和同样的输入数据；不仅看FLOPS，还要看L2命中、kernel launch次数和峰值内存带宽。只有把这几项放在一起，才能判断一个实现是“算法层面的优势”还是“配置层面的运气”。

关于误差，我用BF16参考实现与FP8输出做比对，相对误差在可接受范围内，符合这一精度档次的预期。如果你发现误差异常大，先检查scale初始化和e4m3、e5m2是否混用，这两个问题占比最高。

5. 迁移到自己项目前，需要绕开的五个坑

5.1 形状对齐与padding策略

FP8 GEMM对形状对齐很敏感。M、N、K是16的倍数时性能最稳，达不到时要么由库内部补齐对齐逻辑，要么在外部对张量做padding。能外部pad就外部pad，因为库内部对齐会消耗共享内存和寄存器的有效利用率。

padding的时候注意输出形状也要跟着变，最后要把多余部分切掉，不要让它影响后续算子。如果你的业务场景里M经常是1、2、4这样的值，建议在上层就预先做微小batch合并处理，拉高单次GEMM的有效宽度。

5.2 FP8缩放因子的初始化细节

我最开始测试时偷懒把scale全部设为1，结果输出要么溢出成NaN，要么精度差到没法用。正确的初始化思路是根据输入张量的实际分布估算：统计最大绝对值，然后用目标格式上限除以这个值得到初始scale。

e4m3的有限上限约是448，所以初始scale可以是max_val / 448。如果清楚数据分布比较集中，还可以留一点余量，避免训练或推理过程中偶尔出现的离群值触发溢出。不同张量的scale要独立计算，不能把所有权重共用一个scale。

这里有个额外体会：scale并不是越大越好。scale过大会把有效数值压得太小，尾数位的潜力浪费，精度反而下降。比例尺要贴合数据分布，而不是一味追求“不溢出”。

5.3 集成MoE时的分组约定

MoE场景下，每个专家分到的token数量不固定，这是分段GEMM的主要复杂度来源。接DeepGEMM时，要格外注意group id和token偏移信息的传递。传错的话，轻则性能回退，重则直接算错结果，而且这类错误在数值上不一定立刻看得出来，很可能隐藏到后续层才会放大量级。

我建议先构造一个形状最简单的固定分组用例，比如两个专家、每组各8个token，逐步验证索引映射关系。跑通之后再上动态分组，动态分组的关键是确保每个专家内部的数据是连续的、可以按块搬运的布局，而不是松散排布。

5.4 混精度使用建议

FP8不是万灵药。对精度极其敏感的层，比如输出概率分布附近的计算，直接用BF16可能更稳妥。我的经验是：模型较深层的激活经过多次归一化后分布更稳定，用FP8不容易出问题；而靠近输入和输出的位置，数值变化幅度大，混合使用BF16收益更好。

实际接入时，可以按层配置精度策略：前几层用FP8加速，中间敏感层保留BF16，末尾再回FP8。这种方法能兼顾性能和稳定，也方便逐步扩大FP8覆盖范围。

再补充一个工程细节：切换不同精度时，不要只改dtype，还要重新确认对应的scale机制是否配套。FP8独有的scaling在BF16里没有，反之亦然，很多团队就是在这类“只换一半”的地方跌了交椅。

6. 一些最后想说的

花了两周时间把DeepGEMM拆完并迁移到自己项目里之后，我最深的感触是：优化不是靠一句“换上8位浮点”就能完成的系统工程。它需要你对底层硬件的数据通路有足够理解，也需要对模型本身的数值边界有耐心。把DeepGEMM的源码从头读一遍，比翻十篇综述对你的帮助都大，因为这里面每一个优化点，几乎都能对应到你未来可能在自家框架里遇到的真实瓶颈。如果条件允许，我建议你从Triton路径跑通开始，再对照CUDA路径逐行去看那些“为什么多写一行”“为什么这样调度”的决策，最后再迁移回自己的业务场景验证收益。这样一个流程走完，你得到的不仅是性能，还有一套可迁移的优化思路。