NVIDIA MMA

NVIDIA Tensor Core 通过 MMA（Matrix Multiply-Accumulate，矩阵乘加）执行矩阵 tile 计算。本文先介绍理解 MMA 所需的基础，再以 Ampere、Ada、Hopper 和 Blackwell 为主线，说明不同 GPU 架构能让哪些精度直接进入 Tensor Core，并给出对应的 PTX 指令示例。

从 FMA 到 MMA，再到 GEMM

普通 FMA（Fused Multiply-Add）计算一个标量乘加：

d = a \times b + c .

MMA 是它的矩阵版本：

D = A \times B + C,

其中：

A \in R^{M \times K}, B \in R^{K \times N}, C, D \in R^{M \times N} .

一条 MMA 只处理固定大小的矩阵块。完整 GEMM（General Matrix Multiplication）会把大矩阵切成许多 tile，并沿 $K$ 维反复执行 MMA：

D_{m, n} = \sum_{k} A_{m, k} B_{k, n} + C_{m, n} .

text

FMA  ：一个标量的乘加
MMA  ：一个固定矩阵 tile 的乘加
GEMM ：许多 MMA，加上数据搬运、K 循环和 epilogue

稠密乘法与 2:4 结构化稀疏乘法

稠密（dense）MMA 保存并计算 A、B 的全部元素，即使某个元素恰好为 0 也不会跳过。本文的稀疏（sparse）特指 Tensor Core 支持的结构化稀疏 A：以 2:4 格式为例，A 沿 K 维每连续 4 个位置只保存 2 个非零值和相应的位置 metadata：

text

逻辑 A       ：[a0, 0, 0, a3]
压缩值       ：[a0, a3]
sparse metadata：[0, 3]      // 两个非零值原来的位置

稀疏 MMA 仍计算 $D = A B + C$ ，且只有 A 被压缩，B、C、D 都保持稠密：

对象	稠密 MMA	2:4 稀疏 MMA
A	保存全部 K 个元素	每 4 个逻辑位置保存 2 个值，并附带位置 metadata
B	稠密	稠密
C/D	稠密累加器	稠密累加器
PTX 名称	`mma.sync`、`wgmma.mma_async`、`tcgen05.mma`	`mma.sp`、`wgmma.mma_async.sp`、`tcgen05.mma.sp`

它的优势来自跳过确定为 0 的乘法，而不是改变数值精度：

计算：一条 sparse MMA 通常以两倍 K 覆盖相同类型的逻辑计算，理论有效吞吐最高可达对应 dense MMA 的 2 倍；
存储：A 只保存约一半的 value payload 和少量 metadata，可降低权重容量与搬运量；
调度：固定的 2:4 分组比任意稀疏格式更容易均匀映射到 Tensor Core。

这只是指令级上限。A（通常是 weight）必须先经剪枝、微调和压缩，并满足严格的分组与 metadata 约束；B 读取、C/D 写回、epilogue 等瓶颈也会削弱端到端收益。稀疏指令的 shape 描述逻辑矩阵，因此 K 通常是对应 dense 指令的两倍。

输入/累加/输出

矩阵乘区分 A/B 输入、C/D 累加器和输出：

对象	作用	示例
A/B 输入类型	Tensor Core 接收的两个乘数	FP8、INT8、FP16
C/D 累加器类型	保存沿 $K$ 维的乘积和	FP32、INT32
最终输出类型	epilogue 写回显存的 tensor 类型	FP16、BF16、FP8

例如：

text

FP8 A × FP8 B
    → FP32 accumulator
    → epilogue：scale / bias / activation / cast
    → FP16 output

FP8 输入、FP32 累加和 FP16 输出描述的是三个不同阶段，并不矛盾。低精度输入通常需要更宽的累加器，因为每个输出元素要沿 $K$ 维累加大量乘积：

D_{i, j} = \sum_{k = 0}^{K - 1} A_{i, k} B_{k, j} .

以 INT4 为例，单次 $7 \times 7 = 49$ 已无法放回有符号 INT4，数千项相加更需要 INT32。FP16 也会因有限的动态范围和尾数精度，更容易在长归约中溢出或舍掉小增量，因此训练和高精度推理常选 FP32 累加。

PTX 类型

以下 FP16 指令可用于 SM80 及之后的架构：

text

mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32

