H100 架构介绍（Introduction to H100）

目录 · ← l0 · l2 →

H100 架构介绍（Introduction to H100）

本文档基于课程讲义《1. Introduction to H100.pdf》（Lesson 1）整理，系统介绍 NVIDIA H100 GPU 的 体系结构，作为学习后续 CUDA 专题（TMA、异步执行、Tensor Core、WGMMA、流并发等）的架构基础。 建议先完成 02-CUDA-并行编程与-GPU-体系结构.md 中的同步模型与内存层级基础，再阅读本文。

目录

1. 概述与规格
2. 核心变革：全面异步化
3. TMA（Tensor Memory Accelerator）
4. 同步拷贝 vs 异步拷贝
5. 第四代 Tensor Core 与 WGMMA
6. 其他关键特性
7. GPU 整体组织结构
8. GigaThread Engine
9. 存储层级：HBM3 → L2 → Shared → Registers
10. SM 详解
11. Warp 与指令调度
12. 线程块在 SM/SMSP 间的分布
13. PCIe 5.0 主机接口
14. 总结与学习衔接

1. 概述与规格

H100 是 NVIDIA 基于 Hopper 架构 的旗舰 GPU，专为深度学习与高性能计算设计。

核心定位：H100 相比上一代最大的改进是引入了异步（Asynchronous）特性——这是贯穿整个 Hopper 架构学习的主题（TMA、异步拷贝、异步屏障、WGMMA、流并发等都在此基础之上）。

2. 核心变革：全面异步化

传统 GPU 编程（同步模型）里，数据搬运和计算是串行的：

同步模型：拷贝 → 计算 → 拷贝 → 计算 → ...（步骤串行，延迟无法隐藏）

H100 的核心变革是把数据搬运交给专用硬件在后台完成，让 SM 专注计算：

异步模型：SM 专注计算，TMA 在后台搬运数据，二者重叠执行

这就是 02 文档里”同步执行模型 vs 异步执行模型”一节的硬件基础——Hopper 把”流水线重叠”这一 软件优化思想，做成了架构级的一等公民。

3. TMA（Tensor Memory Accelerator）

TMA 是 Hopper 引入的、每个 SM 都配备的专用数据搬运单元，用于把”张量拷贝”从 SM 上卸载出去。

3.1 为什么需要 TMA

• 目的：加速 memory-bound（访存受限） 操作。
• 过去：global memory ↔ shared memory 的搬运要靠每个线程做地址计算、循环遍历、发多条指令。
• 现在：TMA 通过描述符（descriptor） 异步地发起搬运，单个线程就能发起整个数据搬移， 硬件在后台负责 stride/offset/bounds 与数据搬移。

3.2 对比

TMA 与 cp.async.bulk（后续专题）配合，是 Hopper 上高性能 GEMM/注意力 kernel 的核心搬运动力。

4. 同步拷贝 vs 异步拷贝

这是理解 Hopper 编程模型的最关键概念之一：

• 同步拷贝（Synchronous copy）：发起拷贝的线程（或 warp）必须等拷贝完成后才能继续。 心智模型简单——控制权回到线程时，数据一定已就位。
• 异步拷贝（Asynchronous copy）：发起只启动传输，发起者可以立刻做别的事， 等到真正要用数据前再 wait/check 完成。

4.1 异步拷贝的价值

能够重叠执行 → 用计算隐藏内存延迟

典型场景是流水线 kernel（GEMM/attention）：

计算 tile i 的同时，预取 tile i+1

4.2 代价

异步的唯一代价是簿记（bookkeeping）复杂度上升——你要自己跟踪”哪些数据已到位、何时能安全使用”。 这正是后续专题（异步屏障 cuda::barrier、cuda::pipeline、TMA）要解决的问题。

5. 第四代 Tensor Core 与 WGMMA

每个 SM 有 4 个 Tensor Core（全 GPU 共 528 个）。第四代 Tensor Core 的关键特性：

5.1 WGMMA（Warp Group Matrix Multiply Accumulate）

• 4 个 warp（= 128 线程）协作完成一个 MMA（矩阵乘加）。
• 与传统 wmma（单 warp）相比，WGMMA 通过 warp 组获得更大的矩阵 tile 和更高的吞吐。

