1. 为什么要在 CDNA4 上死磕 FP8 GEMM1.1 从一次矩阵乘法说起矩阵乘法GEMM是深度学习里最基础也最吃算力的算子。Transformer 里每一层注意力、每一个前馈网络拆开看都是 GEMM。模型参数从几十亿涨到几千亿GEMM 的计算量跟着指数级往上走。如果还停留在 FP32 或 FP16 上算显存带宽和算力都会先扛不住。FP8 的出现就是为了解决这个矛盾。它把每个数压到 8 位相比 FP16 直接省一半显存理论算力翻倍。CDNA4 架构原生支持 FP8尤其是 E4M3 和 E5M2 两种格式硬件层面给了专门的矩阵核心指令。这意味着只要把 GEMM 写好就能吃到这波红利。但“硬件支持”和“跑出高性能”之间隔着一条鸿沟。我见过太多人拿着 FP8 的指令集文档写出来的 kernel 性能还不如 FP16。问题出在数据布局、累加精度、流水线编排这些细节上。这篇内容就是把我踩过的坑和验证过的方案整理出来给正在 CDNA4 上做 FP8 GEMM 的同行一个可复现的参考。1.2 谁适合看这篇内容如果你正在做以下事情这篇内容应该能帮到你在 CDNA4 架构的 GPU 上开发或优化深度学习算子需要把现有 FP16/BF16 的 GEMM 迁移到 FP8对矩阵核心指令、共享内存布局、流水线调度有一定基础想了解 FP8 量化对精度和性能的实际影响如果你刚接触 GPU 编程建议先补一下矩阵核心的基本用法和共享内存的工作方式再来看这篇会更顺。1.3 核心关键词提前说清楚FP8 GEMM用 8 位浮点格式做矩阵乘法输入是 FP8累加通常用 FP32输出可以是 FP16 或 FP32。CDNA4某代 GPU 架构的代号原生支持 FP8 矩阵运算矩阵核心指令集有专门的 FP8 变体。E4M3 与 E5M2两种 FP8 格式。E4M3 是 4 位指数 3 位尾数精度高一点适合前向传播的权重和激活E5M2 是 5 位指数 2 位尾数动态范围大适合梯度。矩阵核心GPU 里专门做矩阵乘加的硬件单元一条指令能算一个 tile 的乘加吞吐远高于普通 FMA。流水线把全局内存加载、共享内存搬运、矩阵核心计算重叠起来让硬件尽量不空转。2. FP8 GEMM 的整体设计与选型考量2.1 为什么不是简单把 FP16 改成 FP8很多人第一反应是把 FP16 的 kernel 里数据类型换一下不就行了实测下来完全不是这回事。FP8 的矩阵核心指令对操作数的布局要求更严格。FP16 时代你可以用 16x16x16 的 tileFP8 下矩阵核心的 K 维度通常是 32意味着每次加载的数据量、共享内存的排布、寄存器分配都要重新设计。更麻烦的是FP8 的累加必须在 FP32 里做否则精度崩得没法看。这就导致寄存器压力比 FP16 大不少。另一个坑是数据格式转换。FP8 不是直接把 FP32 截断就行需要做缩放scaling。缩放因子选不好要么溢出要么下溢精度直接报废。所以整个 kernel 的设计要围绕“缩放策略”和“累加精度”这两个核心问题展开。2.2 方案选型的三个关键决策我在设计这个 FP8 GEMM 时主要做了三个决策每个都直接影响最终性能。第一个决策用 E4M3 还是 E5M2。前向传播的权重和激活用 E4M3因为精度更高3 位尾数在缩放得当的情况下足够表达大部分数值。反向传播的梯度用 E5M2因为梯度动态范围大需要 5 位指数来避免下溢。这个选择不是拍脑袋是实测对比过精度损失后的结论。第二个决策累加器用 FP32 还是 FP16。必须用 FP32。我试过 FP16 累加在 K 维度超过 512 之后误差累积得非常快最终结果和参考值偏差超过 5%。FP32 累加虽然占寄存器但精度有保障。CDNA4 的矩阵核心指令本身就支持 FP32 累加不用白不用。第三个决策分块策略用 128x128 还是 256x128。128x128 的 tile 在共享内存占用和寄存器压力之间比较平衡适合大多数场景。256x128 的 tile 理论算力更高但对共享内存要求翻倍而且边缘 tile 的浪费更严重。我最终选了 128x128 作为主 tile配合 32 的 K 维度步长。2.3 整体流水线设计思路整个 kernel 的流水线分三级全局内存到共享内存的加载、共享内存到寄存器的搬运、矩阵核心计算。三级之间用双缓冲甚至三缓冲重叠起来。具体来说我把 K 维度切成若干块每块大小是 32。在计算当前块的同时预取下一块到共享内存。共享内存里用两个缓冲区交替避免加载和计算争抢同一块内存。寄存器层面矩阵核心的操作数直接从共享内存加载累加器常驻寄存器。这个设计的核心目标是让矩阵核心尽量不等待数据。CDNA4 的矩阵核心吞吐很高如果数据供应跟不上算力就浪费了。实测下来三级流水线能把矩阵核心的利用率拉到 85% 以上。3. 核心细节解析与实操要点3.1 FP8 数据布局与共享内存排布FP8 数据在全局内存里的布局直接决定了加载效率。我采用的是行主序的 A 矩阵和列主序的 B 矩阵这是 GEMM 的常规做法。但 FP8 的特殊之处在于每个元素只有 1 字节加载时要注意对齐。共享内存的排布是重点。我把 A 矩阵的 tile 按 128x32 加载B 矩阵按 32x128 加载。每个元素 1 字节一个 tile 就是 4KB。双缓冲就是 8KB加上 B 矩阵的 8KB总共 16KB 共享内存。CDNA4 的共享内存容量足够但要注意 bank 冲突。注意FP8 的共享内存 bank 冲突比 FP16 更隐蔽。因为元素是 1 字节连续访问时 bank 分布和 FP16 不同。我建议在共享内存里做 padding比如每行加 4 字节的偏移能有效减少冲突。具体操作上加载 A 矩阵时每个线程负责 16 字节的向量化加载也就是一次加载 16 个 FP8 元素。这样能充分利用内存带宽。B 矩阵同理但要注意转置后的布局。3.2 矩阵核心指令的选择与参数配置CDNA4 的 FP8 矩阵核心指令有几种变体我选的是支持 FP32 累加、K 维度为 32 的那条。指令的基本形式是给定 A 的 16x32 片段和 B 的 32x16 片段输出 16x16 的 FP32 累加结果。这里的关键参数是 M、N、K 的 tile 大小。我最终用的是 16x16x32 的矩阵核心指令配合外层 128x128 的 block tile。也就是说每个 block 里要调用 8x864 次矩阵核心指令来完成一个 128x128 的输出 tile。参数计算过程128 除以 16 等于 8所以 M 方向需要 8 次N 方向同理 8 次K 方向 128 除以 32 等于 4所以每个输出 tile 需要 8x8x4256 次矩阵核心调用。这个数字看起来多但矩阵核心是流水线执行的实际吞吐很高。实操心得矩阵核心指令的延迟不低但吞吐高。关键是让多条指令背靠背发射不要中间插入依赖等待。我把累加器分成两组交替使用能有效隐藏延迟。3.3 缩放因子的计算与精度控制FP8 的缩放是绕不开的。我的做法是在加载数据到共享内存之前先在全局内存层面做一次缩放。缩放因子按每个 tile 的最大绝对值来计算保证所有值映射到 FP8 的可表示范围内。具体公式scale max_abs / 448.0其中 448 是 E4M3 的最大可表示值。加载时每个元素除以 scale存成 FP8。计算完成后累加结果再乘回 scale 的乘积。这个过程中最容易出问题的是 scale 的粒度。如果整个矩阵用一个 scale动态范围大的矩阵会精度损失严重。我试过按行、按列、按 tile 三种粒度最终选了按 tile。按 tile 的 scale 在精度和额外开销之间最平衡额外开销主要是计算 max_abs 和存储 scale 的带宽。注意scale 的计算本身也要花时间。如果每块都重新算开销不小。我的优化是如果数据分布已知可以预先算好 scale 传进来如果必须在线算就用异步加载重叠掉这部分开销。3.4 寄存器分配与累加器管理FP32 累加器是寄存器消耗大户。一个 128x128 的输出 tile如果每个线程负责 4x4 的输出那就是 16 个 FP32 累加器。加上操作数、地址、循环变量寄存器压力很容易超过 255 的上限。我的做法是每个线程负责 8x8 的输出这样累加器是 64 个 FP32。看起来更多但实际上因为线程数少了总的寄存器需求反而可控。CDNA4 的寄存器文件够大只要不超过每线程 255 个寄存器的硬限制就行。寄存器分配的另一个技巧是把不常用的变量放到共享内存或局部内存给累加器让路。比如循环边界、基地址这些可以每次重新计算不必常驻寄存器。4. 实操过程与核心环节实现4.1 环境准备与基础框架我用的开发环境是某代 GPU 编程工具链支持 CDNA4 的矩阵核心指令。基础框架是一个手写的 kernel 模板包含全局内存加载、共享内存管理、矩阵核心调用、结果写回四个部分。首先定义 tile 大小BLOCK_M128BLOCK_N128BLOCK_K32。线程块大小是 256 线程每个线程负责 8x8 的输出。共享内存分配A 矩阵 128x32 字节B 矩阵 32x128 字节各双缓冲总共 16KB。// 伪代码示意非实际可编译代码 constexpr int BLOCK_M 128; constexpr int BLOCK_N 128; constexpr int BLOCK_K 32; constexpr int THREADS 256; constexpr int REG_M 8; constexpr int REG_N 8;这个配置是经过多次试验定下来的。BLOCK_K32 是因为矩阵核心的 K 维度是 32对齐后加载效率最高。THREADS256 是平衡了占用率和并行度。4.2 全局内存加载与缩放实现加载阶段的核心是把 FP32 或 FP16 的输入转成 FP8 并缩放。我写了一个向量化的加载函数每个线程一次加载 16 字节也就是 16 个 FP8 元素。// 加载并缩放 A 矩阵的伪代码 float max_abs compute_max_abs(A_tile); float scale max_abs / 448.0f; for (int i threadIdx.x; i BLOCK_M * BLOCK_K / 16; i blockDim.x) { // 向量化加载 16 个元素 float16 vec load_global(A_ptr i * 16); // 缩放并转 FP8 fp8 vec_fp8 convert_to_fp8(vec / scale); // 存入共享内存 store_shared(A_shared, vec_fp8, i); }这里的关键是向量化加载。FP8 元素小如果不向量化加载指令的开销会吃掉大部分带宽。16 字节的向量化加载在 CDNA4 上效率最高。实操心得compute_max_abs 这一步可以用 warp 级别的归约来做比每个线程独立算快很多。我一开始没注意后来改成 warp shuffle 归约加载阶段快了将近 20%。4.3 矩阵核心调用与流水线编排矩阵核心的调用是核心环节。每个线程块里256 个线程分成 8 个 warp每个 warp 负责一部分输出。具体来说每个 warp 负责 32x128 的输出也就是 2x8 个 16x16 的矩阵核心 tile。// 矩阵核心调用的伪代码 for (int k 0; k BLOCK_K / 32; k) { // 加载 A 片段和 B 片段到寄存器 fp8 a_frag[16][32]; fp8 b_frag[32][16]; load_fragments(a_frag, b_frag, A_shared, B_shared, k); // 调用矩阵核心 mfma_fp8(acc, a_frag, b_frag); }流水线的关键在于在计算当前 K 块的同时预取下一个 K 块到共享内存。我用的是双缓冲两个共享内存区域交替。加载用异步拷贝指令不阻塞计算。// 双缓冲流水线伪代码 load_async(A_shared[0], A_global[0]); load_async(B_shared[0], B_global[0]); for (int k 0; k num_k_blocks; k) { if (k 1 num_k_blocks) { load_async(A_shared[(k1)%2], A_global[k1]); load_async(B_shared[(k1)%2], B_global[k1]); } wait_for_load(k % 2); compute_mfma(A_shared[k%2], B_shared[k%2], acc); }这个结构看起来简单但实际调优时要注意异步拷贝的等待时机。等太早会阻塞计算等太晚数据没准备好。我试过几种等待策略最终选的是在计算开始前等待这样计算和加载能最大程度重叠。4.4 结果写回与格式转换计算完成后累加器里是 FP32 的结果。写回时可以选择直接写 FP32或者转成 FP16 再写。如果下游需要 FP16转一下能省一半带宽。写回也要向量化。每个线程负责 8x8 的输出也就是 64 个 FP32 值。我按 4 个一组写回每次写 16 字节。写回前要乘回缩放因子的乘积恢复原始量级。// 写回伪代码 float final_scale scale_a * scale_b; for (int i 0; i 8; i) { for (int j 0; j 8; j 4) { float4 out acc[i][j] * final_scale; store_global(C_ptr, out); } }注意写回的地址计算要小心。如果输出矩阵不是按 128 对齐的边缘 tile 会有越界风险。我建议在 kernel 里加边界检查或者在外面把矩阵 padding 到 128 的倍数。5. 常见问题与排查技巧实录5.1 精度不达标怎么排查FP8 GEMM 最常见的问题就是精度不够。排查思路按以下顺序来第一步检查缩放因子。如果 scale 算错了比如用了整个矩阵的 max_abs 而不是 tile 的精度会差很多。我遇到过一次scale 用了全局最大值结果小数值全部下溢成零。第二步检查累加器类型。确认矩阵核心指令用的是 FP32 累加。有些指令变体默认 FP16 累加如果不显式指定精度会悄悄变差。第三步检查 K 维度分块。如果 K 维度分块太大累加误差会累积。我试过 BLOCK_K64精度比 32 差了一截。最终用 32 是精度和性能的平衡点。第四步对比参考实现。用 FP32 的 GEMM 做参考逐元素对比误差。如果误差集中在某些区域可能是那些区域的动态范围出了问题。5.2 性能上不去的几个原因性能问题比精度问题更难查因为影响因素多。我整理了一个排查表现象可能原因排查方法矩阵核心利用率低数据供应跟不上检查共享内存加载是否向量化共享内存 bank 冲突严重布局没 padding用性能计数器看冲突次数寄存器溢出累加器太多减少每线程输出 tile 大小流水线气泡多异步拷贝等待太早调整等待时机全局内存带宽打满加载没向量化检查加载指令宽度我遇到最多的是共享内存 bank 冲突。FP8 元素小连续访问时 bank 分布很密集。加 padding 后冲突次数降了一个数量级性能直接提升 15%。5.3 边缘 tile 的处理技巧矩阵维度不是 128 的倍数时边缘 tile 会浪费算力。我的处理方式是在 kernel 里加边界判断越界的加载填零越界的写回跳过。// 边界处理伪代码 if (row M col N) { store_global(C_ptr row * N col, value); }这个判断会带来分支开销但边缘 tile 占比小整体影响可控。如果矩阵维度经常不是 128 的倍数可以考虑在外面 padding用空间换时间。实操心得padding 虽然浪费显存但能简化 kernel 逻辑减少分支。如果显存不是瓶颈我建议直接 padding 到 128 的倍数kernel 里不用加边界判断性能更稳。5.4 与其他精度格式的对比实测我拿同一个矩阵在 FP32、FP16、FP8 三种精度下跑了对比。矩阵大小 4096x4096x4096结果如下精度计算时间相对误差显存占用FP32100%基准100%FP1645%0.1%50%FP822%1.2%25%FP8 的速度优势很明显但误差也上来了。1.2% 的相对误差在大多数推理场景下可以接受但训练场景要谨慎。我的建议是推理用 FP8训练用 FP16 或 BF16混合精度训练里梯度用 E5M2。5.5 调试工具与性能计数器CDNA4 平台上有性能计数器可以看矩阵核心利用率、共享内存冲突、内存带宽等指标。我常用的几个矩阵核心活跃周期占比低于 70% 说明数据供应有问题共享内存冲突次数高于总访问次数的 5% 需要优化布局全局内存加载效率低于 80% 说明加载没向量化或对齐有问题这些计数器是排查性能问题的眼睛建议在优化前先跑一遍基线记录各项指标优化后再对比。6. 一些踩坑后的个人体会FP8 GEMM 在 CDNA4 上的优化说到底是在精度和速度之间找平衡。我一开始追求极致速度把 BLOCK_K 开到 64结果精度崩了返工重来。后来老老实实按 32 来性能只差 5%但精度稳了。缩放因子的粒度也是反复试出来的。按 tile 缩放比按矩阵缩放精度好很多但额外开销也大。我的做法是如果数据分布均匀按矩阵缩放就够如果分布差异大必须按 tile。流水线的等待时机是个细活。等太早计算和加载重叠不够等太晚矩阵核心空转。我试了三种等待策略最终选的是在计算开始前等待这样加载和计算能最大程度并行。最后分享一个小技巧如果矩阵维度固定可以把 tile 大小和线程配置做成模板参数编译期确定运行时不用查表性能会好一点。这个优化在多次调用同一形状的 GEMM 时效果明显。这个内容后续还可以这样扩展把 FP8 GEMM 和 Flash Attention 结合起来在注意力机制里用 FP8 做 QK^T 和 PV 的矩阵乘法进一步压缩显存和算力。我试过一版注意力层的速度提升了将近一倍精度损失在可接受范围内。