末尾四个类型依次为 D、A、B、C，因此该指令计算 FP16 A × FP16 B + FP32 C → FP32 D。累加类型必须来自硬件支持的签名，不能任意创造 INT16 × INT16 Tensor Core MMA。

后文表格中的 FP16 / FP32 表示两种可选变体。mma.sync 显式传入 C；wgmma.mma_async 和 tcgen05.mma 则让 D 同时充当输入累加器和输出，所以其 C 列表示“执行前的 D”。指令{.sp} 同时代表 dense 和 sparse 版本，shape 按 dense / sparse 排列。

m16n8k16 表示一个 warp 协作完成：

A_{16 \times 16} \times B_{16 \times 8} + C_{16 \times 8} \to D_{16 \times 8} .

即 $16 \times 8 \times 16 = 2048$ 个标量乘加位置。m/n/k 描述一条指令的逻辑 tile，而不是完整 GEMM 的 shape。

精度总表

下表讨论的是稠密 Tensor Core MMA 的 A/B 输入类型：

A/B 输入类型	Ampere SM80	Ada SM89	Hopper SM90	Blackwell SM100	Blackwell SM120	累加器类型
FP64	支持	支持	支持	支持传统 MMA	支持传统 MMA	FP64
FP16	支持	支持	支持	支持	支持	FP16 / FP32
BF16	支持	支持	支持	支持	支持	FP32
TF32	支持	支持	支持	支持	支持	FP32
FP8 E4M3 / E5M2	不支持	支持	支持	支持	支持	非块缩放：FP16 / FP32；block-scaled：FP32
FP6 E3M2 / E2M3	不支持	不支持	不支持	支持	支持	非块缩放：FP16 / FP32；block-scaled：FP32
FP4 E2M1	不支持	不支持	不支持	支持	支持	非块缩放：FP16 / FP32；block-scaled：FP32
INT8 `s8/u8`	支持	支持	支持	支持	支持	INT32
INT4 `s4/u4`	支持	支持	支持传统 MMA	支持传统 MMA	支持传统 MMA	INT32
B1	支持	支持	支持（含 WGMMA）	支持传统 MMA	支持传统 MMA	INT32（逻辑运算后 `popc`）

这里的“支持”只表示架构存在相应签名，不保证任意布局、A/B 类型组合、tile shape 或量化格式都可直接使用。“传统 MMA”特指继承的 warp 级 mma.sync，不表示该架构的新指令族也支持此类型。INT4 与 FP4 是不同格式；INT16 可以参与普通整数指令，但不能直接进入上述 Tensor Core MMA。各指令族的限制见后文架构表。

Ampere：FP16/BF16 与整数 MMA 基线

Ampere SM80（典型卡型为 A100）使用 warp 级 mma.sync，并加入 2:4 结构化稀疏的 mma.sp。下表把 SM80 作为基线，列出该架构可用的底层 PTX MMA 组合；后续架构只列相对前表新增的指令或变体，不重复继承项。表中不展开更高层的 wmma.* fragment 接口。

指令变体	shape（dense / sparse）	A	B	C	D	一句话总结
`mma{.sp}` FP16	`m8n8k4, m16n8k{8,16}` / `m16n8k{16,32}`	FP16	FP16	FP16 / FP32	FP16 / FP32	`.sp` 只把 A 改成 2:4 sparse；`m8n8k4` 还允许 FP16 C → FP32 D
`mma{.sp}` BF16	`m16n8k{8,16}` / `m16n8k{16,32}`	BF16	BF16	FP32	FP32	BF16 乘法、FP32 累加
`mma{.sp}` TF32	`m16n8k{4,8}` / `m16n8k{8,16}`	TF32	TF32	FP32	FP32	输入占 32 bit，但乘法按 TF32 精度执行
`mma` FP64	`m8n8k4` / —	FP64	FP64	FP64	FP64	SM80 没有 sparse FP64 MMA
`mma{.sp}` INT8	`m8n8k16, m16n8k{16,32}` / `m16n8k{32,64}`	`s8/u8`	`s8/u8`	INT32	INT32	A/B 可混合有符号和无符号类型
`mma{.sp}` INT4	`m8n8k32, m16n8k{32,64}` / `m16n8k{64,128}`	`s4/u4`	`s4/u4`	INT32	INT32	整数 INT4，不是 FP4 E2M1
`mma` B1	`m8n8k128, m16n8k{128,256}` / —	B1	B1	INT32	INT32	只有 dense；对 A/B 做 `xor/and` 后 `popc`

