NVIDIA · HOPPER · GH100 DIE · SXM
H200· 数字速查
面向 CUDA 算子 / 框架工程师 · 写 kernel 时一眼查到的关键数字 · 括号内为相对 ×A100 倍数
整机级
CHIP-LEVEL · FULL-DEVICE SPEC| 指标 | H200 SXM | 相对 A100 · 刻度 0–4× | 直觉注解 |
|---|---|---|---|
| 核心规模 | |||
| 发布年份 | 2023.11（SC23） | A100 · 2020.5 | 发布相隔 3.5 年；H200 2024Q2 起交付 |
| 架构 / die | Hopper · GH100 | Ampere · GA100 | 同 die，H200=H100 换 HBM3e |
| 工艺 / 晶体管 / 面积 | TSMC 4N · 80B · 814mm² | 7nm · 54.2B · 826mm² | 4N 专工艺，频率/能效双升 |
| SM 数量 | 132 | ×1.22 | 108→132，多 24% |
| GPC / TPC | 8 GPC / 66 TPC | 7 GPC / 54 TPC | cluster 受限于单 GPC 内 |
| FP32 CUDA Cores | 16,896 (128/SM) | ×2.44 | 每 SM 翻倍 64→128 |
| FP64 Cores | 8,448 (64/SM) | ×2.44 | clock-for-clock 2× |
| INT32 Cores | 8,448 (64/SM) | ×1.22 | 未翻倍，地址算力瓶颈点 |
| Tensor Cores（第四代） | 528 (4/SM) | ×1.22 | 数量同，但单核 2× 吞吐 |
| 存储与带宽 | |||
| HBM 容量 | 141 GB HBM3e | ×1.76 | 80GB→141GB，放得下大模型 |
| HBM 带宽 | 4.8 TB/s (4800 GB/s) | ×2.36 | 2039→4800，HBM 是最大跃升 |
| L2 Cache | 50 MB | ×1.25 | 40→50MB，缓模型切片 |
| L2 分区 | partitioned crossbar | 同 | GPC 直连分区，局部性好 |
| 频率 / 功耗 / 互联 | |||
| GPU Boost Clock | ~1.83 GHz | ×1.30 | 1.41→1.83GHz；1 cyc ≈ 0.55ns |
| TDP | 700 W (max 805W) | ×1.75 | 液冷 SXM，能效仍升 |
| NVLink（第四代） | 900 GB/s · 18 links | ×1.50 | 600→900，all-reduce 3× |
| PCIe | Gen5 · 128 GB/s | ×2.00 | Gen4 64，双向往返 |
| MIG 实例 | 最多 7 × 18GB | 同 | 2代 MIG，含独立 NVDEC |
| Compute Capability | 9.0 | 8.0 | cluster/DSMEM/wgmma/TMA 门槛 |
单 SM 级
PER-SM RESOURCES · CC 9.0 VS CC 8.0| 资源 | H200 (cc 9.0) | A100 (cc 8.0) | 直觉注解 |
|---|---|---|---|
| FP32 Cores | 128 | 64 | datapath 翻倍，FP32 密集 kernel 受益 |
| FP64 Cores | 64 | 32 | 2× clock-for-clock |
| Tensor Cores | 4（第四代） | 4（第三代） | 同数，单核 MMA 吞吐 2× |
| Register File | 256 KB = 65,536 × 32-bit | 256 KB | 同；occupancy 与 reg/thread 权衡核心 |
| Shared Memory（max） | 228 KB | 164 KB | +39%，tiling 能开更大 |
| L1 + SMEM 合池 | 256 KB | 192 KB | 可配比，228KB smem 时 L1 仅 28KB |
| Max threads / SM | 2,048 | 2,048 | 同；64 warps |
| Max thread blocks / SM | 32 | 32 | 同 |
| Max registers / thread | 255 | 255 | 255 而非 256（硬件保留） |
| 每 thread 平均寄存器 @满占用 | 32 | 32 | 65536 / 2048；加 reg 必降占用度 |
| Max thread blocks / cluster | 16 | — | Hopper 新层；实际 cluster ≤ 8 blocks |
OCCUPANCY 速算 · 寄存器与占用度的取舍
每 SM 65,536 个寄存器、最多 2,048 线程。寄存器越多 → 单线程越强，但能并发的 block/warp 越少（占用度下降）。
# 每 thread 用 R 个寄存器时，单 SM 最大并发线程数：
threads = min(2048, floor(65536 / R) * blocksize)
# 例：R=128, blocksize=256 → floor(65536/128)=512 threads/block... 实际按 block 对齐
# R=64 → 每 SM 1024 threads (50% 占用) | R=32 → 2048 (100%)Hopper 上 wgmma 需要较多寄存器存放 accum（每个 warpgroup 4 warps），常主动降到中占用度换 ILP——这是 Hopper kernel 与 Ampere 最大的调参差异之一。
Hopper 架构图
MICROARCHITECTURE SCHEMATIC · GH100FIG. 03 — GH100 BLOCK DIAGRAM · HBM → L2 → GPC → CLUSTER → SM
存储层级
MEMORY HIERARCHY · CAPACITY × BANDWIDTH| 层级 | 容量（每 SM / 全片） | 峰值带宽 | 谁管理 / 注解 |
|---|---|---|---|
| Register File | 256 KB/SM · 33.8 MB/片 | ~13 TB/s/SM（量级） | 编译器分配；最快，决定占用度 |
| Shared Memory / L1 | 228 KB/SM（合池 256KB） | ~4 TB/s/SM（量级） | 程序员显式；bank conflict 防护 |
| L2 Cache | 50 MB（全片共享） | ~12 TB/s（量级） | 硬件管理；可设 residency 控制 |
| HBM（global） | 141 GB | 4.8 TB/s | 最大最慢；用 TMA / cp.async 隐藏 |
| ★ DSMEM（集群内邻居 SM） | 邻居 SM 的 smem | ~local smem 量级 | Hopper 新；SM-to-SM 网络 |
带宽列为架构推导量级 — 见页脚脚注 [1]
精度算力对照
PRECISION × THROUGHPUT · TFLOPS单位 TFLOPS（INT8 为 TOPS）。dense = 稠密峰值（写 kernel 算 roofline 用这个），sparse = 2:4 结构化稀疏（NVIDIA 官方 spec 常引此值）。
DENSE · 主读数（基准口径）SPARSE · 副读数（2:4 参考，非基准）
| 精度 | H200 DENSE | H200 SPARSE | A100 DENSE | A100 SPARSE | ×A100 (dense) |
|---|---|---|---|---|---|
| FP64（非 TC） | 34 | — | 9.7 | — | ×3.50 |
| FP64 Tensor Core | 67 | — | 19.5 | — | ×3.40 |
| FP32（非 TC） | 67 | — | 19.5 | — | ×3.40 |
| TF32 Tensor Core | 494 | 989 | 156 | 312 | ×3.17 |
| BF16 Tensor Core | 990 | 1,979 | 312 | 624 | ×3.17 |
| FP16 Tensor Core | 990 | 1,979 | 312 | 624 | ×3.17 |
| ★ FP8 Tensor Core | 1,979 | 3,958 | — | — | HOPPER 新 |
| INT8 Tensor Core | 1,979 | 3,958 | 624 | 1,248 | ×3.17 |
NVIDIA datasheet 的 TF32/BF16/FP16/FP8/INT8 峰值通常是 sparse（2:4） 值——它要求每 4 个元素里至多 2 个非零，由 Tensor Core 硬件压缩。绝大多数 GEMM/Attention kernel 的权重/激活不满足 2:4 结构，所以 roofline / 性能分析必须用 dense 列（sparse 的一半），否则会把目标定高 2 倍。
关键延迟
LATENCY CROSS-SECTION · CYCLES @ ~1.83 GHZ访问各级存储的延迟量级。cycle 数 @ ~1.83 GHz 换算 ns（1 cyc ≈ 0.546 ns）。资深工程师最易记忆模糊、却对优化最关键的一组数。
| 访问类型 | ~cycles | ~ns @1.83GHz | 直觉注解 |
|---|---|---|---|
| Register 读 | ~1 | ~0.5 | 流水线内零可见延迟；依赖链 FP32~4cyc 才是要 hide 的 |
| Shared Memory（无 bank conflict） | ~28–33 | ~15–18 | SM 内 32 bank×4B 软件管理 |
| L1（global 命中） | ~32–33 | ~18 | 与 smem 共池，命中即此延迟 |
| L2（near 分区，命中） | ~264 | ~144 | GPC 本地分区；H100/H200 的 L2 是两级 |
| L2（far 跨分区） | ~500+ | ~275+ | 数据落远分区 ≈ near 的 2× |
| HBM（global，L2 miss） | ~650 | ~355 | 最慢；TMA / cp.async 必须异步覆盖 |
| ★ DSMEM（集群内邻居 SM） | ~200–210 | ~110–115 | 本地 smem(~33) 的 ~6×，但远低于 global(~650) |
| SMEM bank conflict | ×2 / ×4+ | 倍增 | 32 路冲突最坏 ×32 |
实测来源与可信度说明 — 见页脚脚注 [2]
一次 HBM 访问（~650 cyc）≈ 22 次 shared memory 访问 ≈ 650 次 寄存器访问。这就是为什么 tiling 到 shared memory + 用 TMA 异步预取 是 Hopper kernel 的命门——不是可选项。
Roofline 临界点
ROOFLINE RIDGE POINT · FLOP/BYTEridge point 是硬件常数（每块卡固定）；AI（arithmetic intensity）是算子变量（每字节 FLOP）。判断：算子 AI > ridge → compute-bound；AI < ridge → memory-bound。用 dense 峰值算（见 §5）。
| 精度 (dense) | 峰值 | 临界 ridge (FLOP/byte) | 典型 kernel 归属 |
|---|---|---|---|
| FP8 TC | 1979 T | ~412 | 几乎只有大 GEMM 能达标 |
| FP16 / BF16 TC | 990 T | ~206 | GEMM / FlashAttention 可达 |
| TF32 TC | 494 T | ~103 | 大 GEMM 达标 |
| FP32（非 TC） | 67 T | ~14 | elementwise/归一化多在此 |
| FP64 | 34 T | ~7 | HPC kernel 易达 compute-bound |
实操：用 Nsight Compute 的 achieved AI 对照 ridge point，立即知道瓶颈是算力还是带宽。H200 带宽翻倍（×2.36）后，很多 kernel 从 compute-bound 滑向 memory-bound——这是 H100→H200 移植时最该重检的。
Hopper 关键特性
FEATURE MODULES × 5 · OPERATOR OPTIMIZATION五大算子优化特性，每张卡 = 定义 + 关键数字 + 极简 snippet。点 [详情→] 看机制 / 完整可编译示例 / 陷阱。
TMA · Tensor Memory Accelerator
单线程发起 · 1D-5D 张量 · ≤ smem 容量整块硬件地址生成引擎，用 copy descriptor（张量维度+块坐标）替代逐元素寻址。把 A100 的 LDGSTS 地址循环开销整个卸到硬件。
// host: 建 tensor map（描述符）
CUtensorMap tmap;
cuTensorMapEncodeTiled(&tmap, ...);
// device: 单线程发起异步整块搬运
asm("cp.async.bulk.tensor.2d.shared.global"
" .mbarrier::complete_tx::bytes ...");
mbarrier.wait(); // 等搬运完wgmma · 异步 Tensor Core
1 warpgroup = 4 warps = 128 threads · m64nNk16warp group 级异步 MMA，比 Ampere mma.sync tile 更大且异步执行（fence→commit→wait）。A100 是单 warp 同步 mma.sync，Hopper 改为 4 warp 协作异步，吞吐翻倍。
// 128 线程协作发起一次异步 GEMM
wgmma.fence;
asm("wgmma.mma_async.sync.aligned"
".m64n128k16.f32.f16.f16 ...");
wgmma.commit_group;
wgmma.wait_group 0;Thread Block Clusters
≤ 8 thread blocks · 必须同 GPCCUDA 编程层级新加一层（thread→warp→block→cluster→grid）。cluster 内多个 block 保证并发调度到一组 SM，硬件 barrier 同步，是 DSMEM 的前提。
// launch 时声明 cluster 维度
__global__ void __cluster_dims__(2,2,1)
kernel(float* d) {
namespace cg = cooperative_groups;
cg::cluster_group c = cg::this_cluster();
c.sync(); // 跨 SM 硬件 barrier
unsigned r = c.block_rank();
}DSMEM · 分布式共享内存
集群内 SM 间直访 · 比走 global 快 ~7×cluster 内各 block 的 shared memory 映射进统一地址空间，线程可直接 load/store/atomic 邻居 SM 的 smem，不必绕道 global memory 做数据交换。
cg::cluster_group c = cg::this_cluster();
extern __shared__ int smem[];
// 取邻居 block(rank 1) 的同位置 smem 指针
int* nb = c.map_shared_rank(smem, 1);
int v = *nb; // 直达 SM-to-SM 网络FP8 · 原生低精度
E4M3(精度) + E5M2(范围) · 2× FP16 吞吐两种格式：E4M3 精度高范围小（前向），E5M2 范围大精度低（反向梯度）。相比 FP16 存储减半、吞吐翻倍。eager mode 支持运行时与 FP16/BF16 互转。
#include <cuda_fp8.h>
__nv_fp8_e4m3 w[...]; // 前向权重
__nv_fp8_e5m2 grad[...]; // 反向梯度(大范围)
float acc[...]; // FP32 累加
// 经 wgmma mma_async .e4m3 发射