随着人工智能推理与高性能计算对矩阵运算性能和能效要求的不断提升,GPU 中的专用矩阵运算单元(Matrix Cores)成为提升吞吐的关键。AMD 在 CDNA 系列架构中对矩阵核(Matrix Core)进行了持续演进,从 CDNA 3 到 CDNA 4 扩展了低精度格式、引入块级指数缩放并显著提升了 FP16/FP8 的吞吐能力。对希望在 AMD 平台上实现高性能混合精度计算的开发者来说,理解这些硬件特性、数据布局与编译器内建函数(intrinsics)是关键。本文将以开发实践角度系统介绍在 HIP/ROCm 下如何高效使用 Matrix Core,包括低精度类型的设计权衡、MFMA 指令族的选择、寄存器与线程间的数据分配、以及 CDNA 4 引入的块指数缩放(block exponent scaling)方法。 矩阵核性能与混合精度价值 矩阵核通过硬件级的矩阵乘加(MFMA, matrix fused-multiply-add)指令实现高密度的乘加流水线,专门为深度学习和科研计算中常见的密集矩阵乘法场景优化。在混合精度模式下,输入矩阵采用较低位宽(如 FP16/FP8/FP4),累加器使用更高精度(通常为 FP32),以平衡吞吐与数值稳定性。
CDNA 3 的 Matrix Core 在 FP16 与 FP8 下提供了数倍于 FP32 的峰值性能,而 CDNA 4 在指令集合与硬件吞吐上进一步提升,新增 FP6/FP4 格式并支持块指数缩放,从而在保持精度可控的前提下,把性能潜力推向更高阶。 低精度浮点格式与数值范围 低精度浮点格式不是只看位数,更要关注指数位和尾数位的分配,因为它决定了动态范围与精度。常见格式包括 E5M10 的 FP16、E4M3 与 E5M2 的 FP8、E3M2 或 E2M3 的 FP6,以及更窄的 E2M1(FP4)。不同变体还会对零与 NaN、Infinity 的表达有所差异,例如 E4M3 存在 FNUZ(unsigned zero)和 FN(finite)等变体,影响是否支持负零或无限大。除了元素格式,CDNA 引入的 E8M0 格式可用于作为块缩放(microscaling)因子的编码,等价于按指数尺度对若干元素块统一放缩。 在实际工程中,对格式的选择需要综合考虑模型对动态范围与舍入误差的敏感度、量化策略、以及后端硬件对格式的原生支持。
FP16(E5M10)具有较广的应用基础,兼顾范围与精度;FP8(E4M3/E5M2)在推理场景中可显著提升吞吐,但通常需要更严格的缩放或校正策略;FP4 则适合极度吞吐优先且可以容忍精度退化的场景,并通常结合块缩放来扩大表示范围。 MFMA 指令族与数据分布 在 CDNA 架构中,MFMA 指令以 MxNxK 的形状定义单条指令所处理的矩阵片大小,同时指定输入与输出的数据类型,以及指令的周期数(cycles)。这些参数直接决定单指令的计算量与延迟,从而影响理论峰值算力。MFMA 为 wavefront(在 CDNA 上通常为 64 个线程)级别指令,意味着一个 wavefront 的所有线程协同执行一条 MFMA 指令,寄存器中保存的是输入矩阵 A、B 与累加矩阵 C 的局部片段。理解寄存器数据布局对正确加载/存储与高效实现至关重要。例如 32x32x16 的 FP8->FP32 指令会把 A、B 的多个元素分配到每个线程的寄存器向量中,C 的累加部分也以向量形式保存在每个线程,以便指令调用时能正确组合数据。
编译器内建函数(intrinsics)与 HIP 实践 开发者在 HIP 内核中可以通过 LLVM 提供的一组内建函数直接发出 MFMA 指令,典型的内建函数形式为 __builtin_amdgcn_mfma_f32_32x32x16_fp8_fp8(a,b,c,0,0,0) 或 CDNA 4 的块缩放变体 __builtin_amdgcn_mfma_scale_f32_32x32x64_f8f6f4(a,b,c,Atype,Btype,OPSEL_A,scale_a,OPSEL_B,scale_b)。参数中指定了矩阵形状、输入/输出数据类型以及可选的缩放或广播控制。需要注意的是不同指令对寄存器中向量宽度与对齐有严格要求,常见做法是使用属性 vector_size 或 ext_vector_type 定义固定宽度的寄存器向量类型,并确保在载入内存时以正确的访存模式拼装这些向量。 在实际代码中,FP8/FP4 等小位宽类型通常在内存中以打包字节形式存储,例如 FP8 以 8 位为单元,FP4 以两个元素打包在一个字节中。为配合内建函数,开发者需要在内核里进行打包/解包和位域提取,例如用位运算从 uint8_t 中抽取两个 4bit 的 FP4 元素,再合成编译器期望的 256bit 或 128bit 向量寄存器。HIP 提供的扩展头文件(如 hip_fp8.h 或 hip_ext_ocp.h)中包含了用于创建或提取这些小位宽元素的辅助宏与内联函数,利用这些工具可以减少手写位操作的复杂性。
数据布局和访存模式的影响 高效利用 Matrix Core 的一个关键因素是把数据按照 MFMA 指定的线程间分配规则布局到全局内存中,避免在内核启动后进行复杂的重排操作。大多数 MFMA 指令假设矩阵以行优先(row-major)存放,并根据 wavefront 内线程索引把 A、B、C 的子块分派给每个线程。错误的内存布局会导致每个线程必须随机访问不连续地址,极大地降低访存带宽利用率并增加延迟。建议在主机端预先按目标 MFMA 指令的线程分配方式组织矩阵;如果无法做到,内核中也可以通过协同加载(coalesced loads)或 LDS(本地数据共享)来重排,但这通常增加寄存器与指令开销。 块指数缩放(Block Exponent Scaling)的机理与优势 CDNA 4 引入了块指数缩放(block exponent scaling)机制,使得可以在 MFMA 级别对 A 和 B 的小块应用不同的缩放因子。缩放因子通常以 E8M0 格式编码为单字节,其含义是以 2^(exp-127) 的形式对经过低精度乘加的中间结果进行线性缩放。
块缩放的优势在于它能显著扩展低精度格式的有效动态范围,尤其在 FP8/FP4 这样指数位较窄的格式中更为重要。通过在不改变元素本身编码的情况下,对若干行或列块施加不同的缩放,模型可以在保留高吞吐的前提下减小溢出与下溢的风险,从而提升数值稳定性与实际推理精度。 实现块缩放时需要考虑缩放因子的生成与存储开销。典型做法是在前处理或量化阶段为每个块计算合适的 scale_a 与 scale_b,并将这些单字节的缩放因子与矩阵一起存储或在运行时传入内核。MFMA 内建函数会在完成块内点积后自动将缩放应用到结果,再与累加器相加,从而把缩放逻辑下沉到硬件指令级别,减少了软件层面的显式乘法带来的开销。 工程实践建议与调优方向 在将 Matrix Core 引入现有计算流水线时,首先需要评估模型对低精度的容忍度。
对一些敏感的层可以保留 FP32/FP16,而对大多数矩阵密集层可以尝试 FP8 或更低位宽。量化与缩放策略是成功的关键。常见方法包括逐通道动态缩放、最大绝对值缩放、或基于统计量(如均值与方差)估算的缩放因子。对于 CDNA 4 上可以采用的块缩放,建议把缩放粒度与硬件指令片(例如 32x32x64 指令的 Ax/Bx 维度)对齐,从而最大化硬件加速效果。 在内核实现层面,务必对齐寄存器向量宽度与内建函数的参数期待。例如调用要求 256bit 向量的 MFMA 内建函数时,开发者需要构造等宽的寄存器类型并保证高位或低位按预期填充零或有效数据。
访存访问应尽量做到按 128-bit 或 256-bit 对齐的连续加载,若元素为 FP4,需先在内存中做好打包以便连续载入后解包到寄存器向量。最后,利用 ROCm 的性能分析工具(如 rocprof 或 rocTracer)来分析内核中的算力、带宽与寄存器使用率,针对热点进行寄存器重用、指令流并行化或减少内存副本等优化。 示例场景与性能估算 理解 MFMA 指令的周期数与每指令的浮点运算计数,可以估算理论峰值。例如某条 MFMA 指令定义为 32x32x8,若每个指令需要 32 个 cycles,GPU 上 Matrix Core 的总数与引擎频率已知,则可以通过 2*M*N*K * num_matrix_cores * (max_engine_clock / cycle_count) 来估算理论 TFLOP 级别的吞吐。实际工程中达到理论峰值还需要综合考虑内存带宽、指令并发与数据布局的效率,因此常见的调优工作包括减少内核间同步、提升寄存器/片上存储利用率、以及确保全局内存访问并行化。 未来趋势与生态成熟度 CDNA 4 对低精度指令与块缩放的支持表明了硬件在更极端混合精度场景下的设计方向。
未来若更多软件框架(如 PyTorch、TensorFlow)原生支持这些指令与缩放策略,将极大降低开发者的工程门槛。目前在 ROCm/ROCm-ML 生态中,许多工具库和算子正在加入对 FP8、FP6 与 FP4 的支持,以及对块缩放的量化流水线。开发者应关注上游编译器对 intrinsics 的支持、硬件 ABI 的演进、以及社区关于量化策略与数值稳定性的最佳实践。 总结要点 在 AMD CDNA 3 与 CDNA 4 架构上,Matrix Core 为混合精度矩阵运算提供了强大的硬件加速能力。掌握低精度浮点格式的特性、MFMA 指令族的形状与周期、以及寄存器与线程级别的数据布局,是在 HIP/ROCm 中实现高效矩阵核的基础。CDNA 4 引入的块指数缩放为极低位宽(FP8/FP4/FP6)场景提供了实用的数值扩展手段,通过硬件级缩放可以在更低位宽下保持更稳定的数值表现。
工程实践中需要在量化方案、缩放策略与内核实现上做系统性的权衡与调优,利用硬件原生的内建函数与 ROCm 工具链进行性能剖析与迭代。随着生态逐步完善,基于 Matrix Core 的混合精度计算将在推理与训练任务中日益成为提升性能与能效的核心路径。 。