这些参数也适用于后续 mma.sync：.sync.aligned 要求整个 warp 以一致的谓词和限定符执行；.row.col 表示 A 为 row-major、B 为 column-major；类型后缀依次描述 D、A、B、C。整数指令可用 .satfinite 将溢出值饱和到 INT32 边界，否则按模回绕。稀疏形式还会增加 metadata e 和选择器 f；推荐的 mma.sp::ordered_metadata 要求非零位置索引升序排列。

两个典型签名分别展示浮点与整数路径：

text

// FP16 × FP16 + FP32 → FP32
mma.sync.aligned.m16n8k16.row.col.f32.f16.f16.f32 d, a, b, c;

// UINT4 × INT4 + INT32 → INT32
mma.sync.aligned.m16n8k32.row.col.s32.u4.s4.s32 d, a, b, c;

INT4 MMA 要求 A、B 都是 4 bit 整数，不能直接计算 FP16 activation 与 INT4 weight；常见 W4A16 Kernel 会把当前 weight tile 转换为 FP16/BF16，再执行浮点 MMA。

Ada：`mma.sync` 加入 FP8

Ada SM89（例如 GeForce RTX 4090、L40S）继承 SM80 表中的组合，本代新增的是 FP8 E4M3/E5M2 的稠密与 2:4 稀疏 mma：

指令变体	shape（dense / sparse）	A	B	C	D	一句话总结
`mma{.sp}` FP8	`m16n8k{16,32}` / `m16n8k64`	E4M3 / E5M2	E4M3 / E5M2	dense：FP16 / FP32；sparse：FP32	dense：FP16 / FP32；sparse：FP32	A/B 格式可不同；旧式 sparse FP8 仅提供 FP32 累加

这里的 e4m3/e5m2 是不含 block scale 的裸 FP8。下面的签名表示 E4M3 A × E5M2 B + FP32 C → FP32 D；将 D/C 后缀改为 f16，即可选择 FP16 累加版本。

text

mma.sync.aligned.m16n8k32.row.col.f32.e4m3.e5m2.f32 d, a, b, c;

Hopper：增加 WGMMA

Hopper SM90（例如 H100/H200）保留前两表的 warp 级指令，并新增 FP64 warp MMA 形状和由 4 个 warp、共 128 个线程协作的异步 WGMMA。PTX 中 WGMMA 是架构专用能力，目标写作 sm_90a。

下表中浮点 WGMMA 的 $N = 8, 16, \dots, 256$ ，即从 8 到 256、步长为 8；整数与 B1 的 $N$ 为 8/16/24/32，之后从 48 到 256、步长为 16。

指令变体	shape（dense / sparse）	A	B	C（原 D）	D	一句话总结
`mma.sync` FP64 新形状	`m16n8k{4,8,16}` / —	FP64	FP64	FP64	FP64	SM90 新增更大的 warp 级 FP64 tile，不属于 WGMMA
`wgmma.mma_async{.sp}` FP16	`m64nNk16` / `m64nNk32`	FP16	FP16	FP16 / FP32	FP16 / FP32	sparse 版只把 A 改为 2:4，并把 K 加倍
`wgmma.mma_async{.sp}` BF16	`m64nNk16` / `m64nNk32`	BF16	BF16	FP32	FP32	BF16 只提供 FP32 累加
`wgmma.mma_async{.sp}` TF32	`m64nNk8` / `m64nNk16`	TF32	TF32	FP32	FP32	sparse 版 K 加倍
`wgmma.mma_async{.sp}` FP8	`m64nNk32` / `m64nNk64`	E4M3 / E5M2	E4M3 / E5M2	FP16 / FP32	FP16 / FP32	A/B 可混用两种 FP8；FP32 内部累加低于完整 FP32 精度
`wgmma.mma_async{.sp}` INT8	`m64nNk32` / `m64nNk64`	`s8/u8`	`s8/u8`	INT32	INT32	支持有/无符号混合；WGMMA 最低只到 INT8
`wgmma.mma_async` B1	`m64nNk256` / —	B1	B1	INT32	INT32	只有 dense，且只支持 `and.popc`

