做训练的人大概都体会过这么一种拧巴模型的参数量越来越大显存越占越满算力反而经常喂不饱。我最近在一个基于CDNA4 GPU的测试平台上做矩阵乘法优化第一眼看到FP8 GEMM的收益时多少有点意外——同样的一个GEMM输入从16位浮点换成8位浮点跑出来的吞吐直接往上跳了一大截显存占用也明显下降。这篇东西主要就是记录我在CDNA4上把一个FP8 GEMM从“能算”调到“算得飞快”的全过程包括格式选型、量化缩放、矩阵指令的分块策略以及一堆只有跑过才会知道的坑。适合写过CUDA或HIP内核、想尝试FP8的人也适合还在纠结“FP8到底该怎么落进框架里”的读者。1. FP8 GEMM解决的真实瓶颈带宽、显存、算力三方互相打架1.1 一个最简单的线性层数据要搬多少先别急着把FP8当成“算得更快”的黑科技我自己的体会是它首先解决的是数据搬运问题。一个标准的 GEMM形状是 M×N×K也就是 C[M,N] A[M,K] × B[K,N]浮点运算次数是 2MNK。很多人只看算力峰值却忘了喂给计算单元的数据也要从显存里挪过来。算力越高的卡带宽越容易成为瓶颈这句话在CDNA4上体现得尤其明显。举个例子MNK4096 的方阵乘方阵。读入 A 和 B 两个输入矩阵如果是FP16每个元素2字节A加B总共需要读 2×4096×4096×2 ≈ 128MB换成FP8每个元素1字节同样的矩阵只需要64MB。别小看这一倍的差距在现代GPU架构里从全局内存搬数据的时间经常占掉整个kernel运行时间的一半以上。也就是说FP8 GEMM最直观的收益不是算得快而是让等数据的时间变短了。那么这个规模的计算量是多少2×4096^3 ≈ 137.4 GFLOP。假设某个CDNA4实例的FP16矩阵峰值是标称的1倍FP8峰值是2倍那么在FP16下就算维持400 TFLOPS以上的实际吞吐也要跑将近0.4毫秒但FP8只需要原来一半以上的带宽配合更高的矩阵单元吞吐整体时间能压得更低。后面我会给一份实测数据大体趋势就是如此。1.2 算力受限和带宽受限收益逻辑完全不同需要补充一个容易误判的点FP8不是在所有GEMM上都按同样的比例加速。如果GEMM的M、N、K都很大比如4096以上计算量足够大算力本身是主约束那FP8能吃到峰值翻倍的红利但如果矩阵尺寸偏小比如K只有128整个kernel大部分时间花在启动开销和内存访问上这时候FP8的收益主要来自带宽减半速度提升可能反而不如大矩阵那么夸张。我自己调整内核的时候会把GEMM分成两类来看算力受限型M、N、K都在2048以上单次乘加密度高FP8可以明显提升吞吐。这类形状在LLM的MLP层里最常见。带宽受限型矩阵很小或者M/N严重不均衡比如M128、N4096、K256时间主要消耗在数据供给上。FP8依然有效果但优化重点应该放在减少全局内存访问次数、复用LDS数据上而不是单纯换一种窄位宽格式。还有一个隐藏收益容易被忽略FP8能把更多激活值放进显存。训练大模型的时候前向传播的激活经常要留到反向计算梯度用激活量越大显存越紧张。把激活从FP16切成FP8等于省出一半显存可能让你少做很多激活重计算。所以FP8 GEMM在训练侧的价值很多时候是计算、显存、带宽三重收益叠加的。2. FP8是两套格式不是一套E4M3、E5M2与量化缩放2.1 位宽相同脾气完全不同FP8不是一个统一的格式这一点如果不先搞清楚后面调试精度问题会非常痛苦。8位浮点一共分两种E4M3和E5M2。名字里的E和M分别代表指数位和尾数位。我用一个表把两种格式的关键差异列出来方便对比属性E4M3E5M2指数位45尾数位32最大有限值44857344最小正常数约2^-6≈0.0156约2^-14≈6.1e-5数值精度相对较高相对较低典型用途激活、权重梯度、中间结果粗看指数位越多、范围越大E5M2确实能表示到几万的数量级损失函数里偶尔冒出一个特别大的梯度也不至于直接溢出。但代价是尾数只有2位相对误差大得多。E4M3虽然范围只到448超过就得靠缩放因子拉回来但尾数有3位对于激活和权重这类动态范围相对可控的张量更友好。所以不要“整个模型全用E4M3”或者“全用E5M2”比较好的做法是分对象选格式。我习惯的基准确认是前向传播的激活和权重用E4M3反向传播的梯度用E5M2因为梯度的动态范围真的大。后面第5节我会细说。2.2 缩放因子选多细精度和成本差很远FP8的8位空间不管怎么切也没法直接覆盖FP32的动态范围。所以实际使用FP8 GEMM前必须先把输入张量量化到8位可表示的区间这个过程需要一个缩放因子。最简单的是per-tensor对整个矩阵找一个最大绝对值然后决定用什么比例把最大的数映射到448或者57344附近。公式很直观拿到张量里绝对值的最大值 amax用scale amax / max_val其中 max_val 对E4M3是448、对E5M2是57344。然后每个原始浮点数x都变成x_q clamp(round(x * scale), -max_val, max_val)。反量化的时候再用x ≈ x_q / scale还原。per-tensor的优点是计算便宜、实现简单整个矩阵只需要一个amax统计但缺点也很明显如果矩阵内部不同行或不同列的数值分布差异很大一个全局缩放因子会把部分区间的细节压扁误差明显变大。所以CDNA4这类平台做高性能GEMM时通常会升级成更细的粒度缩放粒度统计范围精度表现开销per-tensor整个矩阵一个scale实现最简单分布差异大时误差高最低per-row / per-columnA按每行、B按每列统计对权重和激活分布不均更友好中等per-block按小块统计更细但量化元数据多较高我自己的经验是在训练场景里至少做到per-row/per-column。A矩阵按行量化B矩阵按列量化这样的话每个输出元素对应的是“一个行缩放因子乘以一个列缩放因子”最终结果的缩放可以在累加完成后统一定标实现复杂度只比per-tensor多一点但误差改善非常明显。这里有个坑必须提醒如果某个行或某个mini-block恰好是全零amax算出来是0缩放因子也是0后续所有量化结果全是0反向传播推导时很容易炸出NaN。我一般会给scale一个下限保护比如if (amax 1e-8f) scale 1.0f;保证全零区域也能正常反量化。2.3 舍入模式RNE与随机舍入量化本质上是给每个数找一个最近的8位浮点数表示。默认的舍入模式是round-to-nearest-even也就是四舍五入到最近的尾数碰到正好在中间的情况取偶数尾数。这个模式对大多数场景足够但在长迭代训练中系统性的舍入偏差会积累。如果发现训练曲线比预期差一点可以先试试随机舍入。原理很朴素对于一个原始值如果它落在两个相邻FP8表示数之间那么以一定的概率向上一档舍入概率大小正比于它离上一档的距离。数学上这样能让量化误差的期望变成零防止误差单向累积。代价是需要生成随机数通常用一个简单的伪随机数函数就能实现性能影响可控。我在调试收敛问题时切换随机舍入往往是改善最明显的一步。不过随机舍入增加了数值结果的不确定性调试时对比误差会有点痛苦。我的做法是训练收敛评估用随机舍入做GEMM正确性验证时切回RNE否则每次输出结果轻微波动你很难判断到底是代码写错了还是随机扰动导致的。3. CDNA4的矩阵执行单元与两类开发入口3.1 MFMA指令是怎么把4096次乘加一口气做完的CDNA4属于计算卡架构所以它和普通显卡不一样里面有专门的矩阵执行单元。这类矩阵指令在CDNA系列上通常叫MFMA指令族一条指令就能完成一个16×16×16的矩阵乘加A是16×16B是16×16乘出来累加到一个FP32的16×16结果块。如果算乘加次数16×16×16 4096次乘加一次性完成。可以这么类比普通SIMD指令就像一个人一口气算4个、8个或者16个独立的乘加矩阵指令则是60多个人排成一个波阵面wavefront每人负责结果矩阵里的几个位置用特殊的数据通路协作完成一整块矩阵运算。这样做的效果是每个周期内完成的浮点运算数量大幅提升。关键点在于MFMA指令要求数据在寄存器里排列成它能识别的形状而不是随便给它两个数组就行。一旦数据摆放不对要么结果错误要么性能极差。这也是我后来发现“分块策略决定成败”的原因只有通过合理的寄存器分块让每次MFMA都操作自己需要的16×16小块矩阵单元才能被喂饱。3.2 快速路线官方矩阵库一把梭如果你只是想让FP8 GEMM先跑起来、拿到一个正确的基准结果最快的路径是用HIP生态里的BLAS类矩阵库它提供了类似传统BLAS的接口专门处理GEMM和各种混合精度。只需要把FP8的输入矩阵指针传进去配置好缩放因子和输出累加器类型库就会自动调度底层矩阵指令。对于多数人来说这是第一步最省事的选择。我最初就是在项目里用库函数跑通了一个FP8版本的矩阵乘验证收敛性和数值误差。优点很明显不用关心分块、不用关心LDS库已经在大部分常见形状上做了调优。但它也不是万能当你需要把量化缩放因子融合进GEMM内部、或者做一些非标准的布局转换时库函数往往不够灵活。性能调优到一定程度你会发现手写kernel的收益反而更大。3.3 进阶路线手写HIP内核如果你决定自己写就绕不开CDNA4的矩阵指令。HIP里可以写内联汇编或者在编译器中直接使用对应的矩阵内建函数。以伪代码示意的话一条FP8 MFMA大约长这样// 伪代码具体助记符以当前指令集手册为准 // 在单个wavefront内发起16x16x16的FP8矩阵乘加 // a_fp8 和 b_fp8 是寄存器里的16x16数据块c_fp32 是16x16累加块 asm volatile( v_mfma_f32_16x16x16_fp8 %0, %1, %2, %0 : v(c_fp32) : v(a_fp8), v(b_fp8) );这里的输入必须是已经量化好的FP8数据。注意FP8的数据在内存/寄存器里也是一个字节一个字节地存着但它不是普通的int8。直接把两个int8矩阵塞进MFMA指令得到的结果和对不上。通常需要先用库或者编译器提供的数据转换路径把FP32/FP16值量化成对应的FP8位格式中间不能跳过转换函数。为什么要折腾手写因为我后面做了几个融合操作把缩放因子的乘除直接嵌到累加器输出阶段把激活函数比如GELU和量化合并到回写全局内存之前。这些操作一旦放进一个kernel又能省掉一次甚至两次完整的全局内存读写性能比库函数更进一步。FD矩阵库能帮你解决80%的问题剩下20%的定制优化才是FP8 GEMM竞争力的来源。4. 分块参数与三段流水让FP8矩阵指令喂满4.1 BM、BN、BK到底怎么选写GEMM内核绕不开三个维度上的分块M方向上每个block算多少行BMN方向上每个block算多少列BNK方向上每次迭代处理多长BK。这三个参数直接影响LDS用量、寄存器用量和MFMA指令的复用率。我一上来直接参考比较保守的配置BM128BN128BK32。一个block负责输出128×128的C子块整个过程沿着K方向循环。为什么是128因为一个wavefront是64个lane处理64×64的输出子块比较自然128×128的block可以交给4个wavefront每个wavefront负责自己的64×64区域。寄存器里每个wavefront要保存16个FP32累加块压力适中。BK的选择其实挺讲究。BK太小比如8或16每次加载的数据太少流水线切换和同步开销占比高BK太大比如128LDS存储占用暴涨能并行跑的block变少占用率下降。我在CDNA4上的合理区间是BK32到64具体看矩阵形状和显存带宽。一个粗略的配置参考如下参数推荐起点说明BM128 或 256越大复用越充分但block内线程增多同步复杂BN128通常与BM保持一致波阵面规划更简单BK32 或 64太大LDS压力大太小流水收益差block内wave数4256线程wavefront是64的倍数别把wave拆碎选参数的时候最好用性能分析工具看一眼矩阵单元利用率。如果不到50%通常不是算力不够而是分块导致的数据供给跟不上只有利用率到了70%以上才是真正开始吃硬件红利。4.2 数据从全局到LDS再到累加器的流水GEMM kernel内部要处理三处数据A和B的大块从全局内存读进LDS共享内存这一步以block为单位大家一起搬。LDS里的数据再按MFMA需要的16×16小块分散到各线程的VGPR寄存器。MFMA指令读取寄存器中的FP8数据块把结果累加到FP32寄存器。第一步和第二步之所以不能合并是因为全局内存访问的粒度比较大而且不同线程需要的数据有重叠。先把数据以连续的方式搬进LDS再做LDS到寄存器的广播能显著提高访存效率。这跟盖房子先把材料搬到工地再分发给各工人的逻辑是一样的。但只是分三步还不够。MFMA计算很快如果每轮循环都老老实实地“先加载再寄存器再计算”内存等待就会拖着矩阵单元一起慢下来。解决办法是流水线化当前这一轮用LDS里的数据计算时下一轮需要的数据已经在后台从全局内存往LDS里搬了。4.3 主循环骨架下面是一个简化版本的主循环不考虑边界填充和所有同步细节重点展示双缓冲和流水结构// 示意伪代码省略了block索引、边界处理和同步barrier float C_local[16][16]; // 累加器留在线程寄存器 __shared__ fp8 sA[2][BM][BK]; __shared__ fp8 sB[2][BK][BN]; // 预加载第0个k块 load_tile_to_lds(0, sA[0], sB[0]); for (int k 0; k K / BK - 1; k) { // 提前加载第k1块与当前计算重叠 load_tile_to_lds(k 1, sA[(k 1) % 2], sB[(k 1) % 2]); barrier(); // 从LDS读入当前块到寄存器 block_sync_and_convert(sA[k % 2], sB[k % 2], a_reg, b_reg); // 用两条FP8 MFMA完成当前BK段的乘加 mfma_fp8_16x16x16(C_local, a_reg_0, b_reg_0); mfma_fp8_16x16x16(C_local, a_reg_1, b_reg_1); barrier(); }这里面有个细节load_tile_to_lds用于把全局内存的连续数据搬到LDS搬的同时计算单元还在做上一轮的MFMA。双缓冲使用k % 2来切换两套LDS空间避免上一轮的数据还没读完就被覆盖。这是FP8 GEMM性能提升的核心动作之一比单纯把数据从FP16换成FP8带来的速度提升还要可观。4.4 LDS布局与bank冲突数据进了LDS不代表万事大吉。LDS为了高带宽内部划分成若干个bank同一时钟周期内如果多个线程访问的地址落在同一个bank的不同位置就会发生冲突导致访问被串行化。这种冲突在FP8 GEMM里很容易被忽视因为数据从16位变成8位后地址跨度变小线程间更容易踩到同一bank。处理办法很土但很有效在LDS数组的每一行末尾补一个元素错开地址对齐。比如原本sA[row][col]的布局在存储时改成sA[row][BK 1]每一行多出一个字节的偏移让不同行数据错开到不同的bank上。代价是浪费一点LDS空间但能够换来稳定的访存带宽非常值得。另外一个容易犯错的地方是很多编译器提供的批量化内存拷贝指令会把连续数据一次性搬到LDS看起来高效但底层依然按bank访问。如果bank冲突严重理论带宽再高也白搭。我的习惯是在调优阶段用性能分析工具专门看LDS bank conflict计数器看到冲突占比降低到很小时访存这块才算过关。5. 数值稳定性FP8跑得快能不能跑得对5.1 累加器留在FP32不要手贱改成FP8FP8 GEMM里最不该动的一条原则是中间累加器用FP32。MFMA指令天然支持把FP8输入乘出来累加到FP32寄存器整个K维度循环中累加过程都在32位下进行。如果为了省寄存器把累加器改成FP8等于每个MFMA之后立即损失精度最终结果的误差会大到不可用。这就像往钱包里数硬币每一枚硬币本身有误差数一百次之后误差可能叠加得面目全非。FP32累加器保证的是误差只来自输入的两次量化而不是来自每一次矩阵乘加运算本身。不过请注意累加顺序不同结果也会有微小差异。FP32本身不是精确运算不同线程、不同block对同一组数据做不同顺序的累加最后一位结果可能差一点点。这是正常现象不要当成bug来查。调试时我固定使用FP64参考实现做对照误差在合理范围内就行不是非得逐位一致。5.2 混合格式与训练收敛的经验如果要做训练我建议不要整个模型一概而论地用E4M3。实际经验是分角色选择格式运算角色推荐格式原因前向激活E4M3尾数位多保留更多有效数字前向权重E4M3权重分布相对稳定动态范围可控反向激活梯度E5M2梯度容易有极端值需要更大动态范围反向权重梯度E5M2梯度变化剧烈优先防溢出初始版本可以把所有张量都切成E4M3跑出来的loss通常会正常下降但和FP16基准相比可能偶尔出现某一步loss突然变大甚至NaN。这个问题大概率出在梯度上梯度分布里偶发巨大值被E4M3的448上限截断后反量化出来仍然失真反向传播越传越大。切到E5M2后动态范围扩大大部分NaN问题会消失。另外别把优化器状态也量化到FP8。优化器里的一阶、二阶动量需要高精度累积FP8存上去会拖累收敛。我自己的做法是优化器保持FP32只在GEMM输入和输出上使用FP8这样既拿到了计算加速又避免优化器精度崩掉。5.3 对照误差的调试流程踩过一次“看起来跑得快、实际结果全错”的坑之后我养成了一个固定调试流程。核心思路是先做小规模、可复现的FP64参考再逐步放大而不是一上来就调4096×4096的大GEMM。具体步骤选择 MNK64 的极小GEMM用FP64算出一个参考C_ref。同一数据用FP8量化后走完整kernel得到C_fp8。对比最大绝对误差、平均相对误差和每行误差分布。如果误差明显偏大先检查缩放因子打印每个A行、B列的amax看是否有全零行或异常大的值。确认缩放没问题后切换到E5M2或随机舍入观察误差是否下降。都正常再逐渐放大到256、1024最后到实际生产规模。在FP8 GEMM这种场景里相对FP32的误差在1e-2到1e-3量级可以接受。如果误差到了1e-1甚至更大基本不是舍入误差问题而是缩放因子、数据布局或者MFMA输入顺序错误。那时候别急着调优性能先回去把数值链路梳理清楚。6. 实测性能与踩坑复盘6.1 计时方法和指标采集性能调优离开测量就是空谈。我的做法是先用HIP的事件API做粗粒度计时比如两次事件之间经过多少毫秒跑几次取中位数。注意要在计时前做warmup连续把同一个kernel跑十几遍让GPU时钟稳定、缓存预热不然第一次跑出来的时间一定偏慢。小时钟只能看整体想看矩阵单元利用率还是要上性能分析工具。CDNA4平台支持从硬件计数器读MFMA相关指标至少需要关注三件事矩阵单元利用率越高越好低于60%优先查数据供给。全局内存命中率/L2命中率命中率低说明分块设计有问题。LDS bank conflict计数冲突多则按第4节的方式调整布局。我在调优时经常发现时间参数看起来差不多但MFMA利用率差出20个百分点。这种差距在代码上看不出来只能靠计数器定位。6.2 几个关键规模的实测数据以下是我在同一款CDNA4测试环境下用FP16和FP8分别跑一组GEMM得到的大致相对数值单卡FP8峰值约为FP16峰值两倍的前提下这篇数据主要是展示趋势不代表所有型号都精确一致GEMM规模FP16耗时FP8耗时实测提速相对FP64参考误差4096×4096×40960.46 ms0.28 ms1.64x约3e-38192×8192×81923.52 ms1.91 ms1.84x约4e-31024×4096×10240.12 ms0.065 ms1.85x约1e-2可以看到规模越大FP8对FP16的优势越明显因为算力受限型GEMM能吃到峰值翻倍的红利。小规模GEMM提速也在但没那么夸张主要收益是带宽减少。误差方面规模越大累计误差略微上升但都在可接受范围。需要说明的是这几个数字不是跑一次就写下来的。我第一次跑8192尺寸时速度反而没达到预期查了半天发现是全局内存的访问模式不够连续分块顺序不对导致很多显存带宽浪费在低效访问上。把BM/BN对调之后速度立刻上来了。所以数字背后的调参过程比数字本身更有价值。6.3 踩坑清单挑几个我真正花时间排过的坑按优先级列出来K不是16的倍数MFMA的K维度天然是16如果实际K1000这种尺寸必须先把K填充到1008或1024再用掩码/边界处理把多余部分排除掉否则程序直接算错或者越界。B矩阵的转置问题FP8矩阵指令对B操作数的数据Layout有要求不是随便给个行主序就行。我踩过“两个库函数矩阵乘都对不上号”的坑最后发现是B矩阵多了一次转置。写代码前先查明白当前指令期望哪种布局。全零行的缩放保护全零tile如果没做保护scale0反量化全零反向传播梯度会出Inf/NaN。解决见第2节的if amax eps保护。库函数转写手写时的数据格式库函数内部会把某个矩阵偷偷转置或者重排你手写时不能直接沿用它的内存布局需要重新按MFMA指令的输入规范组织数据。编译器版本差异不同编译器对FP8内建函数的支持程度不一样。我从旧版本升级工具链后发现同样的内联汇编语法不能编译最后查文档确认助记符变了。如果遇到编译报错优先怀疑指令助记符而不是代码算法。尾数填充垃圾数据从LDS往寄存器搬运FP8数据时如果一次性搬运了4字节向量而实际尾数不足4字节高字节里可能掺着上一块留下的垃圾位。我在调小矩阵时碰到过类似问题表现为偶尔几个元素异常大。处理办法是装载时先清空再按位拼接。6.4 后续还能继续压榨什么如果主循环已经跑顺还有两个方向可以把性能再往上推一层。一个是Split-K当K维度非常大时把K切成几段让多个block并行计算各自的局部累加最后通过原子操作或者独立reduce kernel把部分和合到一起。这个策略还能顺手解决K方向带来的尾部填充浪费问题。另一个是Persistent Kernel传统GEMM按block网格启动一个block算完自己的tile就结束Persistent Kernel只启动固定数量的block常驻GPU不断循环领取新的tile任务。这样可以避免反复启动kernel的开销也能让L2缓存命中率更高。我自己的最终版本就是在这两个方向上继续压榨的效果主要体现在中小规模GEMM和混合形状的工作负载上。最后再补一句个人体会FP8 GEMM不是简单的“换个存储格式”它横跨数值分析、硬件指令、内存系统三个层面。我踩过最多的坑不是指令不会写而是信息差——不知道硬件期望的布局不知道格式之间的差异不知道缩放因子在反向传播里是怎么作用的。如果你正打算在CDNA4上动手做这件事建议先小规模搭建一个数值对照测试台再逐步放开形状和性能优化这条路会比一上来就调大卡稳定得多。