跳转到内容

NVIDIA Ampere — 第三代 Tensor Core 加 TF32 / BF16 / FP64,结构化稀疏 + MIG 重写大模型时代硬件假设

待复核

Ampere 是 NVIDIA 2020 年发布的第八代 GPU 架构,旗舰是 GA100(A100)——一颗专为数据中心训练 / 推理 / HPC 三合一设计的硅片。它在一代之内做了四件改写假设的事:第三代 Tensor Core 加 TF32 / BF16 / FP64 + 2:4 结构化稀疏 + MIG 多实例 GPU + 第三代 NVLink 600 GB/s——最直接的后果是:GPT-3 证明千亿模型门槛后,OPT / BLOOM / LLaMA-1 等非 TPU 路线大模型训练的事实硬件很快变成 A100 集群。

日常类比:volta-architecture-2017 在数据中心车间装了第一台拼乐高 4×4 的机器(Tensor Core FP16);turing-architecture-2018 把它小型化下放消费市场加 INT8。Ampere 是把这台机器升级成”自带四种乐高规格 + 能切成 7 台小机器”——既能承担千亿模型训练里最重的矩阵计算,也可以切成 7 份给 7 个推理客户。Volta 让”训练用 GPU”成立,Ampere 让”训练 / 推理 / 多租户切片”在同一颗硅片上同时成立。

落到硅片:GA100 = 542 亿晶体管、TSMC 7nm N7、826 mm²、108 SM × 64 FP32 = 6912 FP32 核 + 432 第三代 Tensor Core、40 / 80 GB HBM2 / HBM2e、显存带宽 1555 / 2039 GB/s、NVLink 3.0 600 GB/s、PCIe Gen4、SXM4 400 W / PCIe 250 W。代表卡:A100 SXM4 80GB(DGX A100 标配)、A100 PCIe、A30、A40、GeForce RTX 3090 / 3080(GA102)Compute Capability sm_80(GA100)/ sm_86(消费 GA10x)。

不理解 Ampere,下面这些事都没法解释:

  • 为什么 GPT-3 之后的非 TPU 大模型训练默认看 A100 集群——TF32 + BF16 + 600 GB/s NVLink 是事实门槛
  • 为什么 PyTorch 1.7+ 默认 matmul 走 TF32——不改一行代码 FP32 训练快 8×
  • 为什么 云厂商出现”1/7 A100”实例(AWS / Azure / GCP)——MIG 把单卡按硬件切成 7 份
  • 为什么 LLM 训练用 BF16 而不是 FP16——A100 起 BF16 算力 = FP16,且动态范围 ×256
  • 为什么 HPC(流体 / 气候 / 量子化学)2021 后大量上 GPU——FP64 Tensor Core 让双精度也能用 Tensor Core,HPC 算力 ×2

Ampere 在 volta-architecture-2017 / turing-architecture-2018 之上做了 四件事

  1. 第三代 Tensor Core:在 V100 第一代 FP16/FP32、Turing 第二代 INT8/INT4 之上,新增 TF32 / BF16 / FP64。FP16 算力对 V100 ×2.5(312 vs 125 TFLOPS);INT8 ~624 TOPS。TF32(19-bit)= 8-bit 指数 + 10-bit 尾数 + 1 符号——指数和 FP32 一样、尾数和 FP16 一样,作为 FP32 的硬件替身自动跑,用户代码 torch.matmul(a, b) 不改但走 Tensor Core,156 TFLOPS vs 19.5 TFLOPS FP32 = ×8。

  2. 2:4 结构化稀疏:每 4 个权重必须有 2 个为零(训练时 mask),推理时硬件跳过零值,Tensor 算力直接 ×2——TF32 312 / BF16 624 / INT8 1248 TOPS。意义:首次硬件化稀疏,但要求训练阶段就引入 2:4 约束,不是”训完压缩”的免费午餐。

  3. MIG(Multi-Instance GPU):单 A100 硬件分区成最多 7 个独立实例——SM、L2 cache、显存控制器、显存全部物理切分,每个实例对客户端表现为独立 GPU,QoS 互不干扰。意义:数据中心 GPU 第一次像 CPU 一样能”切片售卖”——云厂商 1/7 A100 实例由此而来,把推理客户的成本门槛降到 1/7。

  4. 第三代 NVLink + HBM2e + Async Copy:NVLink 3.0 12 链 × 50 GB/s = 600 GB/s(V100 NVLink 2.0 是 300 GB/s),DGX A100 8 卡全互联;HBM2e 80GB 版带宽 2039 GB/s(V100 是 900 GB/s);新增 cp.async 指令让显存 → shared memory 拷贝绕过寄存器文件,重叠计算与搬运。意义:显存墙在 LLM 时代被这三件事一起推开

  • 第三代 Tensor Core TF32 / BF16,GPT-3 之后的百亿 / 千亿参数训练精度和吞吐很难同时选下来
  • 结构化稀疏,推理端无法在同硬件再翻倍吞吐
  • MIG,A100 卖不进多租户云市场,单价撑不起 R&D
  • NVLink 3.0 600 GB/s,1024 卡集群训练 GPT-3 通信成瓶颈

