H100 Hopper架构白皮书深度解读：SM、Tensor Core与NVLink训练性能

简介：《NVIDIA H100 GPU 架构白皮书》中文版面向AI训练、高性能计算与数据中心架构研发人员，系统讲解Hopper架构的晶体管规模、SM布局与数据格式支持，帮助读者理解H100相较A100在训练和推理中的性能提升及适用场景。资源为单一PDF文档，体积约8.05MB，适合本地检索、标注与团队传阅；目前已有335人学习下载。白皮书内容覆盖TSMC 4N工艺、800亿晶体管、132个SM、FP8/TF32/FP64精度、第四代Tensor Core、Transformer引擎、DPX指令加速、线程块簇、分布式共享内存、TMA、第二代MIG、机密计算、HBM3、50MB二级缓存，以及第四代NVLink和NVSwitch互联等关键技术，并结合大语言模型训练、基因组学分析（如Smith-Waterman加速）、多租户隔离与AI超算集群等场景展开说明。文档按简介、关键特性摘要、架构深度解析、显存子系统、MIG、安全增强和附录等章节组织，适合按需翻阅。具备CUDA基础的工程与科研人员可将其作为架构选型、性能调优与多卡部署的参考。

1. 这份 H100 中文白皮书，值得花一个晚上读完

做 AI 训练的人迟早会遇到一个问题：卡是拿了，但印象只停留在“显存大、算力高”，一问具体数据——单卡 BF16 算力多少、SM 数多少、L2 多大、NVLink 带宽多少——答不上来。如果你也在搭训练集群、调多卡并行或者做性能卡点分析，这份 H100 GPU 架构白皮书的中文版就是要把这些数字给你补全。它讲的是 Hopper 架构的内部细节，覆盖 SM 调度、Tensor Core、显存层级和多卡互联，读完之后你能直接从规格表反推训练配置，而不是靠玄学调参。适合做大模型训练、高性能计算移植，以及想弄明白 GPU 里“算力”到底有多少水分的工程师。

2. Hopper 架构核心解析：SM、Tensor Core 与线程调度

2.1 从框图到数字：SM 内部到底装了什么

白皮书里最值得啃的是 SM（流式多处理器）的架构图。H100 的 GH100 芯片有 144 个 SM，但消费级产品上实际启用的数量要看型号——比如 SXM 版本是 132 个 SM，PCIe 版本是 114 个。这个差异直接决定了你手算峰值算力时的基数。每个 SM 内有 128 个 FP32 CUDA 核心，意味着一条 warp（32 线程）需要四个周期才能完成一次完整的 FMA 吞吐。

真正让 Hopper 和大伙儿熟悉的 Ampere 拉开差距的，是第四代 Tensor Core。白皮书里提到了两个关键点：一是每个 SM 有 4 个 Tensor Core，支持比矩阵乘法更宽的指令格式；二是新增了针对 Transformer 网络的优化路径——既能在硬件指令层面支持 BF16 累加到 FP32，又能在 FP8 精度下把吞吐继续拉高一倍。翻译成训练场景的语言就是：如果只用 FP32 走卷积，你等于开着卡车走胡同；切到 BF16 或者 FP8，Tensor Core 的吞吐才真正跑起来。

另一个容易忽略的点是 SM 里的寄存器堆。Hopper 架构单个 SM 有 256KB 寄存器文件，整颗芯片合计寄存器和缓存容量是上代产品的 1.5 倍以上。寄存器文件大，意味着同一个 kernel 可以承载更高占用率，同时在显存延迟没有隐藏好的时候，能靠线程切换去“扛住”等待。结合我自己的经验，白皮书里这个数字是值得反复核对的——很多训练脚本出现 occupancy 上不去的问题，追根溯源就是没意识到寄存器堆对 block 数量的限制。

2.2 Warp 调度和占用的容量视角

调度器层面，H100 每个 SM 有 4 个 warp scheduler，每个 scheduler 每周期可以发射一条指令。有人会拿“4 scheduler 等于同时执行 4 个 warp”来说事，实际并不准确——一条指令发射到执行单元，能不能一个周期做完，取决于指令类型。FP32 FMA 是 4 周期发射一波，Tensor Core 的 HMMA 指令则跨更多周期。所以调度器不是越多越好，关键是看白皮书里描述的执行流水线深度。

