TMA、WGMMA、FP8/FP4 等硬件能力解释

这几个能力不属于同一层，但经常一起出现：

TMA      = 数据搬运能力
WGMMA    = 矩阵乘计算能力
FP8/FP4  = 数字表示 / 低精度计算能力

放到 GPU 执行链路里，可以这样理解：

HBM 显存
  ↓
TMA：把 Q/K/V、权重 tile 高效搬到 shared memory
  ↓
Shared Memory
  ↓
WGMMA：让更多 warp 协同调用 Tensor Core 做大矩阵乘
  ↓
Tensor Core
  ↓
FP8/FP4：用更少 bit 表示数，提高吞吐、降低带宽

---

1. TMA：Tensor Memory Accelerator

TMA 是 Hopper 上的张量搬运硬件。

以前在 A800/A100 这种 Ampere 上，数据从 global memory 搬到 shared memory，通常靠线程执行 cp.async。

也就是说，线程既要算地址，又要搬数据。

Hopper 的 TMA 更像一个专用 DMA：

以前：
很多线程自己算地址、自己搬 tile

TMA：
一个线程提交 tensor copy descriptor
硬件负责多维 tensor 的地址计算和搬运
其他线程继续算

它解决的是：

Tensor Core 很快，但数据喂不上去。

所以 TMA 的核心价值是：

让搬数据和矩阵计算更好地重叠。

在 FlashAttention-3 里，TMA 很重要，因为 attention 要不断搬：

Q tile
K tile
V tile

如果每次 Tensor Core 都在等 K/V 从 HBM 搬进 shared memory，算力就浪费了。

TMA 就是为了减少这种等待。

---

2. WGMMA：Warp-Group Matrix Multiply-Accumulate

WGMMA 是 Hopper 的新矩阵乘指令路径。

Ampere 上常见的是 mma.sync，通常一个 warp，也就是 32 个线程协同调用 Tensor Core。

Hopper 的 WGMMA 是：

warp group = 4 个 warp = 128 个线程

也就是 128 个线程一起发起更大粒度的矩阵乘。

可以这样理解：

Ampere mma.sync：
32 个线程一组喂 Tensor Core

Hopper wgmma.mma_async：
128 个线程一组喂 Tensor Core
而且可以异步提交

它解决的是：

Tensor Core 更强了，需要更大 tile、更异步的调度方式。

WGMMA 的价值：

1. 更大矩阵 tile
2. 更高 Tensor Core 利用率
3. 可以和 TMA 数据搬运形成流水
4. 更适合 FlashAttention / 大 GEMM

所以 WGMMA 不是一种数据类型，而是：

矩阵乘指令机制。

---

3. FP8：8-bit 浮点

FP8 是 8 位浮点数，不是 INT8。

常见格式有：

E4M3：4 bit 指数，3 bit 尾数
E5M2：5 bit 指数，2 bit 尾数

大概理解：

E4M3：
精度稍好，动态范围小一点

E5M2：
动态范围大一点，精度粗一点

FP8 的本质是：

FP16/BF16：一个数 16 bit
FP8：      一个数 8 bit

好处：

显存占用减半
HBM 带宽压力减半
Tensor Core 吞吐更高
KV cache / activation / weight 更省

代价：

精度下降
动态范围变小
必须做 scale
有些层不能直接 FP8

所以 FP8 真正难的是 scale：

按 tensor scale？
按 channel scale？
按 block scale？
哪些层用 FP8？
哪些层保留 BF16/FP16？

这就是 Hopper Transformer Engine 的价值。

---

4. FP4：4-bit 浮点

FP4 更激进：

FP16：16 bit
FP8 ： 8 bit
FP4 ： 4 bit

理论上 FP4 相比 FP16：

显存占用约 1/4
带宽压力约 1/4
低精度 Tensor Core 吞吐更高

但 4 bit 太少，不能简单粗暴直接用。

它通常要配合 microscaling：

真实值 ≈ 4-bit 编码值 × 局部 scale

也就是一小组数共享一个 scale，每个数只存 4 bit。

这和 INT4 量化有点像，但路线不同：

INT4：
整数低精度量化

FP4 / NVFP4：
浮点低精度 + 硬件 Tensor Core 支持 + microscaling

FP4 主要是 Blackwell 这一代重点强化的能力。

---

5. 为什么它们对大模型重要？

大模型推理的瓶颈经常不是单纯算力，而是：

显存带宽
KV cache 读写
小 batch decode
数据搬运
kernel launch
通信

TMA、WGMMA、FP8/FP4 分别从不同角度解决：

TMA：
降低搬数据成本，让计算和访存重叠

WGMMA：
让 Tensor Core 以更大粒度、更高效率做矩阵乘

FP8/FP4：
减少每个数的 bit 数，降低显存带宽和存储压力

放到 FlashAttention：

QK^T              → WGMMA 加速矩阵乘
K/V 搬运           → TMA 加速数据搬运
P @ V             → WGMMA 加速矩阵乘
Q/K/V/中间数据     → FP8 降低带宽

所以 Hopper/Blackwell 的优势不是一个点，而是一整套：

更快搬数据
更强矩阵乘
更低精度格式
更好流水调度

---

6. 和 A800 的关系

A800 是 Ampere SM80，它有：

Tensor Core
FP16 / BF16 / TF32 / INT8
cp.async
shared memory tiling
FlashAttention-2 可用路径

但它没有：

TMA
WGMMA
FP8 Tensor Core
FP4 / NVFP4 Tensor Core

所以你在 A800 上优化时，不应该按 H100/Blackwell 的思路走。

A800 更应该做：

1. 把小 GEMV 变成更大的 GEMM
2. W1/W3 concat GEMM，而不是手写大而全 kernel
3. 用 FlashAttention-2，不是 FA3
4. 减少 kernel launch：CUDA Graph / 融合算子
5. MoE 用 grouped GEMM，减少 expert 小矩阵碎片
6. 通信和计算 overlap
7. decode 阶段重点优化 KV cache 和 persistent kernel

---

7. 一句话总结

TMA 是搬数据，WGMMA 是算矩阵，FP8/FP4 是降低每个数的 bit 数。

A800 缺的是 Hopper/Blackwell 这些新硬件路径，所以它的优化重点不是追 FP8/FP4，而是：

把计算形状改造成 A800 擅长的 BF16/FP16 大 GEMM。