5.2 稀疏 Tensor Core

• 支持结构化稀疏（structured sparsity，数据以预定义布局置零）的快速 MMA。

5.3 FP8 支持

• 提供针对 FP8 的专用指令，可达 1,979 TFLOPS（相对上一代巨大提升）。
• 吞吐：FP16 密集为 1024 FLOPS/cycle，FP8 为 2048 FLOPS/cycle。

5.4 WGMMA 寄存器布局（重要约束）

FP8/FP16 时，WGMMA 期望数据打包进寄存器（例如 2 个 FP16 装进 1 个 32 位寄存器）。 这些约束是后续学习 wgmma 专题时必须记住的硬性要求。

6. 其他关键特性

• Transformer Engine：面向 Transformer 的混合精度（FP8/FP16）加速。
• 4th-Gen NVLink：GPU 间高速互连（详见 §7.5）。
• DPX 指令：动态规划加速（如 Smith-Waterman 序列比对）。
• 50 MB L2 Cache：大容量统一二级缓存（详见 §9.3）。

7. GPU 整体组织结构

整个 GPU die 划分为：

• NVIDIA GigaThread Engine（线程块分发）
• 8 个 GPC（Graphics Processing Cluster）
• HBM3 堆栈 + 内存控制器
• PCIe 5.0 主机接口
• NVLink 交换机 / 端口 / Hub
• L2 缓存切片

8. GigaThread Engine

GigaThread Engine 是把一次 kernel 启动分发成线程块（CTA）的硬件：

• 跟踪哪些 CTA 未启动 / 运行中 / 已完成。
• 当某个 SM 有容量容纳下一个 CTA 时，GigaThread Engine（及相关前端逻辑）就把下一个 CTA 分给它。
• 强制占用率（occupancy）限制；在 Hopper 上还理解 cluster（线程块簇）（后续专题）。

简单理解：它是”线程块调度器”，决定哪个 block 去哪个 SM。

9. 存储层级：HBM3 → L2 → Shared → Registers

H100 的数据流遵循经典的越近越快、容量越小层级：

HBM3（大容量、高带宽）→ L2 Cache（统一 50MB）→ Shared Memory（低延迟）→ Registers（最快）

9.1 HBM3（片外显存）

• 80 GB，3.35 TB/s，5 个堆栈（实际 6 个，1 个因良率关闭）。
• 5120 位总线宽度 → 使 TMA 能一次搬 128 字节。
• 通过 10 个独立的 512-bit 内存控制器连接。
• 数据路径：SM → (L1D/coalescer) → L2 → memory partition/crossbar → memory controller → HBM 堆栈。
• 内存控制器负责 DRAM 协议与调度（activate/precharge/读写时序、重排以最大化行命中、地址 → channel/bank/row/col 映射）。

9.2 GPC 与 TPC

• GPC = 一组 SM 的分组（讲义记为每 GPC 18 个 SM；H100 共 8 个 GPC、132 个启用 SM）。
• 每个 GPC 连接自己专属的 L2 分片；取数到 shared memory 时走 HBM → L2 → L1。
• GPC 内部支持 DSMEM（分布式共享内存，distributed shared memory）：同一 GPC 内的 SM 可访问彼此的 shared memory；GPC 之外没有 DSMEM。
• TPC 内放 2 个 SM，作用就是让这 2 个 SM 之间的 DSMEM 通信特别快。

9.3 L2 Cache

• 50 MB，分成 25 MB 分区。
• L2 是分区的：同一 GPC 内的 SM 到”直连的那个 L2 分区”路径更近——访问大多命中邻近分区时，有效延迟/带宽更好。
• 128 字节 cache line，32 字节 sector：一次内存请求可触及单条 128B line 的 1–4 个 sector。 访存”不合并”会导致 sector/request 爆炸。
• L2 吸收并合并来自 SM 的零碎小写，转成干净的大块高效写到 HBM。

9.4 统一 Shared Memory + L1（每 SM）

