架构演进：Volta 到 Blackwell，以及别家的路¶
推理领域这几年的热点——FP8、FP4、MoE、投机解码、PD 分离、MLA——看上去是算法和系统的创新，背后其实是同一件硬件事实：**算力涨得比带宽快**。这一章沿着 Volta、Turing、Ampere、Ada、Hopper、Blackwell 一路看过来，每一代新增了什么、为什么加，再用两个小模型算清楚"算力和带宽的增速不同"给推理带来了什么后果，最后对照 AMD 和昇腾，看看别家怎么解同一道题。
自测：能答上来就可以跳过本章
- 从 V100 到 B200，同精度的算力涨了多少倍，带宽涨了多少倍？屋脊点怎么变的？
- 为什么每一代都要加一种更低的精度？FP4 这种只有 16 个可表示值的格式怎么可能用来算矩阵乘？
- decode 时，batch 要多大才能从带宽受限变成算力受限？加大 batch 能让注意力也变成算力受限吗？
- Hopper 相对 Ampere，对推理影响最大的三个新特性是什么？Blackwell 呢？
- AMD 的 MI300X 和华为昇腾在架构上与 NVIDIA 最大的不同是什么？
自测参考答案（先自己答，再展开对照）
- BF16/FP16 稠密算力从 125 到 2250 TFLOPS，18 倍；带宽从 0.9 到 8 TB/s，约 9 倍。屋脊点（算力÷带宽）从 139 涨到 281，也就是说，同样从显存读 1 字节，现在要做两倍多的运算才能把算力喂饱。算上每代新增的低精度，最低精度的屋脊点从 139 涨到 1125。
- 因为加宽数据通路、堆更多计算单元的代价越来越高，而位宽减半几乎免费地让吞吐翻倍、让权重和 KV 占的显存和带宽减半。FP4 能用，靠的是块缩放：每 16 或 32 个元素共享一个缩放因子（NVFP4 用 FP8 的因子，MXFP4 用 8 位指数），每个块内部的动态范围很小，4 比特够用；累加仍然在 FP32 里做。
- 权重是被整个 batch 共用的，算术强度 = 2 × batch ÷ 每个参数的字节数，BF16 权重时要达到 H100 的屋脊点 295，batch 要到 295；FP8 权重减半到 148，4 比特权重 74。但注意力不一样：KV 是每个请求自己的，加大 batch 并不会让某个请求的 KV 被复用，它的算术强度只等于"每个 KV 头被多少个查询头共用"，MHA 是 1、GQA 8 组是 8、MLA 吸收后约 128，都远低于屋脊点。所以 decode 的注意力部分永远是带宽受限的，只能靠减少 KV 的字节数（GQA、MLA、量化 KV）来加速。
- Hopper：FP8 Tensor Core（算力翻倍、KV 和权重减半）、TMA + 线程块集群（异步搬运，撑起 FlashAttention-3 这类 warp 专门化 kernel）、NVLink 4 与 900 GB/s 的 NVSwitch（撑起大规模张量并行和 EP）。Blackwell：FP4 与块缩放格式、Tensor Memory 与单线程发起的 tcgen05（让更大的矩阵块可行）、NVLink 5 与 NVL72 机柜（一个 NVLink 域 72 张卡，大规模专家并行不再走网络）。
- AMD MI300X：一块 GPU 由多个计算小芯片（XCD）+ 多个 I/O 小芯片拼成，显存 192 GB、带宽 5.3 TB/s（同期 H100 是 80 GB / 3.35 TB/s），另有 256 MB 的 Infinity Cache 挡在显存前面；软件栈是 ROCm/HIP，和 CUDA 源码级接近但生态差一截。昇腾是达芬奇架构：一个 AI Core 里有专做矩阵乘的 Cube 单元、做向量的 Vector 单元和 Scalar 单元，编程模型不是 SIMT 而是显式的多级缓冲流水（Ascend C），软件栈 CANN + HCCL。
一张表：每一代加了什么¶
| 架构 | 计算能力 | 代表卡 | 新增的硬件能力 | 对推理意味着什么 |
|---|---|---|---|---|
| Volta | 7.0 | V100 | 第一代 Tensor Core（FP16）、独立线程调度 | 矩阵乘有了专用单元；warp 内不能再假设天然同步 |
| Turing | 7.5 | T4 | INT8 / INT4 Tensor Core | 量化推理的开端；T4 至今是最低的部署门槛 |
| Ampere | 8.0 | A100 | BF16、TF32、2:4 稀疏、cp.async、164 KB 共享内存、MIG |
BF16 成为训练与推理的主流格式；异步拷贝让多级流水成为可能 |
| Ada | 8.9 | L4、L40S | FP8 Tensor Core、大 L2（72 MB） | 便宜的推理卡开始支持 FP8，但没有 HBM，带宽是短板 |
| Hopper | 9.0 | H100、H200 | FP8 + Transformer Engine、TMA、线程块集群与分布式共享内存、wgmma、NVLink 4 + NVSwitch、DPX |
FlashAttention-3 式的 warp 专门化；FP8 推理普及；机内 8 卡全互联撑起 TP 和 EP |
| Blackwell | 10.0 | B200、B300、GB200 | FP6 / FP4 与块缩放格式、Tensor Memory、tcgen05（单线程发起、双 SM 协作）、双芯片封装、NVLink 5 与 NVL72、第二代 Transformer Engine |
FP4 推理；一个 NVLink 域 72 张卡，大规模 EP 不再走网络 |
（RTX 50 系列的计算能力是 12.0，虽然也叫 Blackwell，但没有 tcgen05 和 Tensor Memory，编程模型更接近 Ada。）
三条主线贯穿其中：
- 精度一路往下：FP16 → INT8 → BF16/TF32 → FP8 → FP6/FP4。每降一档，Tensor Core 吞吐翻倍，权重和 KV 占的显存与带宽减半；
- 搬数据的方式一路异步化：普通读取 →
cp.async（绕开寄存器）→ TMA（一条指令搬一整块）；同时同步原语从__syncthreads()走向 mbarrier 和线程块集群； - 规模一路往上：单芯片 → 双芯片封装 → 机内 8 卡 NVSwitch → 机柜 72 卡 NVLink 域。
算力涨得比带宽快¶
把几代 SXM 卡的算力、带宽、屋脊点放在一起：
先看一眼各代的屋脊点和 decode 要攒的 batch：
# 每一代数据中心卡（SXM 形态）的算力、带宽各涨了多少，屋脊点（算力÷带宽）怎么变。
# 屋脊点的含义：每从显存读 1 字节，要做多少次运算才能把算力喂饱
gens = [
# 名称, 年份, FP16/BF16 稠密 TFLOPS, 该代最低精度的稠密 TFLOPS, 精度名, 带宽 GB/s, 显存 GB
("V100", 2017, 125, 125, "FP16", 900, 32),
("A100", 2020, 312, 312, "BF16", 2039, 80),
("H100", 2022, 989, 1979, "FP8", 3350, 80),
("B200", 2024, 2250, 9000, "FP4", 8000, 180),
]
print("GPU 年份 BF16算力 较上代 带宽 较上代 BF16屋脊点 最低精度算力 屋脊点")
prev = None
for name, year, bf16, low, dtype, bw, mem in gens:
up = (f"{bf16 / prev[0]:5.1f}x", f"{bw / prev[1]:5.1f}x") if prev else (" -", " -")
print(f"{name:5s} {year} {bf16:7.0f} {up[0]} {bw / 1000:6.2f} TB/s {up[1]} "
f"{bf16 * 1e12 / (bw * 1e9):9.0f} {low:11.0f} {dtype:4s} {low * 1e12 / (bw * 1e9):7.0f}")
prev = (bf16, bw)
print()
print("同一个 8B 模型（BF16 权重 16 GB，KV 每 token 128 KB）在各代上的理论下限：")
print("GPU decode 一步(batch=1) prefill 4K(BF16，算力用满一半) 能放的 KV")
for name, year, bf16, low, dtype, bw, mem in gens:
decode = 16e9 / (bw * 1e9) * 1e3 # 读一遍权重
prefill = 2 * 8e9 * 4096 / (bf16 * 1e12 * 0.5) * 1e3
kv = (mem * 1e9 - 16e9) / (128 * 1024) / 1e3
print(f"{name:5s} {decode:12.1f} ms {prefill:26.0f} ms {kv:9.0f}K token")
GPU 年份 BF16算力 较上代 带宽 较上代 BF16屋脊点 最低精度算力 屋脊点
V100 2017 125 - 0.90 TB/s - 139 125 FP16 139
A100 2020 312 2.5x 2.04 TB/s 2.3x 153 312 BF16 153
H100 2022 989 3.2x 3.35 TB/s 1.6x 295 1979 FP8 591
B200 2024 2250 2.3x 8.00 TB/s 2.4x 281 9000 FP4 1125
同一个 8B 模型（BF16 权重 16 GB，KV 每 token 128 KB）在各代上的理论下限：
GPU decode 一步(batch=1) prefill 4K(BF16，算力用满一半) 能放的 KV
V100 17.8 ms 1049 ms 122K token
A100 7.8 ms 420 ms 488K token
H100 4.8 ms 133 ms 488K token
B200 2.0 ms 58 ms 1251K token
看两件事：
- 同精度（BF16）下，算力涨了 18 倍，带宽涨了 9 倍，屋脊点从 139 涨到 281。V100 时代，每从显存读 1 字节做 139 次运算就能喂饱算力；现在要 281 次。如果把每代新增的最低精度也算上，H100 的 FP8 是 591、B200 的 FP4 是 1125，差距还要再大一截。
- decode 和 prefill 的改善完全不同步：8B 模型 batch 为 1 的 decode 一步，从 V100 的 17.8 ms 到 B200 的 2.0 ms，快了 9 倍（就是带宽的倍数）；同一个模型 prefill 4K，从 1049 ms 到 58 ms，快了 18 倍。硬件每一代都在把"计算"变得更便宜，而"把数据搬进来"只是稍微便宜了一点。
推理系统这几年的主要努力，几乎都可以读成"把带宽的账做小、把算力的账做满"：
| 方向 | 做的是什么 | 对应上面哪个数 |
|---|---|---|
| 权重量化（FP8、AWQ、GPTQ、FP4） | 让每个参数占的字节更少 | 直接减少 decode 每一步读的字节 |
| GQA、MLA、KV 量化 | 让每个 token 的 KV 更小 | 同上，且长上下文时占比更大 |
| 连续批处理、大 batch | 让读进来的权重被更多 token 共用 | 提高算术强度，往屋脊点靠 |
| MoE | 每个 token 只激活一部分专家 | 参数量涨了，但每步读的字节没有同比例涨 |
| 投机解码 | 一步验证多个 token | 一次读权重、多产出几个 token |
| PD 分离 | prefill（算力受限）和 decode（带宽受限）分开部署 | 两类卡各用其长 |
| CUDA Graph、重叠调度、megakernel | 减少 GPU 空转 | 让已经很贵的算力不被 CPU 拖住 |
batch 要多大才够¶
decode 一步里，权重被这一批的所有 token 共用：读 1 字节权重要做 2 × batch 次运算。要达到屋脊点，batch 要多大？
# decode 要多大的 batch 才能从"带宽受限"变成"算力受限"？
# 一步里每个参数：读 w 字节，做 2 × batch 次运算（batch 个 token 共用这一份权重）
# 所以算术强度 = 2 × batch / w（FLOP/字节），它要达到屋脊点 算力÷带宽
gens = [("A100", 312, 2039), ("H100", 989, 3350), ("H100 FP8", 1979, 3350),
("B200", 2250, 8000), ("B200 FP4", 9000, 8000)]
weights = [("BF16 权重", 2), ("FP8 权重", 1), ("4 比特权重", 0.5)]
print("GPU 屋脊点 " + " ".join(f"{n}" for n, _ in weights))
for name, tf, bw in gens:
ridge = tf * 1e12 / (bw * 1e9)
row = "".join(f"{ridge * w / 2:11.0f}" for _, w in weights)
print(f"{name:9s} {ridge:6.0f}{row}")
print()
print("权重之外，decode 还要读 KV。KV 是每个请求自己的，加大 batch 不会让它被复用：")
print("每个 KV 元素被组里的 g 个查询头各用一次乘加，所以注意力部分的算术强度 = g（BF16 KV，FLOP/字节）")
print("每组查询头数 算术强度 和 H100 BF16 屋脊点 295 相比 注意力方案")
for name, g in [("MHA（每个头一份 KV）", 1), ("GQA 8 组", 8), ("MLA（吸收后，128 个头共用）", 128)]:
print(f"{g:11d} {g:9d} {295 / g:20.0f} 倍之差 {name}")
GPU 屋脊点 BF16 权重 FP8 权重 4 比特权重
A100 153 153 77 38
H100 295 295 148 74
H100 FP8 591 591 295 148
B200 281 281 141 70
B200 FP4 1125 1125 562 281
权重之外，decode 还要读 KV。KV 是每个请求自己的，加大 batch 不会让它被复用：
每个 KV 元素被组里的 g 个查询头各用一次乘加，所以注意力部分的算术强度 = g（BF16 KV，FLOP/字节）
每组查询头数 算术强度 和 H100 BF16 屋脊点 295 相比 注意力方案
1 1 295 倍之差 MHA（每个头一份 KV）
8 8 37 倍之差 GQA 8 组
128 128 2 倍之差 MLA（吸收后，128 个头共用）
H100 上 BF16 权重要 batch 295 才算力受限，FP8 权重 148，4 比特权重 74——量化不但省显存，还降低了"变成算力受限"的门槛。反过来，B200 的 FP4 屋脊点 1125，即使权重也是 4 比特，batch 也要 281，实际部署里很难在单卡上凑到，所以 Blackwell 上更依赖大 batch、MoE 和投机解码把算力用起来。
注意力部分则完全不同：KV 是每个请求自己的，加大 batch 不会让它被复用。它的算术强度只取决于"一个 KV 头被多少个查询头共用"，MHA 是 1、GQA 8 组是 8、MLA 吸收后约 128，全都远低于屋脊点。decode 的注意力永远是带宽受限的，唯一的出路是让 KV 更小（GQA、MLA、量化 KV、稀疏注意力），这就是各家模型在注意力结构上反复折腾的原因（见推理手册的 MLA 推理和稀疏注意力）。
低精度为什么可行¶
FP4 只有 16 个可表示的值，怎么用来算矩阵乘？关键是块缩放：一组元素共享一个缩放因子，实际的值是"4 比特的尾数 × 这个块的缩放因子"。
| 格式 | 位宽 | 每块多少个元素 | 缩放因子 | 用在哪 |
|---|---|---|---|---|
| FP8 E4M3 / E5M2 | 8 | 整个张量或每行 | FP32 | Hopper 起的推理与训练 |
| 分块 FP8 | 8 | 128 × 128 或 1 × 128 | FP32 | DeepSeek-V3 的训练与推理 |
| MXFP8 / MXFP6 / MXFP4 | 8 / 6 / 4 | 32 | 8 位指数（2 的幂） | 开放计算标准，Blackwell 原生支持 |
| NVFP4 | 4 | 16 | FP8 E4M3，另加一个张量级 FP32 | Blackwell 推理的主推格式 |
块越小，块内的动态范围越窄，低位宽就越够用；代价是缩放因子本身要占空间（NVFP4 每 16 个元素一个 FP8 因子，等于每个元素多 0.5 比特）和额外的处理。Blackwell 的 Tensor Core 直接认这些格式，硬件在做乘加时自动应用缩放因子，不用软件反量化。
累加仍然是 FP32：矩阵乘要把 K 个乘积加起来，低精度累加的误差会随 K 快速累积。所以"FP4 推理"指的是输入和权重是 FP4，中间累加和输出都还是高精度。精度怎么保住，见推理手册的量化部署实战和超大 MoE 的低比特推理。
别家怎么解同一道题¶
AMD Instinct（CDNA 架构）。MI300X 用小芯片（chiplet）拼成：8 个计算小芯片（XCD，每个 38 个计算单元 CU）加 4 个 I/O 小芯片，堆 8 个 HBM3，显存 192 GB、带宽 5.3 TB/s，同期 H100 只有 80 GB / 3.35 TB/s，大显存是它最大的卖点（一张卡就能放下 70B 的 BF16 权重）。I/O 小芯片上还有 256 MB 的 Infinity Cache，挡在显存前面。计算单元里的 Matrix Core 对应 Tensor Core，MI300X 的 BF16 稠密算力 1307 TFLOPS、FP8 2615 TFLOPS。软件栈是 ROCm：HIP 在源码层面和 CUDA 高度相似（hipify 能自动转大部分代码），RCCL 对应 NCCL，vLLM 和 SGLang 都支持。差距主要在生态：算子库的覆盖和成熟度、新模型的首日支持、调试与性能工具。后续的 MI355X 加了 FP6 / FP4 和更大的 HBM3e。
华为昇腾（达芬奇架构）。它不是 SIMT：一个 AI Core 里有三种单元——Cube（矩阵乘，一条指令做 16×16×16）、Vector（向量运算）、Scalar（标量与控制），数据在 Global Memory、L1、L0A/L0B/L0C 几级缓冲之间显式搬运。写 kernel 用 Ascend C，要自己安排"搬进来、算、搬出去"的多级流水，更像在写 DSP，而不像写 CUDA。软件栈是 CANN（对应 CUDA），集合通信是 HCCL（对应 NCCL），推理引擎有 MindIE，同时 vLLM 和 SGLang 都有对应的硬件后端插件（见多硬件 platform 抽象）。国内的推理岗位经常要求了解这一套。
Google TPU。用的是脉动阵列：一个大的乘加阵列，数据像波一样从一边流进去、结果从另一边流出来，控制逻辑极少，非常适合大矩阵乘；没有 SIMT，也没有共享内存这一层，编程通过 XLA 编译器而不是手写 kernel。TPU 的互连是 3D 环面拓扑，不需要交换机就能组成很大的集群。
三家的共同点是：都在同一个方向上走——加大矩阵单元、降低精度、加大片上缓冲、加宽互连，因为屋脊点上升这件事对谁都一样。
怎么读一张新发布的规格表¶
新卡发布时，按这个顺序看，就能很快判断它对推理意味着什么：
- 显存容量和带宽 → decode 的上限（读满一遍要多久）、能放多大的模型和多少 KV；
- 各精度的稠密算力（别看稀疏的数字）→ prefill 的上限、屋脊点；
- 屋脊点 = 算力 ÷ 带宽 → batch 要多大才算力受限，这张卡适合 prefill 还是 decode；
- 互连：NVLink / 以太网带宽、一个域里有多少张卡 → 能做多大的 TP、EP，跨机要不要走网络；
- 新的硬件单元（新精度、新的搬运或同步机制）→ 需要新版 CUDA、新版算子库，以及大概多久会被 vLLM / SGLang 支持；
- 形态与功耗：SXM 还是 PCIe、整机功耗和散热，决定了能不能装进现有机房。
怎么讲清楚
讲架构演进或"为什么现在都在做量化和 MoE"：核心事实是算力涨得比带宽快——V100 到 B200，同精度算力 18 倍、带宽 9 倍，屋脊点从 139 涨到 281，算上新增的低精度到 1125。于是 decode（带宽受限）改善得比 prefill（算力受限）慢得多，推理系统的主要工作就是把带宽的账做小、把算力的账做满：权重与 KV 量化、GQA/MLA、大 batch 与连续批处理、MoE、投机解码、PD 分离。定量地说：BF16 权重要 batch 295 才在 H100 上变成算力受限，FP8 148、4 比特 74；而注意力部分的算术强度只等于每组查询头数（MHA 1、GQA 8、MLA 约 128），加大 batch 没用，只能减少 KV 字节。再按代讲新特性：Ampere 的 BF16 与 cp.async、Hopper 的 FP8 与 TMA 与集群、Blackwell 的 FP4 块缩放与 Tensor Memory 与 NVL72。最后能提一句别家：AMD 靠大显存和 ROCm，昇腾是达芬奇架构、Cube/Vector 单元加显式多级流水，TPU 是脉动阵列加 XLA。
练习¶
1. 屋脊点与选卡。 H20 的 BF16 稠密算力 148 TFLOPS、带宽 4.0 TB/s；H100 SXM 是 989 TFLOPS、3.35 TB/s。各自的屋脊点是多少？一个 BF16 权重的模型在两张卡上 decode，batch 要多大才算力受限？为什么说 H20 适合 decode、不适合 prefill？
参考答案
H20 的屋脊点 148 / 4.0 = 37，H100 是 295。BF16 权重时，batch 分别要到 37 和 295 才变成算力受限。H20 的带宽比 H100 还高 20%，decode（读一遍权重）比 H100 快；但算力只有 H100 的 15%，prefill 慢 6 倍多。所以异构的 PD 分离里，用算力强的卡做 prefill、带宽高的卡做 decode，正是按这两个数分工。
2. 量化的两笔账。 一个 70B 模型，BF16 权重 140 GB。把权重量化到 4 比特后：decode 一步读的字节少了多少？prefill 的时间会不会同比例缩短？为什么？
参考答案
权重从 140 GB 降到 35 GB，decode 一步读的字节少了四分之三，按带宽算快 4 倍（KV 的部分不变，长上下文时收益会被摊薄）。prefill 不会同比例缩短：prefill 是算力受限的，权重只读一遍，主要时间花在矩阵乘上；如果算子是"先反量化成 BF16 再乘"，算力和原来一样，prefill 几乎不变；只有当硬件原生支持这个低精度的 Tensor Core（比如 FP8、FP4），算力真的翻倍，prefill 才会变快。这就是"权重量化主要加速 decode，激活也量化才加速 prefill"的原因。
3. 一代的账。 假设下一代卡算力翻倍、带宽只涨 30%，其余不变。对以下三件事各有什么影响：（a）batch 为 1 的 decode 延迟；（b）prefill 吞吐；（c）要达到算力受限所需的 batch？部署时你会怎样调整策略？
参考答案
（a）decode 延迟只随带宽改善，快约 23%；（b）prefill 吞吐翻倍；（c）屋脊点变成 2 / 1.3 ≈ 1.54 倍，所需 batch 也涨到 1.54 倍。
策略上：更激进地量化权重和 KV（把带宽的账做小）、把 batch 做得更大（更多并发、更长的排队窗口，代价是延迟）、更值得做投机解码（一次读权重多产出几个 token）、prefill 和 decode 分开部署并给 decode 配带宽更高的卡。
4. 判断一个说法。 "我们的模型从 BF16 换成 FP8，算力翻倍，所以推理速度会翻倍。" 这句话在什么场景下成立，什么场景下不成立？
参考答案
prefill（或大 batch 的计算受限阶段）成立：FP8 Tensor Core 的吞吐是 BF16 的两倍，如果算子确实用 FP8 计算且能达到相近的利用率，时间接近减半。
batch 小的 decode 不成立：这时是带宽受限，速度取决于每一步读多少字节。如果只是计算用 FP8 而权重仍以 BF16 存放，读的字节没变，速度不变；如果权重也存成 FP8，读的字节减半，这才快一倍——加速来自"权重变小"，而不是"算力翻倍"。长上下文时 KV 占比大，权重那部分的收益还会被进一步摊薄。
小结¶
- 三条主线：精度一路往下（FP16 → FP8 → FP4，靠块缩放）、搬数据一路异步化（cp.async → TMA）、规模一路往上（单芯片 → 双芯片 → NVSwitch → NVL72）。
- V100 到 B200，同精度算力 18 倍、带宽 9 倍，屋脊点从 139 到 281（算上低精度到 1125）：prefill 改善得比 decode 快得多。
- 量化、GQA/MLA、大 batch、MoE、投机解码、PD 分离，都是在"把带宽的账做小、把算力的账做满"。
- BF16 权重要 batch 295 才在 H100 上算力受限，FP8 148、4 比特 74；注意力的算术强度只等于每组查询头数，decode 的注意力永远带宽受限。
- AMD 靠小芯片、大显存和 ROCm；昇腾是达芬奇架构的 Cube/Vector 单元加显式多级流水；TPU 是脉动阵列加 XLA。大家在同一个方向上走。
- 读规格表的顺序：容量与带宽 → 稠密算力 → 屋脊点 → 互连 → 新单元 → 形态与功耗。
发现错误或有更好的写法？在 GitHub 上编辑此页提交问题