真正的约束来自“占用率目标”。做 kernel 调优时，大家习惯把 occupancy 调到 50% 以上，目的就是确保有足够多的 warp 来隐藏 TMA（Tensor Memory Accelerator）异步拷贝的延迟和显存访问延迟。Hopper 引入的 TMA 是白皮书里值得单拎出来的部件——它允许数据从全局内存直接异步搬运到共享内存，不需要每个线程单独发 load 指令。想要用好它，得保证 SM 里同时驻留的线程块数量足够。比如一个 kernel 每个 block 用 256 线程、64KB 共享内存，那你最多只能在一个 SM 上放 3 个 block（因为 Hopper 每 SM 共享内存上限是 228KB），再多就溢出，只能靠寄存器换共享内存。

参数对应关系我一般这样记：如果一个 block 的共享内存申请量在 100KB 以上，就别追求 100% 占用率，能用 50% 把 TMA 流水灌满就是胜利。

2.3 手把手核对一张 SM 规格表

拿到白皮书，建议把关键参数抄下来做一张对照表，以后查起来不用反复翻 PDF：

注意：白皮书给出的是架构最大值，实际设备行为要用 CUDA 接口去查，别直接拿表里的 SM 数去写硬编码。

验证方式很简单，用 CUDA Runtime API 就能读到真实数字。下面的代码可以丢到 Linux 下编译运行，它会打印设备名、SM 数量、时钟频率、显存带宽上限等一揽子信息：

#include <cuda_runtime.h> #include <cstdio> int main() { int dev = 0; cudaDeviceProp prop; cudaGetDeviceProperties(&prop, dev); printf("Device: %s\n", prop.name); printf("SMs: %d\n", prop.multiProcessorCount); printf("Max threads per SM: %d\n", prop.maxThreadsPerMultiProcessor); printf("Clock rate: %.2f GHz\n", prop.clockRate * 1e-6); cudaDeviceReset(); return 0; }

编译命令用nvcc -arch=sm_90 device_query.cu -o device_query，跑出来的 SM 数量基本就是准确上限。注意maxThreadsPerMultiProcessor是硬件层面的绝对上限，但实际还要看寄存器分配——如果你在 kernel 里用__launch_bounds__限制每线程寄存器数（比如 128 个），那么每 SM 能放的线程数就不是白皮书理论值，而是65536 / 128 = 512线程。

2.4 白皮书里容易被误解的“计算单元复用”

一个非常经典的坑：H100 的 FP32 核心是 128 个，但 TF32 的 Tensor Core 峰值算力只有 BF16 的一半。原因不在 Tensor Core 数量，而在数据路径——TF32 输入是 19 位格式，需要先从 FP32 转一次格式，转换单元在硬件上是有限的。白皮书里的图会标注成“same tensor core, different precision”，很多人会误以为不同精度只是指令参数不同，实际上延迟和吞吐都变了。

如果你的模型用 FP32 训练，Tensor Core 根本不会参与——FP32 矩阵乘法走的是传统 CUDA 核心 FMA 路径。这时候白皮书里的 989 TFLOPS（BF16）跟你没有任何关系，你的算力天花板是 67 TFLOPS 的 FP32 峰值。这一点在评估训练框架时经常被人忽略，我见过有团队拿着 FP32 训练的 benchmark 说“H100 不如上一代”，其实是没开 AMP 混合精度。

3. 显存与数据搬运实战：HBM3、NVLink 和 PCIe 的带宽边界

3.1 HBM3 和“看起来很高”的 3.35TB/s

白皮书里的 3.35TB/s 来自 HBM3 显存系统，由 8 颗 HBM3 堆叠组成，总位宽 5120 位。但这只是理论峰值——生效带宽取决于访问模式，而 HBM 对连续块访问、stride 访问、随机小粒度访问的效率差异非常大。

实测跑法可以参考 CUDA 自带的 bandwidthTest 工具，或者自己写一个 memcpy D2D 基准。HBM3 的连续大块拷贝能达到大约 2.9~3.1TB/s，含系统调用开销。但是当内核用随机访问模式，比如按 4KB 粒度读写，有效带宽很可能掉到 300GB/s 以下。原因很简单：HBM3 的 burst 长度是固定的，短请求会导致 bus turn-around 频繁切换。

