架构演进

本文直接对比数据中心 Ampere、Hopper 与 Blackwell 中影响大模型推理的硬件能力，主线是 Tensor Core 计算、Kernel 数据流水线、显存以及多 GPU 互联。代表产品分别取 A100、H100/H200 和 B200；RTX Blackwell 属于不同的 Compute Capability，不在本文范围内。

1. 架构总表

分类	对比项	Ampere	Hopper	Blackwell
基础	代表产品	A100	H100、H200	B200、GB200
基础	Compute Capability	8.0	9.0	10.0
基础	编译目标	`sm_80`	基线 `sm_90`；WGMMA 等专属能力使用 `sm_90a`	基线 `sm_100`；`tcgen05` 等能力使用 `sm_100a` 或 `sm_100f`
矩阵计算	Tensor Core	第三代	第四代	第五代
矩阵计算	原生浮点输入	TF32、BF16、FP16	TF32、BF16、FP16、FP8 E4M3/E5M2	TF32、BF16、FP16、FP8、FP6、FP4，以及 FP4/FP6/FP8 Mixed-input
矩阵计算	低精度缩放	主要由 Kernel 或框架管理 Tensor/Channel Scale	第一代 Transformer Engine 管理 FP8 Scale 与格式选择	第二代 Transformer Engine；支持 MXFP8、MXFP4、NVFP4 等 Micro-block Scaling
矩阵计算	主力 MMA 指令	Warp 级 `mma.sync`	Warpgroup 级异步 `wgmma.mma_async`，同时继承 `mma.sync`	异步 `tcgen05.mma`，同时继承传统 `mma.sync`
矩阵计算	MMA 协作范围	1 个 Warp、32 个线程	1 个 Warpgroup、4 个 Warp、128 个线程	单线程提交；Tensor Core 支持 1-CTA 或 2-CTA MMA
矩阵计算	MMA 累加器	Register File	Register File	Tensor Memory（TMEM）
矩阵计算	FP4 Weight	FP4 Weight Tile 通常先反量化为 FP16/BF16；INT4 MMA 只适合输入也满足整数指令契约的场景	WGMMA 不接收 FP4；Weight Tile 通常转为 FP8/FP16/BF16	`tcgen05.mma` 可直接接收 FP4 及其 Scale，原生支持 MXFP4、NVFP4
矩阵计算	2:4 结构化稀疏	`mma.sp`	`mma.sp`、`wgmma.mma_async.sp`	`mma.sp`、`tcgen05.mma.sp`
数据流水线	Global → Shared Memory	`cp.async`；线程计算地址并提交 Copy	TMA 根据 Tensor Map 完成 1D～5D 双向搬运和地址生成	延续 TMA，并与 TMEM、`tcgen05` 组成 Tensor Core 流水线
数据流水线	计算与搬运重叠	`cp.async` + Split Arrive/Wait Barrier	TMA + Transaction Barrier + 异步 WGMMA	TMA + Transaction Barrier + 异步 `tcgen05` + TMEM
数据流水线	跨 SM 协作	Block 之间通常经 Global Memory 交换数据	Thread Block Cluster、Cluster Barrier、Distributed Shared Memory、TMA Multicast	延续 Cluster/DSMEM；增加两个相邻 CTA/SM 共同执行的 2-CTA MMA
片上存储	Unified L1 / Shared Memory	192 KiB / SM	256 KiB / SM	256 KiB / SM
片上存储	最大 Shared Memory	164 KiB / SM	228 KiB / SM	228 KiB / SM
片上存储	Tensor Memory	无	无	每个 SM 新增专用 TMEM，保存 MMA Accumulator、Scale 及部分操作数
GPU Memory	代表产品 HBM	A100 80 GB HBM2e	H100 80 GB HBM3；H200 141 GB HBM3E	B200 180 GB HBM3E
GPU Memory	代表产品带宽	A100：约 2.0 TB/s	H100：3.35 TB/s；H200：4.8 TB/s	B200：最高 8 TB/s
GPU Memory	L2 Cache	A100：40 MB	H100：50 MB	数据中心 Blackwell：最高约 126 MB
多 GPU	NVLink	第三代，A100 最高 600 GB/s 双向	第四代，H100/H200 最高 900 GB/s 双向	第五代，B200 最高 1.8 TB/s 双向
推理	Prefill / 大 Batch GEMM	以 FP16/BF16/TF32 Tensor Core 为主；支持 INT8 与 2:4 稀疏	FP8、WGMMA 和 TMA 提高 GEMM、Attention 的计算与供数能力	FP4/FP6、`tcgen05`、TMEM 与 2-CTA MMA 进一步提高低精度矩阵吞吐
推理	Decode / 小 Batch	通常受权重读取和 KV Cache 访问限制	更高 HBM 带宽与 FP8/Weight-only Kernel 缓解带宽压力	更高 HBM 带宽、原生 FP4 与更大显存进一步减少权重流量并容纳更多 KV Cache
推理	Tensor Parallel / MoE	NVLink 3 承载 AllReduce、AllGather、All-to-All	NVLink 4 提高单机 TP 与 MoE 通信带宽	NVLink 5 继续提高带宽，并扩展机架级 NVLink Domain
Kernel	典型主循环	`cp.async → Shared Memory → mma.sync → Register`	`TMA → Shared Memory → wgmma → Register`	`TMA → Shared Memory → tcgen05 → TMEM → Register Epilogue`

