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 之上做了 四件事:
-
第三代 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:4 结构化稀疏:每 4 个权重必须有 2 个为零(训练时 mask),推理时硬件跳过零值,Tensor 算力直接 ×2——TF32 312 / BF16 624 / INT8 1248 TOPS。意义:首次硬件化稀疏,但要求训练阶段就引入 2:4 约束,不是”训完压缩”的免费午餐。
-
MIG(Multi-Instance GPU):单 A100 硬件分区成最多 7 个独立实例——SM、L2 cache、显存控制器、显存全部物理切分,每个实例对客户端表现为独立 GPU,QoS 互不干扰。意义:数据中心 GPU 第一次像 CPU 一样能”切片售卖”——云厂商 1/7 A100 实例由此而来,把推理客户的成本门槛降到 1/7。
-
第三代 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 时代被这三件事一起推开。
这四件事怎么互为支柱
Section titled “这四件事怎么互为支柱”- 没 第三代 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 torchtorch.backends.cuda.matmul.allow_tf32 = True # PyTorch 1.7+ 默认 Truea = 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 默认采纳的关键。
案例 2:BF16 训练 GPT-3 / LLaMA
Section titled “案例 2:BF16 训练 GPT-3 / LLaMA”# Megatron-LM / DeepSpeed on A100model = GPT(...).to(dtype=torch.bfloat16) # BF16 而非 FP16optimizer = 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 张卡卖”# 把 A100 切成 7 个 1g.5gb 实例 (每实例 1/7 SM + 1/7 显存)nvidia-smi mig -cgi 19,19,19,19,19,19,19 -Cnvidia-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 卡。
案例 4:A100 vs V100 同代对比
Section titled “案例 4:A100 vs V100 同代对比”V100 SXM2 (Volta): 32GB HBM2 900 GB/s, 125 FP16 Tensor TFLOPS, 0 TF32, 0 BF16, NVLink 300 GB/sA100 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。
-
TF32 ≠ FP32:尾数 10 bit(FP32 是 23 bit),梯度累积久了会发散。强收敛任务(科学计算 / 收敛阈值严的优化)需
torch.backends.cuda.matmul.allow_tf32 = False关闭。 -
结构化稀疏不是”训完压缩”:必须训练时就插 2:4 mask(如
apex.contrib.sparsity),fine-tune 后再启用 sparse Tensor Core。直接把训好的稠密权重塞进去精度大跌。 -
MIG 切完 NVLink 失效:MIG 实例之间不共享 NVLink——切了 1/7 就只能用那 1/7 的显存带宽,需要单卡满算力的训练任务千万别开 MIG。
-
BF16 vs FP16 选错:FP16 mantissa 10 bit 精度高但范围窄;BF16 mantissa 7 bit 精度低但范围宽。LLM 选 BF16,CV 旧代码选 FP16——硬背规则会翻车。
-
FP64 Tensor Core 仅 sm_80+:CUDA 11 起
wmma::fragment<..., double>才合法;旧 CUDA 10 代码不会自动用上,HPC 库(cuBLAS / cuSOLVER)需重新链接 11.0+。 -
SXM4 vs PCIe 形态差极大:A100 PCIe 版 NVLink 只 64 GB/s(两卡桥接),SXM4 才有 600 GB/s 全互联。多卡训练买 PCIe 版会撞通信墙——DGX A100 / HGX A100 都是 SXM4。
适用 vs 不适用场景
Section titled “适用 vs 不适用场景”适用:
- 大模型训练 —— 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 才舒服
历史小故事(可跳过)
Section titled “历史小故事(可跳过)”- 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+
- 专用化继续深化:volta-architecture-2017 加 Tensor Core、turing-architecture-2018 加 RT Core,Ampere 给 Tensor Core 再加四种规格 + 稀疏——专用化不止”加新单元”,也包括”现有单元加规格”
- 硬件兼容性的”自动加速”模板:TF32 让 FP32 老代码免改获得 8×——这是黄金硬件升级路径,后续 H100 FP8 / Blackwell FP4 都试图复制(但难度更大)
- 多租户切片是 GPU 商业模式转折:MIG 让 A100 进入”按 1/7 卖”的云市场,单卡 ROI 模型彻底改变;后续 H100 沿用、AMD MI300X 跟进
- 显存墙是 LLM 时代主矛盾之一:HBM2e ×2.3 带宽、NVLink ×2、
cp.async——A100 把”显存喂不饱算力”的窗口推后两年,给 GPT-3 / PaLM 时代留出空间 - 架构师为 18-24 个月后铺路:A100 上市时 GPT-3 还没正式发,TF32 / BF16 / 600 GB/s NVLink 是为”还没出现的千亿模型”提前埋的路。架构决策永远超前于应用 18 个月
- 白皮书:NVIDIA A100 Tensor Core GPU Architecture(83 页,2020)
- 第三代 Tensor Core 详解:NVIDIA Ampere Architecture In-Depth(NVIDIA Blog 2020-05)
- TF32 / BF16 PyTorch 文档:Numerical accuracy
- MIG 用户指南:NVIDIA Multi-Instance GPU User Guide
- turing-architecture-2018 —— 直接前代消费线,第二代 Tensor Core + RT Core
- volta-architecture-2017 —— Tensor Core 第一代发源地,A100 直接继承数据中心定位
- pascal-architecture-2016 —— HBM2 + NVLink 第一代鼻祖
- maxwell-architecture-2014 —— SMM 4 分区骨架延续到 Ampere SM
- kepler-architecture-2012 —— SMX 大分区组织
- fermi-architecture-2010 —— ECC + L1/L2 cache 起点
- tesla-architecture-2008 —— SIMT + warp = 32 鼻祖
- tesla-architecture-2008 —— SIMT 鼻祖,Ampere 沿用 warp = 32
- fermi-architecture-2010 —— ECC + L1/L2 cache 在 Ampere 全面继承
- kepler-architecture-2012 —— SMX 4 分区组织延续到 Ampere SM
- maxwell-architecture-2014 —— SMM 4 分区骨架被 Ampere SM 直接继承
- pascal-architecture-2016 —— HBM2 + NVLink 1.0 起点,Ampere 升 HBM2e + NVLink 3.0
- volta-architecture-2017 —— Tensor Core 第一代,Ampere 第三代加 TF32/BF16/FP64 + 稀疏
- turing-architecture-2018 —— 第二代 Tensor Core 加 INT8/INT4,Ampere 在数据中心继续扩展规格
- attention —— Transformer 大模型在 A100 集群上扩张,靠的是 BF16 + 600 GB/s NVLink
- chinchilla —— Chinchilla 实验也是 A100 集群跑出来
- cuda —— Compute Capability 8.0 = Ampere,CUDA 11 起支持 TF32 / BF16 / FP64 Tensor
- 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)