案例 1:TF32 让 PyTorch FP32 训练免改代码加速 8×

Section titled “案例 1:TF32 让 PyTorch FP32 训练免改代码加速 8×”
import torch
torch.backends.cuda.matmul.allow_tf32 = True # PyTorch 1.7+ 默认 True
a = torch.randn(8192, 8192, device="cuda", dtype=torch.float32)
b = torch.randn(8192, 8192, device="cuda", dtype=torch.float32)
c = a @ b # 在 A100 上自动走 TF32 Tensor Core, 156 TFLOPS
# V100 / Turing 上仍走 FP32 CUDA Core, 19.5 TFLOPS

意义:老代码、老 dtype、新硬件——这是”硬件兼容性”的教科书示例,也是 A100 上市后立刻被 PyTorch / TF / JAX 默认采纳的关键。

# Megatron-LM / DeepSpeed on A100
model = GPT(...).to(dtype=torch.bfloat16) # BF16 而非 FP16
optimizer = torch.optim.AdamW(model.parameters())
# A100: BF16 312 TFLOPS = FP16 312 TFLOPS, 但动态范围 1e-38 ~ 1e38 (FP16 是 6e-5 ~ 6e4)
# 结果: 大模型训练不再需要 loss scaling 调参, 收敛更稳

意义:BF16 是 LLM 训练事实标准——根因就是 Ampere 让 BF16 与 FP16 算力齐平,且免去 FP16 loss scaling 这个工程坑。

案例 3:MIG 切片把 A100 当 7 张卡卖

Section titled “案例 3:MIG 切片把 A100 当 7 张卡卖”
Terminal window
# 把 A100 切成 7 个 1g.5gb 实例 (每实例 1/7 SM + 1/7 显存)
nvidia-smi mig -cgi 19,19,19,19,19,19,19 -C
nvidia-smi -L
# GPU 0: A100-SXM4-40GB (UUID: ...)
# MIG 1g.5gb Device 0: ... <- 7 个独立 GPU
# MIG 1g.5gb Device 1: ...
# ...

意义:GPU 进入”硬件级多租户”时代——AWS p4d / GCP a2-highgpu / Azure NDA100 的 1/7 A100 实例由此而来,推理客户花 1/7 价格独占 1/7 卡。

V100 SXM2 (Volta): 32GB HBM2 900 GB/s, 125 FP16 Tensor TFLOPS, 0 TF32, 0 BF16, NVLink 300 GB/s
A100 SXM4 80GB: 80GB HBM2e 2039 GB/s, 312 FP16/BF16 Tensor, 156 TF32, 19.5 FP64 Tensor, NVLink 600 GB/s, +MIG +2:4 稀疏
百亿 / 千亿模型训练: V100 能跑但调参和通信压力大; A100 把 BF16 + 80GB + NVLink 变成后续集群默认组合