表中的 HBM、L2 和 NVLink 数值来自代表产品，不是架构固定常量；不同容量、封装和系统形态会有不同配置。更完整的产品参数见卡型对比。

2. Tensor Core

三代架构的计算精度主线是：

text

Ampere：FP16 / BF16 / TF32
                 ↓
Hopper：增加 FP8 与 Transformer Engine
                 ↓
Blackwell：增加 FP4 / FP6、Micro-block Scaling 与原生 Block-scaled MMA

2.1 Ampere

Ampere 第三代 Tensor Core 为 LLM 推理提供 FP16、BF16、TF32、INT8 和 INT4 MMA，并加入 2:4 结构化稀疏。它的主力 PTX 接口是一个 Warp 协作执行的 mma.sync。

Ampere 的 INT4 Tensor Core 指令要求 A、B 都满足整数 MMA 的类型和布局契约。常见 W4A16 模型只有 Weight 是 INT4、Activation 仍是 FP16/BF16，因此 Kernel 通常在主循环内将 Weight Tile 反量化，再执行 FP16/BF16 MMA；这不属于原生 FP4 Tensor Core 计算。

2.2 Hopper

Hopper 第四代 Tensor Core 增加 FP8 E4M3 和 E5M2，并由第一代 Transformer Engine 管理张量的数值范围、Scale 与格式选择。FP8 同时减少矩阵输入的存储流量，并提高 Tensor Core 的低精度吞吐。

Hopper 新增 wgmma.mma_async。它将 MMA 协作范围从一个 Warp 扩大到由 4 个 Warp 组成的 Warpgroup，并允许从 Shared Memory Descriptor 读取矩阵 Tile。WGMMA 异步提交后，Kernel 可以继续推进 TMA 数据搬运或其他独立工作，读取 Accumulator 前再等待计算完成。

WGMMA 的最低原生浮点输入是 FP8，不直接接收 FP4。Hopper 上的 W4A16、W4A8 Kernel 仍需在 Tile 内完成解包、反量化或类型转换，再把 FP8、FP16 或 BF16 输入送入 Tensor Core。

2.3 Blackwell

Blackwell 第五代 Tensor Core 使用 tcgen05.mma 作为数据中心 SM100 的新主路径。它支持 FP4 E2M1、FP6 E2M3/E3M2、FP8 E4M3/E5M2，以及不同窄浮点类型之间的 Mixed-input MMA。

Blackwell 还把量化 Scale 纳入 Tensor Core 指令。将 K 维划分为多个量化 Group 后，Block-scaled MMA 计算为：

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

其中， $g$ 表示 K 维 Group， $S_{A, g}$ 和 $S_{B, g}$ 分别是 A、B 的 Block Scale。Scale 在每个 Group 进入累加前应用，而不是等完整 GEMM 结束后统一处理。

Blackwell LLM 推理常见的两种 FP4 格式为：

格式	数值	Block Scale	Scale 粒度	Tensor Scale
MXFP4	FP4 E2M1	UE8M0	每 32 个元素	无
NVFP4	FP4 E2M1	FP8 E4M3	每 16 个元素	FP32

NVFP4 的近似还原关系为：

x = x_{E 2 M 1} \times s_{b l o c k} \times s_{t e n s o r} .

更细的 Scale 粒度减少了一个大 Scale 同时覆盖极大值和极小值造成的量化误差，使 FP4 更适合直接用于大模型权重与 Activation。具体格式和指令签名见 NVFP4 与 NVIDIA MMA。

3. Kernel 数据流水线

三代架构的核心差别不仅是 Tensor Core 峰值，还包括数据能否持续送入 Tensor Core。

3.1 Ampere 流水线

text

HBM
  → cp.async
  → Shared Memory：A/B Tile
  → mma.sync
  → Register Accumulator
  → Epilogue
  → HBM

cp.async 可以绕过通用寄存器，把数据直接从 Global Memory 搬到 Shared Memory。但每次访问的地址仍由线程生成，Kernel 需要由多个线程循环提交 Copy。

3.2 Hopper 流水线

text

