企业官网建设流程全解析

1. 为什么要在 CDNA4 上死磕 FP8 GEMM

1.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 的 tile，FP8 下矩阵核心的 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 里要调用 8x8=64 次矩阵核心指令来完成一个 128x128 的输出 tile。

参数计算过程：128 除以 16 等于 8，所以 M 方向需要 8 次；N 方向同理 8 次；K 方向 128 除以 32 等于 4，所以每个输出 tile 需要 8x8x4=256 次矩阵核心调用。这个数字看起来多，但矩阵核心是流水线执行的，实际吞吐很高。

实操心得：矩阵核心指令的延迟不低，但吞吐高。关键是让多条指令背靠背发射，不要中间插入依赖等待。我把累加器分成两组，交替使用，能有效隐藏延迟。

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_M=128，BLOCK_N=128，BLOCK_K=32。线程块大小是 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_K=32 是因为矩阵核心的 K 维度是 32，对齐后加载效率最高。THREADS=256 是平衡了占用率和并行度。

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[(k+1)%2], A_global[k+1]); load_async(B_shared[(k+1)%2], B_global[k+1]); } 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_K=64，精度比 32 差了一截。最终用 32 是精度和性能的平衡点。

第四步，对比参考实现。用 FP32 的 GEMM 做参考，逐元素对比误差。如果误差集中在某些区域，可能是那些区域的动态范围出了问题。

5.2 性能上不去的几个原因

性能问题比精度问题更难查，因为影响因素多。我整理了一个排查表：

我遇到最多的是共享内存 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，结果如下：

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 的矩阵乘法，进一步压缩显存和算力。我试过一版，注意力层的速度提升了将近一倍，精度损失在可接受范围内。