意义:V100 → A100 的代差不在 FP32 算力(19.5 vs 15.7),而在 Tensor Core 新增四种规格 + 显存带宽 ×2.3 + NVLink ×2——A100 接过 V100 的接力棒,成为 LLM 早期扩张的主力 GPU。

  1. TF32 ≠ FP32:尾数 10 bit(FP32 是 23 bit),梯度累积久了会发散。强收敛任务(科学计算 / 收敛阈值严的优化)需 torch.backends.cuda.matmul.allow_tf32 = False 关闭。

  2. 结构化稀疏不是”训完压缩”:必须训练时就插 2:4 mask(如 apex.contrib.sparsity),fine-tune 后再启用 sparse Tensor Core。直接把训好的稠密权重塞进去精度大跌。

  3. MIG 切完 NVLink 失效:MIG 实例之间不共享 NVLink——切了 1/7 就只能用那 1/7 的显存带宽,需要单卡满算力的训练任务千万别开 MIG。

  4. BF16 vs FP16 选错:FP16 mantissa 10 bit 精度高但范围窄;BF16 mantissa 7 bit 精度低但范围宽。LLM 选 BF16,CV 旧代码选 FP16——硬背规则会翻车。

  5. FP64 Tensor Core 仅 sm_80+:CUDA 11 起 wmma::fragment<..., double> 才合法;旧 CUDA 10 代码不会自动用上,HPC 库(cuBLAS / cuSOLVER)需重新链接 11.0+。

  6. SXM4 vs PCIe 形态差极大:A100 PCIe 版 NVLink 只 64 GB/s(两卡桥接),SXM4 才有 600 GB/s 全互联。多卡训练买 PCIe 版会撞通信墙——DGX A100 / HGX A100 都是 SXM4。

适用

  • 大模型训练 —— OPT / BLOOM / LLaMA-1 等非 TPU 路线常用 A100 集群,GPT-3 属于 V100 时代的前奏
  • 大模型推理 —— BF16 / INT8 + 2:4 稀疏吞吐够支撑 GPT-3 在线服务
  • HPC 双精度 —— FP64 Tensor Core 让 LAMMPS / OpenFOAM / 量子化学包速度 ×2
  • 多租户云 GPU —— MIG 把推理 SaaS 成本切到 1/7
  • TF32 自动加速 —— 老 FP32 训练代码免改获得 ×8

