AMD CDNA 架构针对 HPC 和 AI 优化，其 Matrix Core 通过指令融合实现 GEMM 高吞吐。在 MI300X 上，FP16 峰值超 10 PFLOPS，稀疏加速翻倍。这种融合将矩阵乘累加（MFMA）浓缩为单指令，wavefront 级并行执行，远超 SIMD 向量单元。
观点一：MFMA 指令融合机制是 GEMM 加速核心。传统 GEMM 分解为循环外积，Matrix Core 融合 A×B+C 为 V_MFMA_FXxMxNxK，单周期处理 MxNxK 规模，支持 FP64 到 INT8。CDNA2/CDNA3 统一寄存器文件（512×64bit），消除 CDNA1 双 RF 开销，提升 ILP。
证据：ROCm 文档显示，__builtin_amdgcn_mfma_f32_16x16x4f32 内置函数单波前计算 16×16 输出，A(16×4)/B(4×16) 映射寄存器通道。“MFMA 在 wavefront 上操作，输入矩阵条目分布于矢量寄存器。”
落地参数：优先 16x16x4（FP32），32x32x4（BF16）。修饰符 cbsz=0/1（广播 A 块），blgp=0（默认布局）。编译：hipcc -O3 -mamdgpu-matrix-core=cdna3。寄存器阈值<192/256。
观点二：张量操作调度确保高占用率。调度器管理 wavefront 发射，MFMA 后 4 周期延迟需双缓冲隐藏。稀疏 SMFMAC 用 K 矩阵编码 2:4 模式，跳过零值，MI300 吞吐 2x。
证据：Matrix Instruction Calculator 查询 gfx942 --detail-instruction，MFMA_F32_32x32x4F32 吞吐 1024 OPS/CU，寄存器 64 个。
落地清单：
- 布局：A 行优先，线程(tx,ty)→A[ty][tx%K]，用 --register-layout 验证。
- 调度：K 循环展开 4/指令，tile 128x128x16（rocBLAS 默认）。
- 稀疏：阈值 0.5 剪枝，SMFMAC_F16_16x16x16BF16 峰值 2048 FLOPS/CU。
- 监控：rocm-smi mfma_util>85%，stall_wm>5% 调 blgp=1。
观点三：实战 GEMM 参数与稀疏加速。高吞吐需 CU 占用>90%，HBM3e 带宽>5TB/s。Transformer GEMM 受益最大，BERT 推理 3x 提升。
证据：CDNA3 支持 FP8/TF32+sparsity，rocm.blogs.amd.com 性能表：FP16 Matrix 1024 vs 向量 128，8x 加速。
落地参数表：
| 参数 |
值 |
应用 |
| 指令 |
V_MFMA_F16_16x16x4FP16 |
推理 |
| 稀疏 |
SMFMAC_F32_32x32x4F32 |
训练 |
| 缓冲 |
2 层共享内存 |
延迟隐藏 |
| Tile |
256x256x32 |
大矩阵 |
回滚：融合失效用 rocWMMA；溢出 FP32。风险：寄存器饥饿（finline-limit=2000），波 stall（-munsafe-fp-atomics）。
ROCm 集成：rocblas_gemm_ex 自动 MFMA，hip kernels 嵌入 intrinsics。MI300 集群部署，Llama 训练吞吐 4 PFLOPS/GPU。
资料来源：
（正文 920 字）