HBM
  → TMA：硬件地址生成与多维搬运
  → Shared Memory：A/B Tile
  → wgmma.mma_async
  → Register Accumulator
  → Epilogue
  → HBM

TMA 使用 Tensor Map 描述 Tensor 的基址、Shape、Stride 与布局，由硬件完成 1D～5D 地址生成和批量搬运。WGMMA 与 TMA 都是异步操作，因此 Hopper GEMM、Attention Kernel 通常把 Warp 分为 Producer 与 Consumer：Producer 提交 TMA，Consumer Warpgroup 执行 WGMMA。

Thread Block Cluster 还允许多个相邻 SM 上的 Block 使用 Cluster Barrier、Distributed Shared Memory 和 TMA Multicast。它主要用于多个计算 Tile 共享同一份输入，或者单个 Block 的 Shared Memory 容量不足的场景。

3.3 Blackwell 流水线

text

HBM
  → TMA
  → Shared Memory：A/B Tile
  → tcgen05.mma，支持 1-CTA / 2-CTA
  → Tensor Memory：Accumulator / Scale
  → tmem_load
  → Register：Epilogue Fragment
  → HBM

Blackwell 新增 Tensor Memory，将大型 MMA Accumulator 和 Block Scale 从通用 Register File 移到专用片上存储。tcgen05.mma 由单线程提交，并可让一个 CTA 独立执行，或者让两个相邻 CTA/SM 共同执行更大的矩阵 Tile。

TMEM 降低了 Tensor Core 主循环的寄存器压力，同时把 Mainloop Accumulator 与 Epilogue 的寄存器 Fragment 分开。Kernel 只在 Epilogue 需要处理某个输出子块时，才通过 tmem_load 将相应数据读入寄存器。

4. 推理影响

4.1 Prefill

Prefill 一次处理多个输入 Token，QKV Projection、Attention 与 MLP 通常形成较大的 GEMM，较容易进入计算受限区域：

Ampere：以 FP16/BF16 mma.sync 为主。
Hopper：FP8、WGMMA 与 TMA 提高 Tensor Core 计算和数据供给能力。
Blackwell：FP4/FP6、Block-scaled tcgen05、TMEM 与 2-CTA MMA 提高更低精度下的计算密度。

实际收益仍取决于 GEMM Shape、Batch、序列长度、量化格式和 Kernel 实现。只有框架真正选择相应指令路径时，架构能力才会转化为性能。

4.2 Decode

Decode 每步只新增少量 Token，矩阵 M 维较小，权重和 KV Cache 的读取更容易成为瓶颈：

显存带宽决定每个 Token 可以多快读取模型权重和 KV Cache。
显存容量决定单卡可以容纳的模型、KV Cache 和并发请求数量。
FP8、FP4 与 Weight-only 量化通过减少权重字节数降低每步 Decode 的数据搬运量。
Blackwell 原生 NVFP4/MXFP4 可以避免先把完整 Weight Tile 展开成 FP16，再进入 Tensor Core。

因此，Decode 不应只比较 Tensor Core FLOPS。H200 相比 H100 没有更换 Hopper 架构，但更大的 HBM3E 容量与更高带宽仍能显著影响大模型推理，这也是产品参数与架构能力需要分开理解的原因。

4.3 多 GPU 推理

大模型无法放入单卡或需要更低延迟时，Tensor Parallel、Pipeline Parallel、Expert Parallel 会引入 AllReduce、AllGather、Send/Recv 或 All-to-All：

text

Ampere NVLink 3：最高 600 GB/s
        ↓
Hopper NVLink 4：最高 900 GB/s
        ↓
Blackwell NVLink 5：最高 1.8 TB/s

更高 NVLink 带宽可以降低通信占比。Blackwell 还通过新的 NVSwitch 扩大 NVLink Domain，使更多 GPU 可以在高速 All-to-All Fabric 内组成 Tensor Parallel 或 Expert Parallel Group。

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编程

通信

特性

专题

架构演进 ​

1. 架构总表 ​

2. Tensor Core ​

2.1 Ampere ​

2.2 Hopper ​

2.3 Blackwell ​

3. Kernel 数据流水线 ​

3.1 Ampere 流水线 ​

3.2 Hopper 流水线 ​

3.3 Blackwell 流水线 ​

4. 推理影响 ​

4.1 Prefill ​

4.2 Decode ​

4.3 多 GPU 推理 ​

5. 参考 ​

架构演进

1. 架构总表

2. Tensor Core

2.1 Ampere

2.2 Hopper

2.3 Blackwell

3. Kernel 数据流水线

3.1 Ampere 流水线

3.2 Hopper 流水线

3.3 Blackwell 流水线

4. 推理影响

4.1 Prefill

4.2 Decode

4.3 多 GPU 推理

5. 参考