所以在设计数据流水时，我一般会要求所有显存内的张量搬运尽量走cudaMemcpyAsync的连续段，不要用cudaMemcpy2D做大量窄条拷贝。白皮书并没有直接告诉你这个限制，但通过理解它的 DRAM 组织结构可以推出来：数据是按 channel 交错的，32B 对齐访问只命中一个 channel，效率天然打折。

3.2 NVLink 4 和环形拓扑下的 AllReduce 代价

NVLink 4 的单条链路带宽是 50GB/s 双向，SXM 版卡上有 18 条链路，合计 900GB/s。这组数字在官方宣传里很好看，但做分布式训练的人得把它换算成 AllReduce 时间。

一个常见的做法是用 ring-allreduce：卡数越多，通信数据被切成 N 份，每张卡只和邻居交换。H100 的 NVLink 拓扑里，8 张卡在 fully connected 拓扑下任两卡之间都是 2-hop 以内，但这不等于延时可忽略。NUMA 绑定、PCIe switch 路径、NVLink 域划分都会影响实际吞吐。

跑 NVLink 实测，可以用 NCCL 自带的all_reduce_perf工具。下面的命令对 8 卡做 1GB 数据的 AllReduce，看 bus bandwidth：

mpirun -np 8 ./build/all_reduce_perf -b 1G -e 1G -f 2 -g 1 -t 1

参数解释：-b和-e是起始和结束字节数，这里固定 1GB；-f 2表示从 2 的幂次递增（此处仅一次）；-g 1是每个进程用一张 GPU；-t 1表示测试 AllReduce 模式为 ring。如果输出里的 bus bandwidth 明显低于 500GB/s，先怀疑 NVLink 是否跑在降速状态——直连链路的 link rate 可能被驱动配置限制成 25GB/s 而不是 50GB/s。

3.3 PCIe Gen5 和 NVLink 的边界转移

对多卡训练来说，重点不是算力而是数据能不能喂得饱。H100 SXM 版本的 NVLink 总带宽 900GB/s，而 PCIe Gen5 x16 的单向带宽是 63GB/s。如果你的数据预处理、embedding 查询、checkpoint 加载都走 PCIe，那瓶颈会非常明显——一次加载 10GB 模型权重的时间，用 NVLink 是 11ms，用 PCIe 是 160ms。差一个数量级。

白皮书里也提到 NVLink 支持多卡共享显存寻址（NVLink C2C 模式），但在软件层面，CUDA 的cudaMallocManaged并不能自动把数据智能分配到各卡，仍然需要显式控制 prefetch。所以这块我的经验是：训练循环开始前做一次 warmup 数据搬运，把常用权重打到各卡显存，训练中只同步梯度。

实测脚本可以用nvidia-smi dmon -s p -d 5盯总线利用率，如果发现 PCIe read 持续跑到 90% 以上，说明你的数据流水线在拖后腿，优先把数据增强、tokenize 挪到 CPU 侧异步，或者直接换用 NVLink 桥接的存储方案。

提示：白皮书里的架构图只体现硬件拓扑，实际驱动还会在 Device Memory 和 Host Memory 之间做双份拷贝。所以所有带宽数字都要默认乘以一个 0.7~0.8 的衰减系数。

4. 从白皮书数字到性能模型：Roofline 与动手算训练配比

4.1 把白皮书规格换算成有效 FLOPs

我们做性能评估，第一步就是把 H100 白皮书里的表转成自己的“计算上限表”。常见帧：

• FP32 稠密算力：约 67 TFLOPS
• TF32 Tensor Core：约 495 TFLOPS（稠密）
• BF16 Tensor Core：约 990 TFLOPS（稠密）
• FP8 Tensor Core：约为 BF16 的两倍

这些数字的依据是 SM 数 × 每 SM 操作数 × 主频。以 BF16 为例：132 SM × 4 Tensor Core × 每个 Tensor Core 每周期 512 FLOPs × 1.98GHz 约等于 990 TFLOPS。这就是白皮书尾巴上那组数字的来源。实际上手写 kernel 时，FFMA（浮点乘加）在每个周期能发射的指令数有限制，所以 CUDA CORE 能跑出的效率远低于 Tensor Core。别拿 CUDA 核心去算大矩阵乘法，那是给 SIMD 型负载准备的。

