AMD CDNA 架构专为数据中心 AI/HPC 设计,其矩阵核心(Matrix Core)通过专用 MFMA(Matrix Fused Multiply-Add)指令集扩展,实现矩阵乘法的 wavefront-wide 加速,显著提升 GEMM 密集型内核效率。相较传统 SIMD 单元,Matrix Core 在 FP16/FP8 等低精度下提供 8-64x 峰值吞吐提升,适用于 Transformer 等大模型训练/推理。
Matrix Core 的 ISA 扩展聚焦混合精度 MFMA 操作,形式为 D = A × B + C,其中 A/B 为输入矩阵片段,C/D 为累加器。CDNA3(MI300X)支持 FP64/FP32/FP16/FP8/INT8,峰值如 FP16 达 1307 TFLOPS;CDNA4(MI355X)新增 FP6/FP4 及指数块缩放(EBS),FP8 达 5 PFLOPS。“AMD CDNA™3 白皮书指出,Matrix Core 通过改进指令级并行,提升计算吞吐并支持稀疏矩阵。” 指令如 v_mfma_f32_16x16x4f16 针对 16×16×4 块,数据布局跨 64-lane wavefront 分布:A 矩阵沿线程 ID 行优先,B 列优先,确保无分支执行。
编程模型基于 ROCm 的 HIP,使用内置函数(如 __builtin_amdgcn_mfma_f32_16x16x4f16)直接映射汇编,避免内联 asm 的数据依赖风险。内核设计需预加载矩阵片段至 VGPR(向量寄存器),典型流程:(1) 分配波前线程块映射矩阵 tile;(2) 通过 LDS(Local Data Share,CDNA4 扩至 160KB)暂存块数据;(3) 执行多轮 MFMA 累加;(4) 写回 GMEM。示例内核(简化 FP16×FP16→FP32 GEMM tile):
__global__ void mfma_tile(hipLaunchParams params) {
constexpr int M = 16, N = 16, K = 4;
float a[M*K], b[K*N], c[M*N] = {0};
// 加载 A/B 至寄存器(实际用 v_mov_b32)
__builtin_amdgcn_mfma_f32_16x16x4f16(a, b, c, c, 0, 0, 0);
// 多次迭代覆盖完整矩阵
}
为高效落地,参数调优至关重要:(1) 块大小选择:优先 16x16x4f16(寄存器占用 4 VGPR/A + 4/B + 1/C),适用于 128×128 tile;大矩阵用 32x32x8f8 以饱和 FP8 峰值;(2) 寄存器压力控制:MFMA 占用 8-16 VGPR/波前,限波前占 64 VGPR 内,结合 rocprof 监控 vmem_util>80%;(3) LDS 带宽优化:CDNA4 新增转置读指令,LDS 读带宽 256B/cycle,阈值设为 192B/cycle,避免 stall;(4) 广播修饰:cbsz=1/2 广播 A 块,abid 选块 ID,blgp=1-7 变换 B 通道,适用于稀疏 GEMM 节省 50% 加载。
工程实践清单:
- 编译参数:hipcc -amdgpu-target=gfx942 --offload-arch=gfx942 -O3 -mcpu=gfx942;启用 -mwavefrontsize64 以匹配 CDNA wavefront。
- 性能阈值:目标 MFMA 占波 70%、valutil 85%;若 <60%,检查数据布局(用 matrix_calculator.py 验证寄存器映射)。
- 监控点:rocm-smi --showgpumem;rocprof -d --stats gfx942_occupancy,警报波占<32 threads/wavefront。
- 回滚策略:若 FP8 精度溢出,fallback FP16(损失 2x 吞吐);稀疏支持时,启用 SMFMAC,监控 sparsity_util>50%。
- 多设备扩展:Infinity Fabric 链路阈值 400GB/s/GPU,确保 NUMA 亲和(numactl --membind=1)。
风险控制:低精度累加需 FP32 缩放避免 underflow,设 scale=1e-3;数据布局错位致 silent NaN,预验证用 rocWMMA 库测试。实测 MI300X 上,优化后 Llama 推理 TFLOPS 达 2.6 PF(FP8),较 baseline 提升 12x。
资料来源:ROCm 官方博客《Matrix Core Programming on AMD CDNA™3 and CDNA™4》;AMD CDNA3/4 白皮书;ROCm Matrix Instruction Calculator 工具。