FP64 warp MMA 可用 .rn/.rz/.rm/.rp 指定舍入；整数 WGMMA 可用 .satfinite 控制溢出；稀疏 WGMMA 额外接收 sp-meta 和 sp-sel。WGMMA 的几个 scale 参数容易混淆：scale-d 决定计算 $A B + D$ 还是 $A B$ ，imm-scale-a/b 只表示是否取负，并非量化 scale；FP16/BF16 descriptor 版本还可用 imm-trans-a/b 控制转置。

a-desc/b-desc 是 shared-memory matrix descriptor，编码地址、stride 和 swizzle 等布局。异步序列需要依次执行 wgmma.fence、MMA、wgmma.commit_group，并在读取 D 前执行 wgmma.wait_group。

例如，FP8 mixed-input、FP32 累加可以写成：

text

wgmma.mma_async.sync.aligned.m64n8k32.f32.e5m2.e4m3
    d, a-desc, b-desc, scale-d, imm-scale-a, imm-scale-b;

WGMMA 不接受 INT4，但 Hopper 仍能用继承的 mma.sync 执行 INT4。以 FP4 保存的 weight 也不能直接交给 WGMMA，Kernel 需要先在 tile 内将其转换为 FP8、FP16、BF16 或 INT8。

Blackwell：FP4/FP6 与 block-scaled MMA

Blackwell 首次让 FP4 E2M1 和 FP6 E3M2/E2M3 直接进入 Tensor Core。数据中心 SM100 使用 tcgen05.mma，消费级 SM120 则扩展 warp 级 mma.sync；两者的编程接口不同，但都支持 FP4/FP6/FP8 mixed-input 和 block scaling。

SM100：`tcgen05.mma`

SM100（例如 B200/GB200）用第五代 Tensor Core 指令取代 WGMMA 作为新主路径，对应 PTX 架构专用目标 sm_100a。tcgen05 是异步、单线程发射的指令：发射线程启动整个 MMA，D 写入 Tensor Memory（TMEM），并不要求 128 个线程同时执行同一条指令。以下只列 SM100 新增的 tcgen05，不重复旧 mma.sync：

这里用 F8/F6/F4 简写 {E4M3, E5M2, E3M2, E2M3, E2M1}。指令名中的 {.sp} 表示同时存在稠密和 sparse A 版本；mma.ws{.sp} 具有相同的数据类型，但只用于 weight-stationary 卷积。

指令与 kind	K（dense / sparse）	A	B	C（原 D）	D	一句话总结
`mma{.sp}` / `mma.ws{.sp}` `kind::f16`	`16 / 32`	FP16	FP16	FP16 / FP32	FP16 / FP32	FP16 输入可选择 FP16 或 FP32 累加
`mma{.sp}` / `mma.ws{.sp}` `kind::f16`	`16 / 32`	BF16	BF16	FP32	FP32	BF16 复用 `kind::f16`，但只允许 FP32 累加
`mma{.sp}` / `mma.ws{.sp}` `kind::tf32`	`8 / 16`	TF32	TF32	FP32	FP32	TF32 乘法、FP32 累加
`mma{.sp}` / `mma.ws{.sp}` `kind::f8f6f4`	`32 / 64`	任一 F8/F6/F4	任一 F8/F6/F4	FP16 / FP32	FP16 / FP32	A/B 可不同，例如 FP4 × FP8；这是裸窄浮点路径
`mma{.sp}` / `mma.ws{.sp}` `kind::i8`	`32 / 64`	`s8/u8`	`s8/u8`	INT32	INT32	新整数路径只有 INT8，没有 `tcgen05` INT4
`mma{.sp}` `kind::mxf8f6f4.block_scale`	`32 / 64`	MXFP8/6/4 + UE8M0 scale	MXFP8/6/4 + UE8M0 scale	FP32	FP32	A/B 可为不同 MX 窄浮点，按 block 32 缩放
`mma{.sp}` `kind::mxf4.block_scale`	`64 / 128`	MXFP4 E2M1 + UE8M0	MXFP4 E2M1 + UE8M0	FP32	FP32	专用 MXFP4 路径
`mma{.sp}` `kind::mxf4nvf4.block_scale`	`64 / 128`	MXFP4 + UE8M0，或 NVFP4 + UE4M3	与 A 相同的格式族	FP32	FP32	支持 MXFP4 或 NVFP4，但两侧不能混用不同 scale 格式