4.2 用 Roofline 判断模型是算力密集还是访存密集

Roofline 模型把时间分成两个区间：计算密集区和访存密集区。H100 的 BF16 Tensor Core 峰值算力约 990 TFLOPS，显存带宽 3.35TB/s，那么转折点的算术强度就是 990 / 3.35 ≈ 295 FLOP/Byte。如果模型的算术强度低于 295，实际时间会被显存延迟卡死，你怎么调 Kernel 都白搭。

拿 LLM 训练举例：一个 batch size 为 8、sequence length 为 2048、hidden size 为 8192 的 attention 计算，每字节的 FLOPs 大概是 200 出头——明显低于转折点，属于访存密集型。这时候应该把注意力放在 batch size 是否足够大以及是否用了 FlashAttention 来减少 HBM 访问，而不是去优化矩阵乘法的指令发射。白皮书的价值就在于让你提前算清楚这一点。

4.3 动手算一个 7B 模型的理想训练时间

我一般会写个小脚本来做粗算，思路很简单：定义模型规模、精度、FLOPs 利用率和卡数，输出预估训练时间。这样在做实验规划时，几分钟就能判断一个方案是否靠谱。

def train_time_hours(model_params, tokens, peak_tflops, mfu, gpu_count): total_flops = 6 * model_params * tokens effective_tflops = peak_tflops * mfu * gpu_count seconds = total_flops / (effective_tflops * 1e12) return seconds / 3600 # 7B 模型，3000 亿 token，H100 BF16 全卡算力 990TF，MFU 40%，256 卡 print(train_time_hours(7e9, 3e10, 990, 0.40, 256))

大致算下来，25 天左右。这里 MFU（Model FLOPs Utilization）取 0.4 是基于 Mamba 等非 transformer 结构一般到不了 0.5；如果你用的是优化到位的 transformer 加 FlashAttention，MFU 可以到 0.45~0.5。脚本里的关键参数是mfu，它把 float 峰值做了折扣——建议别拍脑袋定，而是先用 ncu profiling 一段真实训练来标定。

4.4 算力效能的另一个边界：显存容量与精度混合

白皮书里没有直接回答“该不该用 FP8”这个问题，但架构上给出了答案：Hopper 的 FP8 Tensor Core 吞吐是 BF16 的两倍，代价是尾数位太少。结果就是 FP8 在 loss 缩放和 emulated FP32 累积上要花额外算力，很可能抵消收益。我的建议：先把主训练精度设为 BF16，开启混合精度；只有在调通梯度裁剪和 loss 缩放后，再考虑小规模试验 FP8。

另外，显存容量决定了你能开多大的 batch 和序列长度。7B 模型 BF16 权重占 14GB，Adam 优化器状态接近 42GB，激活值则依赖 batch。用 80GB 单卡做推理可以，但做全量微调，单卡一定爆显存。这时候白皮书的共享显存（NVLink 跨卡访问）能缓解问题，但不建议依赖它——跨卡访问的带宽与本地显存带宽差一个量级，硬塞进去只会把 time-to-train 拉长。

5. 坑与排查：读白皮书时最容易踩中的五个盲区

5.1 现象：照白皮书 SM 数写核函数，运行时错乱

有人照着白皮书里“144 SM”或者“132 SM”直接把 kernel 的 grid 大小写死，结果在 H100 上运行后性能不升反降，甚至 launch failure。原因：不同型号的 SXM 与 PCIe 使能的 SM 数不同，白皮书正文给的是 GH100 芯片全量，而不是某个具体 SKU 的激活数量。解决方法：运行device_query读取multiProcessorCount，所有依赖 SM 上限的 grid 配置都要从设备属性里取，不写硬编码。

5.2 现象：BF16 实测吞吐始终到不了峰值的一半

一股脑用“990 TFLOPS”去估计大 GEMM 的耗时，结果 kernel 跑了三倍时间。原因：Tensor Core 的峰值是在全 L2 命中、无 bank conflict、指令完全流水的情况下测的，真实矩阵乘法里还得看数据布局（row-major vs column-major）、split-K 的归约开销、以及是否走 cuBLAS 的 heuristics。解决：直接调cublasGemmEx并分别测CUBLAS_COMPUTE_32F_FAST_TF32和CUBLAS_COMPUTE_16F两种路径，选更快的精度组合；不要自己写 sgemm 去挑战库的调度。