• 总共 256 KB、33 TB/s 带宽，分成 32 个 bank（每 bank 4 字节）。
• 可配置的 shared memory 上限 = 228 KB/SM；每 block 上限 = 227 KB（CUDA 保留 1 KB）。
• 无冲突时 load/store 达 128 B/cycle；TMA 异步拷贝能逼近峰值带宽并与计算重叠。
• bank 按连续的 4 字节字交错。
• L1 cache 充当合并缓冲（coalescing buffer）：收集 warp 请求的数据并高效交付。
• cache line = 128 B，sector = 32 B，可逐 sector 填充。

shared memory 与 L1 此消彼长：调大 shared，留给 L1 缓存的空间就变小，反之亦然。 因此 tile 大小需要权衡（见下）。

9.5 关于 shared memory 的更多细节

• 静态 shared memory（编译期数组）架构上限 48 KB（为兼容旧 compute capability）。 要超过它必须用动态 shared memory：extern __shared__。
• GEMM 的最优 tile 通常 64–128 KB/block；过大的传输会损害延迟隐藏。

9.6 Registers（寄存器）

• 每线程私有的片上寄存器：最快带宽、最低延迟，每 SM 256 KB，最多 128 读/写每 cycle。
• 寄存器经常是限制因子：若 kernel 每线程用 R 个寄存器，最大驻留线程数 ≈ floor(65536 / R)。 例如 128 regs/thread → 每 block 最多 512 线程。
• 32 位是寄存器基本单位；用 FP16/FP8 时需用打包类型把 2/4 个元素塞进 1 个寄存器。
• 功耗效率：寄存器比 shared memory 高效 30–50×，比 HBM3 高效 1000×+。
• 寄存器用量在编译期确定，不是运行时。
• 寄存器不够时溢出（spill）到 local memory（很慢）；CUDA 13.0 新增”先溢出到 shared memory、 shared 不够才回落 local memory”。
• 粒度陷阱：即使 63 → 65 regs/thread 的小变化，也可能因内部粒度取整而丢掉驻留 warp，导致性能下降。

10. SM 详解

SM 是 GPU 的基本执行单元，执行 CUDA kernel 的线程块。其组件包括：

• FP32 CUDA Core、INT/FP64 单元
• 第四代 Tensor Core
• Shared Memory / L1 cache
• L1 指令缓存
• Warp Scheduler（warp 调度器）
• Dispatch Unit（分发单元）
• Registers（寄存器）
• L0 指令缓存

10.1 Quadrant / SMSP（四个子分区）

每个 SM 划分为 4 个完全相同的子分区，叫 Quadrant 或 SMSP。每个 quadrant 含：

• 每个 quadrant 还有自己的 warp 调度器、L0 缓存、分发单元、寄存器文件。
• 每个 quadrant 每个 cycle 都能向本地单元发指令，最多同时调度 16 个 warp。

10.2 SFU（Special Function Unit，特殊函数单元）

• 每 SM 16 个 SFU（每 SMSP 4 个），每个每 cycle 发 1 条指令 → 每 SM 16 ops/cycle。
• 负责复杂数学函数：sin、cos、log、exp、sqrt、倒数等（在 CUDA Core 上算会很贵）。
• 用多项式近似 + 查找表 + 插值，牺牲一点精度换巨大吞吐。
• 注意：若所有线程同时调用数学函数，SFU 会停顿（stall）。

10.3 LSU（Load/Store Unit）

• 执行每线程的内存指令：load、store、atomic。
• 每 quadrant 8 个，共 32 个/SM，直连 L1、L2。
• warp 执行 ld/st/atom 时，LSU 合并 32 个线程的地址，形成 cache-line/sector 请求并查询 L1： 命中则快速回寄存器；未命中则继续到 L2/DRAM，并可能回填 L1。
• 合并访存时，LSU 把 warp 的 32 个地址合并成最少的 cache line，减少请求/重放/停顿，提升有效带宽。

10.4 INT / FP32 / FP64 单元

• INT 单元可与浮点数据通路并行执行：一边算地址（INT）、一边算数据（FP），互不阻塞。

10.5 L0 与 L1 指令缓存