四个指令名可由两个正交特征理解：.sp 表示 A 为 2:4 sparse，并增加 [sp-meta-tmem]；.ws 表示用于卷积的 weight-stationary 形式，否则为 activation-stationary。普通 mma{.sp} 支持 cta_group::1/2，.ws 只支持 cta_group::1 且没有 block-scaled kind。最终合法的 M/N/K、CTA group 和 sparsity 组合由 idesc 限定。

下面分别给出裸窄浮点和 block-scaled 形式。idesc 编码 M/N/K、A/B 类型与布局，因此 FP4 × FP8 不需要单独的指令名：

text

// 裸 FP4/FP6/FP8 mixed-input
tcgen05.mma.cta_group::1.kind::f8f6f4
    [d-tmem], a-desc, b-desc, idesc,
    {disable-output-lane}, enable-input-d;

// Block-scaled FP4/FP6/FP8
tcgen05.mma.cta_group::1.kind::mxf8f6f4.block_scale.scale_vec::1X
    [d-tmem], a-desc, b-desc, idesc,
    [scale-a-tmem], [scale-b-tmem], enable-input-d;

[d-tmem] 是 D/累加器的 TMEM 地址；A 可来自 shared memory 或 TMEM，B 来自 shared memory。enable-input-d 决定是否累加原 D，普通 FP16/TF32 还可用 scale-input-d 缩放原 D；disable-output-lane 用于屏蔽部分输出列，block-scaled 形式则在该位置传入 A/B scale。.scale_vec::1X/2X/4X 控制 scale vector 的读取与复用；collector 限定符提示硬件复用 A 或 B，tcgen05.commit...mbarrier 用于跟踪异步完成。

其核心计算为：

D = C + \sum_{g} (A_{q, g} S_{A, g}) (B_{q, g} S_{B, g}) .

Scale 在每个 $K$ block 的部分和进入最终累加之前应用，而不是等完整 GEMM 结束后只乘一次。

SM120：扩展的 `mma.sync`

SM120（例如 GeForce RTX 50 系列）不是把 SM100 的 tcgen05 原样下放，而是扩展 32-thread warp 级 mma.sync。它继承传统 FP16/BF16/TF32/FP8/整数 warp MMA，以下只列 SM120 新增的窄浮点和 block-scaled 变体；F8/F6/F4 沿用 SM100 表中的定义。

指令与 kind	shape（dense / sparse）	A	B	C	D	一句话总结
`mma{.sp}` `kind::f8f6f4`	`m16n8k32` / `m16n8k64`	任一 F8/F6/F4	任一 F8/F6/F4	FP16 / FP32	FP16 / FP32	支持 FP4 × FP8；`.sp` 只把 A 改成 2:4 sparse
`mma{.sp}` `kind::mxf8f6f4.block_scale`	`m16n8k32` / `m16n8k64`	MXFP8/6/4 + UE8M0	MXFP8/6/4 + UE8M0	FP32	FP32	block 32 的 dense / sparse MX 路径
`mma{.sp}` `kind::mxf4.block_scale`	`m16n8k64` / `m16n8k128`	MXFP4 E2M1 + UE8M0	MXFP4 E2M1 + UE8M0	FP32	FP32	专用 MXFP4 dense / sparse 路径
`mma{.sp}` `kind::mxf4nvf4.block_scale`	`m16n8k64` / `m16n8k128`	MXFP4 + UE8M0，或 NVFP4 + UE4M3	与 A 相同的格式族	FP32	FP32	支持 MXFP4/NVFP4，sparse 版只压缩 A

这些新签名固定为 A row-major、B column-major，稀疏版使用 mma.sp::ordered_metadata。裸 kind::f8f6f4 将 FP4/FP6 放入带 padding 的 8 bit 容器，mxf4/mxf4nvf4 则紧凑打包 E2M1。编译目标应使用 sm_120a；SM100 的 tcgen05 不能当作 SM120 指令使用。

一个 FP4 × FP8、FP32 累加的例子是：

text

mma.sync.aligned.m16n8k32.row.col.kind::f8f6f4.f32.e2m1.e4m3.f32
    d, a, b, c;

Block-scaled 形式会在 d, a, b, c 后增加 A/B scale 及其 byte/thread ID；stype 选择 UE8M0 或 UE4M3，scale_vec 控制 scale vector 的读取与复用。实际 Kernel 仍需满足 sub-byte 对齐、A/B layout 和 scale swizzle 契约。

MXFP4 与 NVFP4

裸 FP4 E2M1 的动态范围很小，模型部署更常使用带 block scale 的格式：