5.3 现象：连锁了 NVLink，但 AllReduce 带宽跑不满

配置了 8 卡 NVLink，跑 NCCL 全归约却只有 300GB/s 的 bus bandwidth。原因：NVSwitch 在其中参与了拓扑路由，如果使用环境是 4 卡直连 NVLink、4 卡走 PCIe 的混合拓扑，NCCL 会退化到 tree algorithm；另外 GPU 频率和 NVLink 频率没有拉满，也可能导致降速。解决：先跑nvidia-smi topo -m看拓扑图，再设置NCCL_ALGO=Ring或NCCL_PROTO=Simple做对比；必要时用NVLink 带宽测试工具（nvbandwidth）分开测点到点。

5.4 现象：Shared Memory 一开大，占用率直接砍半，kernel 变慢

Hopper 的 TMA 功能很诱人，一上来就申请 160KB shared memory 作为 staging buffer，结果发现 SM 上的 block 数只剩 1，占用率变低。原因：共享内存是每个 SM 的有限资源，你开 160KB 后，按 228KB 上限算只能容纳 1 个 block，导致延迟隐藏能力骤降。解决：把大块 staging 改成按 tile 分块搬运，或用cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, size)提升动态共享内存额度，配合__launch_bounds__限制线程数为 1024，留出 double-buffer 余量。

5.5 现象：白皮书推荐的 griddim 在 V100/A100 上不可用

把 Hopper 特有的 grid sync / cluster 特性搬迁到老卡上，报错 invalid device function。原因：cudaGridDependencySynchronize()是 SM90 引入的，老架构没有对应硬件支撑，白皮书不会标注这类向下兼容问题。解决：查阅 CUDA 编程指南的“Compute Capability”对比表，对所有架构特有 API 做宏分支：

#if __CUDA_ARCH__ >= 900 // 启用 grid sync 或 cluster launch #endif

这样代码才能在 A100/H100 等异构集群里安全运行。

6. 进阶用法：把白皮书规格表反推成训练配置验证单

白皮书的终极用途，是把“宣传数字”翻译成“当前训练的期望值”并反向查问题。我的做法是每次跑大任务前，先做一轮“规格核对 + 实测基线”的验证，耗时不超过半小时，但能过滤掉大量环境配置错误。

第一步，进入 Python 直接校验当前 GPU 的白皮书宣称值：

import pynvml pynvml.nvmlInit() handle = pynvml.nvmlDeviceGetHandleByIndex(0) mem = pynvml.nvmlDeviceGetMemoryInfo(handle) clocks = pynvml.nvmlDeviceGetClockInfo(handle, pynvml.NVML_CLOCK_SM) print(f"SM clock: {clocks} MHz, Total mem: {mem.total // (1<<20)} MB")

pynvml会返回真实时钟与显存总量，跟白皮书标称做比对。如果 SM 时钟比标称低 20%，确认是否处于省电模式或温度墙。

第二步，用一个标准 GEMM 基准测试建立“本机”实际算力：

./gemm_bench --m=8192 --n=8192 --k=8192 --precision=b float16 --alg=1

拿到结果后与白皮书峰值对比，得到这台机器的算力系数（我一般记做 0.75~0.85）。之后每次跑训练，实际 MFU 超过算力系数一半就算健康。

第三步，针对热点 kernel 用 Nsight Compute 抓瓶颈：

ncu --set full --kernel-name regex:matmul_kernel ./train_script

看 frame 数据里的 Memory Throughput 和 Compute Throughput 两个数值。如果 Compute 已到 90% 以上，表示算力吃满，应该去优化通信流水；如果 Memory Throughput 高但 Compute 只有 30%，说明访存模式有毛病，优先处理 layout、shared memory 和 TMA 路径。

这三步跑下来，基本就能定位性能问题到底在哪一层。之后再去微调 batch size、gradient checkpoint 和 tensor core 路径才有的放矢，而不是猜。在那以后我每次接新项目都会强制自己先走一遍这份“白皮书核对 + 实测基线”的流程，因为被规格表骗过一次，就不想再被骗第二次了。希望帮到你。

本文还有配套的精品资源，点击获取