引言 在现代 AI 推理与高性能计算领域,矩阵乘加运算(Matrix Fused-Multiply-Add,MFMA)构成了绝大多数计算负载的核心。AMD 在 CDNA 系列 GPU 中引入了专用的 Matrix Cores,通过硬件级别的矩阵乘加指令大幅提升混合精度计算吞吐,尤其对低精度数据类型如 FP16、FP8、FP6、FP4 的支持,使得算力与能效比获得跨代提升。本文面向开发者与工程团队,系统梳理在 CDNA3 与 CDNA4 架构上编写 Matrix Core 程序的要点,包括浮点表示、MFMA 指令语义、LLVM/ HIP 的编译器内建函数(intrinsics)、常见数据布局与实战优化建议,旨在帮助读者快速上手并获得可观的性能收益。 Matrix Cores 的价值与架构差异 Matrix Cores 专门针对大规模矩阵乘加进行了硬件级优化。与逐元素的标量乘加相比,Matrix Cores 能在单条指令内完成多个乘累加运算,其带来的性能提升在混合精度场景下尤其显著。在 CDNA3(如 Instinct MI325X)中,使用 FP16 输入矩阵即可带来近 8 倍的峰值算力提升,FP8 时更能接近 16 倍。
CDNA4(如 Instinct MI355X)进一步扩展了矩阵维度支持、引入更多低精度格式,并在架构级别提升了吞吐,使得 FP16 与 FP8 的峰值分别能到达比 CDNA3 更高的倍数,且新增 FP6、FP4 等极低精度格式,理论上能实现更大的性能跃升。 低精度浮点类型要点 理解低精度浮点的位域划分对程序正确性与数值稳定性至关重要。浮点数由符号位、指数位与尾数(或称为显式小数位)构成。指数位决定数值范围,尾数位决定精度。FP16(E5M10)广为人知,FP8 则分为 E4M3 与 E5M2 两类,且存在若干变体(例如 E4M3FN 与 E4M3FNUZ)在指数偏置与特殊值(零、NaN、无穷大)的表示上有所不同。CDNA3 通常采用 FNUZ 变体,而 CDNA4 支持 OCP 标准的 E4M3FN 与 E5M2。
FP6 有 E2M3 与 E3M2(后者类似 BF6),FP4 的常见表示为 E2M1。另有特殊 8 位格式 E8M0 并不表示常规元素数据,而用于作为块级缩放(block exponent scaling)的尺度因子,其数值按 2^(exponent - bias) 进行解释。 选择合适的低精度格式需要在吞吐与数值范围之间权衡。指数位更宽的格式(例如 E5M2)可以覆盖更大的动态范围,适用于权重与激活存在大幅差异的场景。尾数位更多则提升表示精度,减少量化误差。CDNA4 对多种组合的原生支持,让工程师可以在精度与性能间灵活取舍。
MFMA 指令族与数据布局概念 AMD 的 MFMA 指令以操作尺寸(MxNxK)与输入输出类型为特征,典型语义为 D := A * B + C(或原地累加 C := A * B + C)。指令粒度通常为 wavefront 级别(CDNA 的 wavefront 大小为 64 个线程),指令执行时需要将 A、B、C、D 的子块分发到每个线程寄存器中。不同的 MFMA 指令对线程间数据映射与寄存器布局有严格规定,开发者必须按照 ISA 的数据布局去组织全局内存到寄存器的加载顺序,否则内建函数的行为将无法匹配期望的数学运算。 以若干代表性例子说明,诸如 32x32x2 的 FP32 MFMA 在单个 wavefront 内每个线程只需保存 A 与 B 的少量元素,但要保存输出矩阵 C 的更多元素;而较大的 32x32x16 的 FP8 情形中每个线程会保存更多的 A/B 元素与相对固定数量的输出累加子块。CDNA4 新增的 block-scaled MFMA 可以在点积后乘以每个块的尺度(由 E8M0 表示),这对动态范围有限的超低精度格式(例如 FP4)尤其重要。 编译器内建函数与 HIP 编程实践 在 HIP/ROCm 生态中,LLVM 提供了一组内建函数来直接映射到硬件 MFMA 指令,函数命名约定包含输出类型、矩阵尺寸与输入类型描述,例如 __builtin_amdgcn_mfma_f32_16x16x16f16 表示以 FP16 输入、FP32 输出执行 16x16x16 MFMA。
CDNA4 对于带缩放的 MFMA 有专门的内建函数,如 __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4,参数中包含输入类型标识、OPSEL 码以及尺度向量(E8M0 类型)。需要注意内建函数的参数类型约束,例如某些 MFMA 内建函数期望 A、B 操作数以 256 位向量或特定整型传入,而尺度则通常以 8 位编码的 E8M0 值传入。 在设备代码层面,使用合适的存储类型与向量化类型非常关键。HIP 提供了如 __hip_fp8_storage_t、__amd_fp8_storage_t、__amd_fp4x2_storage_t 等用于表示低精度存储单元的类型定义。因为内存访问粒度最低为字节级,需要用位打包与解包函数来从字节中提取单个 FP4 或将两个 FP4 打包成一个字节以便加载。内建函数通常期望累加向量先置零(表示纯乘积求和),或者传入已有累加值以启用原地累加。
数据装载顺序与寄存器布局往往是可移植性与性能的关键。访问模式应尽量保证 coalesced 的 global memory 读取,在可能时优先使用 s_load 或 copy 程序将数据从全局内存搬入寄存器,然后由内建 MFMA 指令在 wavefront 内完成计算与累加。由于 MFMA 是同步到 wavefront 的原子级协作操作,线程内的索引计算必须严格符合 ISA 指定的映射,以确保每条寄存器向量包含正确的子矩阵片段。 性能估算与硬件参数 在进行性能评估与算法选择时,理解 MFMA 的周期计数与硬件并行度至关重要。给定一条 MFMA 指令的 M、N、K,以及其 cycle count,可用公式估算理论峰值 FLOP:2 * M * N * K * num_matrix_cores * (max_engine_clock / cycle_count) / 10^6。实际部署时需注意峰值通常难以达到,因为内存带宽、指令带宽、寄存器重命名与控制依赖以及其他硬件资源约束都会影响性能。
使用更低精度的数据类型能提高每周期的计算量,但也会因需要额外的缩放、格式转换或更多的指令序列而引入开销。 实践建议与常见陷阱 在迁移或新开发矩阵核代码时,建议首先从相对成熟且数值稳定的 FP16 开始验证功能与性能,再逐步尝试 FP8、FP6 与 FP4。量化误差与数值下溢、上溢需要通过校准和缩放策略来控制,CDNA4 的块级缩放特性为超低精度的实用化提供了强有力的工具。为最大化吞吐,应尽量设计线程索引与内存布局以符合 MFMA 的原生分块规则,避免在线程间做大量的数据重排列或序列化操作。 调试和验证方面,单步验证每个 MFMA 指令的输出与浮点高精度参考结果一致是保证正确性的基本手段。由于低精度格式的舍入与区间限制,逐元素比较可能无法通过,应采用相对误差度量并基于下游任务容忍度评估是否可接受。
将尺度编码与索引管理封装成可复用的加载/解包函数有助于降低工程复杂度。 工具链与生态兼容性 要在 CDNA 系列 GPU 上使用 Matrix Core 功能,需配套使用支持相应内建函数的 LLVM/Clang 与 ROCm 版本。CDNA4 的 block-scaled MFMA 支持通常出现在较新的 ROCm 发行版(例如 7.0 及后续版本)中,且编译器、头文件(如 hip_fp8.h、hip_ext_ocp.h)会提供相应的存储类型与位操作辅助函数。为保证跨设备可移植性,建议在构建系统中加入运行时硬件检测,并在不支持的设备上回退到标量或向量化实现路径。 结语 Matrix Core 编程在 AMD CDNA3 与 CDNA4 平台上为 AI 推理与高性能矩阵计算提供了明显的性能优势。理解低精度浮点格式的内部表示、掌握 MFMA 指令族与内建函数的参数语义、并按硬件要求组织数据布局,是发挥 Matrix Cores 算力的关键。
CDNA4 所引入的块级缩放和对 FP6/FP4 的支持进一步扩大了精度与性能权衡的空间,为需要极高吞吐的应用场景提供了新的选项。结合谨慎的数值校准、良好的内存访问模式以及与 ROCm/LLVM 工具链的紧密配合,开发者可以在 HIP 内核中高效利用 Matrix Cores,显著提升推理延迟与吞吐表现。对于追求极限性能的工程团队,建议在性能关键路径上采用逐层验证、自动化基准与硬件感知的编译选项,以便在不同 CDNA 代际间获得稳定而可重复的加速效果。 。