• L0 I-Cache：指令流的微缓冲，只装很少的指令（通常几个紧凑循环）。
  • 作用：以硬件速度喂指令，避免饱和 shared memory 带宽；只求速度、不求容量。
  • 激进循环展开 / 内联可能超出 L0 容量。
  • SMSP 0 的代码不能用 SMSP 1 的 L0。
• L1 I-Cache：缓冲 SASS 机器码，让 warp 调度器始终有指令可发；解耦执行单元与存储层级 （否则每次取指令要几百周期，造成大量停顿）。

11. Warp 与指令调度

11.1 Warp（线程束）

• 一组 32 个线程：
  • 从线程块一起创建（线程 0–31 = warp 0，32–63 = warp 1，……）。
  • 每 cycle 共享一个 warp 调度器的调度。
  • 寄存器私有，但指令流（基本）共享。
• 为什么是 32：在”每个线程的控制粒度”与”每条指令的工作量”之间取平衡；32 线程也让常见访存模式天然对齐缓存/总线粒度、易于合并。

11.2 Warp Scheduler（warp 调度器）

负责指令发射：

• 每 cycle 从就绪 warp 中选一个发射指令。
• 处理 warp 级分支与分歧（divergence）。
• 维护 scoreboard，跟踪哪些 warp 真正就绪。
• 确保源寄存器在发射前已从寄存器文件读出或由流水线前递。
• 处理结构限制（如一个 warp 同时能有多少长延迟操作在飞、执行管道的可用性等）。

11.3 Warp 如何执行

warp 调度器扫描分配给该 SMSP 的 16 个 warp → 挑 1 个就绪的
→ 取指令 → 把选中的 warp ID + 程序计数器发给分发单元 → 发射到对应执行管道
→ 等待分发反馈

• warp 调度器总是尽量发射指令、让空闲 warp 忙起来。
• 实践建议：让 block 维度是 4 个 warp 的倍数（如 128 线程），以便工作均匀分布在 4 个子分区上。

11.4 Dispatch Unit 与 Dispatch Port

• Dispatch Unit：物理上把 warp 的操作发往对应功能单元（每个 cycle 发到执行管道）。
• Dispatch Port：分发单元连接具体执行管道（Tensor Core / CUDA Core 等）的物理连接点。
  • 多个执行管道竞争同一个 port → 不能同时接收指令。
  • 每 cycle 一个共享 port 只能发 1 条指令 → 分发单元需串行化发射。
  • port 只在”发射指令”时需要，不占整个执行周期。

12. 线程块在 SM/SMSP 间的分布

• GigaThread Engine 选择有足够资源的 SM 来容纳线程块。
• 硬件不会把整个 block 塞进单一 SMSP（除非 block 特别小）；而是把 block 切成 warp， 轮转（round-robin）分配到各 SMSP。
• __syncthreads() 由 shared memory 中的屏障逻辑处理（所有 SMSP 共享）： 其它 warp 进入休眠，直到所有 warp 都完成执行。

13. PCIe 5.0 主机接口

• 主 128 GB/s 数据通道，连接 H100 与主机 CPU / 系统内存。
• 连接 GPU 与网卡（NIC）的关键物理桥梁，支持 GPUDirect RDMA：网卡直接读写 GPU 内存，不占用 CPU。

14. 总结与学习衔接

14.1 H100 的本质（PDF 结论）

1. 决定性特征是”全面异步执行”：从简单同步拷贝，转向用 TMA 在后台处理数据搬移，让 SM 专注计算。
2. 专用化（Specialization）：针对每个瓶颈用专用单元——
   • TMA → 内存带宽；
   • 第四代 Tensor Core → 矩阵运算；
   • SFU → 复杂数学函数。
3. 高效的数据流：HBM3（高带宽）→ L2（50MB 统一）→ Shared Memory（低延迟）→ Registers（最高速）。

14.2 与本课程其他内容的衔接

14.3 一句话记忆

H100 = 全面异步化 + 专用单元（TMA 搬数据 / Tensor Core 算矩阵 / SFU 算函数）+ 分级存储。 学习 Hopper 编程，核心就是学会”让 TMA 在后台搬数据、让 Tensor Core 在前台算矩阵、让 SM 专注计算”。

参考来源：1. Introduction to H100.pdf（Lesson 1 - Introduction to H100，35 页）。