不适用

  • FP8 训练 —— Hopper H100 起才有,A100 仅到 BF16
  • Transformer Engine —— H100 专用,A100 无
  • 消费图形 / 光追 —— GA100 无 RT Core,消费 GA102(RTX 3090)才有第二代 RT Core
  • 极致单卡推理延迟 —— L4 / L40S(Ada)SM 频率更高、能效更优
  • 千亿参数稠密推理 —— 80 GB 不够,H100 80GB / Grace Hopper / B100 才舒服
  • 2020-05 GTC(线上):黄仁勋”厨房演讲”发布 A100,疫情中云端首发
  • 2020-05 DGX A100:8 × A100 SXM4 + 第三代 NVLink Switch,单机 5 PFLOPS BF16
  • 2020-06 MLPerf 0.7:A100 把 ResNet-50 训练时间纪录推到分钟级
  • 2020-06 GPT-3 论文:1024 × V100 集群训练(论文写于 A100 发布前后);后续千亿模型训练全切 A100
  • 2020-09 RTX 3090 / 3080:GA102 消费版上市,第二代 RT Core + 第三代 Tensor Core 下放
  • 2021-06 A100 80GB:HBM2e 升级,2039 GB/s 带宽,配合 GPT-3 时代显存压力
  • 2022-03 Hopper H100 发布:第四代 Tensor Core + Transformer Engine + FP8,A100 让位但兼容性长尾延续到 2024+
  1. 专用化继续深化volta-architecture-2017 加 Tensor Core、turing-architecture-2018 加 RT Core,Ampere 给 Tensor Core 再加四种规格 + 稀疏——专用化不止”加新单元”,也包括”现有单元加规格”
  2. 硬件兼容性的”自动加速”模板:TF32 让 FP32 老代码免改获得 8×——这是黄金硬件升级路径,后续 H100 FP8 / Blackwell FP4 都试图复制(但难度更大)
  3. 多租户切片是 GPU 商业模式转折:MIG 让 A100 进入”按 1/7 卖”的云市场,单卡 ROI 模型彻底改变;后续 H100 沿用、AMD MI300X 跟进
  4. 显存墙是 LLM 时代主矛盾之一:HBM2e ×2.3 带宽、NVLink ×2、cp.async——A100 把”显存喂不饱算力”的窗口推后两年,给 GPT-3 / PaLM 时代留出空间
  5. 架构师为 18-24 个月后铺路:A100 上市时 GPT-3 还没正式发,TF32 / BF16 / 600 GB/s NVLink 是为”还没出现的千亿模型”提前埋的路。架构决策永远超前于应用 18 个月
  • aurora-exascale-2024 —— Aurora 2024 — 不用 NVIDIA 也能造 2 EFLOPS 超算
  • awq-2023 —— AWQ 2023 — 把 70B 大模型权重压到 35GB
  • big-little-2011 —— big.LITTLE — 让一颗芯片同时装快核和省电核
  • biggan-2018 —— BigGAN — 把 GAN 暴力放大到 ImageNet 512×512
  • blackwell-architecture-2024 —— NVIDIA Blackwell — 双 die NV-HBI + 第二代 Transformer Engine + FP4 让万亿参数训练日常化
  • brook-2004 —— Brook for GPUs — 让显卡第一次能用人话编程
  • burgess-2020-turing-rt —— Burgess 2020 RTX ON — Turing 把光线追踪做进硅片
  • cell-be-2005 —— Cell BE — 一颗 CPU 里塞 8 个加速核
  • cuda —— CUDA — 把显卡变成通用并行计算平台
  • cuda-streams-concurrency-2018 —— CUDA Streams 并发量化研究 — 为什么 SM 利用率拉不满
  • cudnn-2014 —— cuDNN — 把卷积写成矩阵乘,让所有深度学习框架共享底层加速
  • cutlass-2020 —— CUTLASS — 把 SOTA GEMM 拆成可组合的 C++ 模板层级
  • dash-numa-1992 —— Stanford DASH — 第一台真跑起来的目录式 CC-NUMA 多处理器
  • deering-1988-triangle-processor —— Deering 1988 Triangle Processor — 现代 GPU 的祖先架构
  • dlrm-2019 —— DLRM — Meta 把工业推荐模型拆成 4 个标准积木
  • fastertransformer-2021 —— FasterTransformer 2021 — NVIDIA 第一代开源 LLM 推理引擎
  • flexsc-2010 —— FlexSC — 把系统调用从同步陷入改成异步队列
  • goldsmith-1987-bvh —— Goldsmith-Salmon 1987 — 让计算机自己给场景搭层次包围盒
  • gpu-cache-coherence-2013 —— GPU 缓存一致性 — 用时戳代替失效消息
  • hopper-architecture-2022 —— NVIDIA Hopper — Transformer Engine + FP8 + TMA + Thread Block Cluster 把硅片为 LLM 量身定制
  • karis-2014-taa —— Karis 2014 TAA — 让游戏每帧只采一次也能 4K 不锯齿
  • karis-2014-ue4-pbr —— Karis UE4 PBR — 把电影质感塞进游戏的 33 毫秒
  • kokkos-2014 —— Kokkos — 一份 C++ 代码同时跑 CPU、GPU、Xeon Phi
  • llm-int8-2022 —— LLM.int8() — 大模型激活值里藏着几个超大异常通道
  • mueller-2022-instant-ngp —— Instant-NGP — 把 NeRF 训练从几小时压到 5 秒
  • newcombe-2011-kinectfusion —— KinectFusion — 用消费级深度相机实时重建三维世界
  • nickolls-dally-2010-cuda-era —— Nickolls-Dally 2010 — GPU 怎么从画三角形变成跑 AI
  • ntk-2018 —— NTK — 把无限宽的神经网络变成一个可解的核方法
  • nvlink-nvswitch-2018 —— NVLink 2.0 + NVSwitch — 把 16 块 GPU 拼成一台机器
  • opencl-2010 —— OpenCL 2010 — 一份代码同时跑 CPU/GPU/DSP/FPGA 的开放标准
  • quantum-supremacy-2019 —— Quantum Supremacy 2019 — 量子机用 200 秒做完超算 1 万年的事
  • ring-allreduce-2017 —— Ring All-Reduce — 把 HPC 的环形规约搬进深度学习
  • smoothquant-2023 —— SmoothQuant 2023 — 把激活的烫手山芋扔给权重
  • sparsegpt-2023 —— SparseGPT — 175B 大模型一次过剪 50%,不重训
  • sycl-cpp-2020 —— SYCL 2020 — 用一份标准 C++ 让 GPU/CPU/加速器一起跑
  • tensorrt-llm-2023 —— TensorRT-LLM — NVIDIA 把 FT 升级成可调度的官方推理栈
  • thrust-2010 —— Thrust — 让 GPU 编程像写 STL 一样一行调用
  • tomasulo-1967 —— Tomasulo 算法 — 让 CPU 自己决定指令的执行顺序
  • unified-memory-2014 —— CUDA Unified Memory — 让 CPU 和 GPU 共享一张内存地图
  • williams-1983-mipmap —— Williams 1983 mipmap — 提前烤好金字塔,纹理过滤变 O(1)