格式	数据元素	Scale 类型	稠密情况下每个 scale 覆盖的 $K$ 元素数
MXFP4	FP4 E2M1	UE8M0	32
NVFP4	FP4 E2M1	UE4M3	16

Blackwell 可在 MMA 路径中直接读取 FP4 与对应 scale，无需先把完整 weight tensor 反量化为 FP16。FP4 不等于 INT4：传统 INT4 继续使用继承的 mma.sync，SM100 的 tcgen05.mma 整数 kind 只有 kind::i8。

量化配置如何映射到 MMA

W4A16、W8A8 和 W4A8 描述 Weight 与 Activation 的存储或量化位宽，不能单独确定 Tensor Core 实际接收的类型。

量化配置	实际 MMA 输入 → 累加器	映射条件
W4A16 INT4 Weight-Only	FP16 × FP16 → FP16/FP32，或 BF16 × BF16 → FP32	weight tile 先反量化；通常不是 INT4 MMA
W8A8 INT8	INT8 × INT8 → INT32	本文所列 SM80 及之后架构均可直接映射
W8A8 FP8	裸 FP8 × 裸 FP8 → FP16/FP32；MXFP8 × MXFP8 → FP32	裸 FP8 从 SM89 开始；硬件 block-scaled MXFP8 需要 Blackwell
W4A8，Hopper	INT4 weight tile 转为 INT8 后执行 INT8 × INT8 → INT32，或转为 FP8 后执行 FP8 × FP8 → FP16/FP32	WGMMA 不接受 INT4；转换通常融合在 mainloop 中
W4A8，Blackwell FP4	裸 FP4 × 裸 FP8 → FP16/FP32；MXFP4 × MXFP8 → FP32	分别匹配 `kind::f8f6f4` 和 `kind::mxf8f6f4` 的格式、scale 与布局
W4A4，Blackwell NVFP4	NVFP4 × NVFP4 → FP32	两侧都是 E2M1，使用 UE4M3 scale 和 block 16

这里的“原生”指 A/B 的低精度类型和 scale 能被 MMA 签名直接接收。即使 INT4 → FP8/FP16 的转换与主循环融合，只要实际指令接收的是 FP8/FP16，就不属于原生 INT4 mixed-input MMA。NVFP4 也不能与 FP8/MXFP8 直接组成原生 block-scaled MMA。

判断一个低位模型是否使用原生低精度 MMA，应依次检查：

Weight 和 Activation 的实际数值格式，而不只看位宽；
目标架构、MMA 指令族，以及 A/B 类型、shape 和 layout 是否匹配；
Scale、zero-point、packing 是否符合 Kernel 契约，并确认编译产物实际生成的 MMA 指令。

LLM

Decoder Layer

FFN

Qwen

DeepSeek V3

Attention

量化基础

量化配置

典型方案

Weight-Activation

PTQ

Weight-Only

Weight-Activation

Prefix Cache

KV Cache

Recurrent State Cache

DP（Data Parallelism）

GEMM

投机采样

方式

Draft Model

优化

Prefill CUDA graph

模型实现

架构

硬件结构

逻辑结构

通信

CUDA

指令体系

矩阵指令

CUDA编程

通信

特性

专题

NVIDIA MMA ​

从 FMA 到 MMA，再到 GEMM ​

稠密乘法与 2:4 结构化稀疏乘法 ​

输入/累加/输出 ​

PTX 类型 ​

精度总表 ​

Ampere：FP16/BF16 与整数 MMA 基线 ​

Ada：mma.sync 加入 FP8 ​

Hopper：增加 WGMMA ​

Blackwell：FP4/FP6 与 block-scaled MMA ​

SM100：tcgen05.mma ​

SM120：扩展的 mma.sync ​

MXFP4 与 NVFP4 ​

量化配置如何映射到 MMA ​

参考 ​

NVIDIA MMA

从 FMA 到 MMA，再到 GEMM

稠密乘法与 2:4 结构化稀疏乘法

输入/累加/输出

PTX 类型

精度总表

Ampere：FP16/BF16 与整数 MMA 基线

Ada：`mma.sync` 加入 FP8

Hopper：增加 WGMMA

Blackwell：FP4/FP6 与 block-scaled MMA

SM100：`tcgen05.mma`

SM120：扩展的 `mma.sync`

MXFP4 与 NVFP4

量化配置如何映射到 MMA

参考