MI300X · CDNA3 · GFX942 · 特性专题

MFMA · Matrix Core

★ CDNA3 计算主线 · 累加器在 AccVGPR · FP64 也有矩阵路径（NVIDIA 的 2.4×）

是什么

DEFINITION

v_mfma_* 是 CDNA Matrix Core 的矩阵指令族，角色 ≈ wgmma/tcgen05，但语义是同步的：一个 wave（64 线程）发一条 MFMA，操作数 D/A/B 全部在 VGPR / AccVGPR 里，指令退休时结果就绪。每 CU 4 个 Matrix Core 挂在 VALU 发射口上——MFMA 可以与普通 VALU 指令同发射，矩阵运算的阴影里还能跑标量/向量活^[1]。

AccVGPR 是 MFMA 的专属累加器空间：每 wave 最多 256 个（与 256 常规 VGPR 合计 512 上限），且可以从 memory 直接载入（A 操作数不经 VGPR 中转）^[2]。

为什么（vs wgmma / tcgen05）

RATIONALE · VS NVIDIA MMA

维度	H200 wgmma	B200 tcgen05	MI300X MFMA
执行语义	异步（commit/wait）	异步（mbarrier）	同步（退休即就绪）
发射主体	warpgroup 128 线程	单线程	单 wave 64 线程
操作数位置	A/B 在 SMEM，D 在 RF	A/B 在 SMEM/TMEM，D 在 TMEM	全在 VGPR/AccVGPR
FP64 矩阵	67 TFLOPS	37 TFLOPS（弱化）	163.4 TFLOPS（×2.4 H200）^[1]
tile 尺寸	m64nNk16	最大 128×256×16	32×32×8 / 16×16×16（BF16）
延迟隐藏手段	异步流水	异步 + TMEM	wave 并行（32/CU）+ 与 VALU 同发射
2:4 稀疏	有	有	有（SMFMAC：I8/F8/F16/BF16/TF32）^[1]

AMD 的哲学：不造异步引擎，靠高 wave 并发 + 同发射把 MFMA 的延迟吃满。写惯 Hopper 双缓冲的人要换一套心智模型——这里是「足够多的 wave 各自同步推进」，不是「一个 warpgroup 异步流水」。

关键数字

KEY NUMBERS · SHAPES × RATES

项	值	注
Matrix Core 数	4/CU · 1,216 全片	挂 VALU 发射口^[1]
BF16 吞吐密度	2,048 FLOPs/clk/CU	FP8/INT8 4,096 · TF32 1,024 · FP64 256^[1]
MFMA 形状（BF16/FP16）	16×16×16 · 32×32×8（+4 块变体）	ISA Table 28^[2]
MFMA 形状（FP8/INT8）	16×16×32 · 32×32×16	FP8 含 BF8 变体
MFMA 形状（FP64）	16×16×4 · 4×4×4×4B	FP64 矩阵路径的依据
AccVGPR	≤256/wave（与 VGPR 合计 ≤512）	可从 memory 直载（A 操作数）^[2]
依赖链延迟（实测）	FP8 ~24.6ns · FP16 ~24.7ns · FP64 ~33.2ns	MI300A 实测，GPU event 计时^[5]
稀疏指令	v_smfmac_*（2:4）	I8 16×16×64 · F16 16×16×32 等

可编译示例

COMPILABLE EXAMPLE · VERBATIM

逐字对齐 CDNA3 ISA 语义；生产代码用 hipBLASLt / Composable Kernel，裸 asm 只在写算子库时出现。

① 单 wave 发一条 BF16 MFMA（D=A×B+C 全在寄存器）

// 16×16×16 BF16 输入 → FP32 累加：
// A/B 每 lane 4×BF16（v[0:3], v[4:7]），C/D 每 lane 4×FP32（v[8:11] / v[8:11]）
asm("v_mfma_f32_16x16x16f16 a[0:3], v[0:3], v[4:7], a[0:3]");
//                              ^AccVGPR 累加器：直载直用，不占常规 VGPR

② FP8（含 BF8）与稀疏变体

// FP8 e4m3 × e4m3 → FP32，K 维翻倍（32）：
asm("v_mfma_f32_16x16x32_bf8_bf8 a[0:3], v[0:1], v[2:3], a[0:3]");
// 2:4 结构化稀疏（SMFMAC）：B 半零，K 再翻倍：
asm("v_smfmac_f32_16x16x32_bf16 a[0:3], v[0:3], v[4:11], a[0:3]");

③ 实用路径：hipBLASLt / Composable Kernel

// hipBLASLt GEMM（等价 cuBLASLt API 形状）：
hipblasLtMatmul(handle, &desc, &alpha,
                  A, B, &beta, C, D, &algo, workspace);
// 自研算子：Composable Kernel 的 tile 编程
// （CK 对 MFMA/AccVGPR/LDS 布局做了完整抽象）

注意事项 / 陷阱

PITFALLS · CHECKLIST

CAUTION · 陷阱清单

MFMA 是同步指令：发射后流水线停等结果就绪——吞吐靠同时驻留的 32 个 wave 轮转，单 wave 串行发 MFMA 是性能自杀。
操作数必须先进 VGPR/AccVGPR：没有「A/B 直接读 LDS/SMEM」的路径（对比 wgmma/tcgen05）——LDS→VGPR 的搬运本身要算进 roofline。
AccVGPR 与常规 VGPR 合计 512 上限：大累加器（32×32 tile）会把 wave 的 VGPR 预算吃光，occupancy 掉一半是常事。
形状选错直接损失峰值：32×32×8 与 16×16×16 吞吐相同，但 寄存器占用和 wave 映射不同——按 CK 文档的形状选择表来。
2:4 稀疏用 SMFMAC 且只对 B 半边压缩，元数据布局有专用格式；权重不满足 2:4 时没有收益还多一道转换。
FP64 MFMA（163T）是 MI300X 对 H200（67T）/B200（37T）的最大优势项——HPC 迁移评估时重点看它，但注意 B200 的 FP64 vector 也只剩 37T。

数据源

SOURCES

CDNA3 ISA Reference（MI300）（v_mfma/v_smfmac 全形状表 + AccVGPR 语义）
CDNA3 White Paper（Matrix Core 组织 / 吞吐密度）
Composable Kernel — Block GEMM on MI300（形状选择 / tile 实践）
arXiv 2602.10262（MFMA 依赖链延迟实测）