第 2 章 Hopper 微架构地图

读懂一个现代 GPU kernel 的前提，是先读懂它被写给的那颗芯片。 这一章的硬件数字都可以在 NVIDIA H100 Whitepaper、H100 产品规格页与 Hopper Tuning Guide 里逐条对上。

2.1 拆开一颗 H100

如果把一颗 H100 SXM5 GPU 物理拆开，你看到的是这样一块芯片：

                    H100 SXM5 物理架构
┌──────────────────────────────────────────────────────────┐
│                  GH100 GPU Die (814 mm²)                 │
│  ┌────────────────────────────────────────────────────┐  │
│  │ 8 × GPC (Graphics Processing Cluster)              │  │
│  │  每个 GPC 包含:                                    │  │
│  │   - 9 × TPC (Texture Processing Cluster)           │  │
│  │     每个 TPC 包含 2 个 SM                          │  │
│  │   合计: 8 × 9 × 2 = 144 SM 物理                    │  │
│  │   良品率筛选后启用: 132 SM (SXM5) / 114 SM (PCIe)  │  │
│  └────────────────────────────────────────────────────┘  │
│  L2 Cache: 60 MB 全尺寸 (H100 SXM5/PCIe 启用 50 MB)      │
│  HBM3 Memory Controller: 6 stacks × 16 GB = 96 GB        │
│       (SXM5 启用 5 stacks = 80 GB, 见 80GB 版本)         │
│  NVLink 4: 18 链接, 总带宽 900 GB/s                      │
│  PCIe Gen5: x16, 128 GB/s 双向                           │
└──────────────────────────────────────────────────────────┘

来源：NVIDIA H100 Tensor Core GPU Architecture Whitepaper (2022) 的 GH100 GPU 硬件架构一节与规格表。

注意 132 这个数字：H100 SXM5 实际启用 132 个 SM，但 die 上物理存在 144 个。其余 12 个被禁用是为了良品率——制造一颗 814 mm² 的 GPU die，几乎不可能全部 144 个 SM 都没缺陷，禁用一部分缺陷 SM，能把"整片可用"的苛刻要求换成"够 132 个好 SM 就行"，良品率因此大为改善。这是高端芯片的常见做法（A100 同样：物理 128 SM，启用 108 SM）。

把视角再往里拉一层：

flowchart TB
  subgraph H100 [H100 GPU]
    direction TB
    subgraph GPC8 [8 个 GPC]
      direction LR
      GPC1[GPC 0] -.- GPC2[GPC 1] -.- GPCDots[...] -.- GPC8b[GPC 7]
    end
    subgraph TPC9 [每 GPC 内 9 个 TPC]
      TPC1[TPC 0]
      TPC2[TPC 1]
      TPCDots[...]
    end
    subgraph SM2 [每 TPC 内 2 个 SM]
      SMA[SM A]
      SMB[SM B]
    end
    GPC8 --> TPC9 --> SM2
    L2[L2 Cache 50 MB]
    HBM[HBM3 80 GB / 3.35 TB/s]
    SM2 --> L2
    L2 --> HBM
  end

实际写 CUDA 代码时，GPC 和 TPC 这两个层级你基本感知不到——它们影响的主要是 L2 cache 的拓扑（L2 按分区组织，每个分区就近服务与它直连的 GPC 里的 SM），CUDA 编程模型里也没有 GPC/TPC 这两层抽象——唯一的例外是 §2.3 的 Cluster，它被保证调度在同一个 GPC 内。真正影响代码的是从 SM 这一层往下。

接下来我们就从 SM 开始往内部拆。

2.2 一颗 SM 的内部结构

Hopper SM 是整颗 GPU 中最关键的执行单元。它的内部结构如下：

                Hopper Streaming Multiprocessor (SM)
┌─────────────────────────────────────────────────────────────┐
│                                                             │
│  ┌─────────┐ ┌─────────┐ ┌─────────┐ ┌─────────┐           │
│  │  Sub-   │ │  Sub-   │ │  Sub-   │ │  Sub-   │           │
│  │  Core 0 │ │  Core 1 │ │  Core 2 │ │  Core 3 │           │
│  │         │ │         │ │         │ │         │           │
│  │ Warp    │ │ Warp    │ │ Warp    │ │ Warp    │           │
│  │ Sched   │ │ Sched   │ │ Sched   │ │ Sched   │           │
│  │         │ │         │ │         │ │         │           │
│  │ 32 FP32 │ │ 32 FP32 │ │ 32 FP32 │ │ 32 FP32 │           │
│  │ 16 INT  │ │ 16 INT  │ │ 16 INT  │ │ 16 INT  │           │
│  │ 16 FP64 │ │ 16 FP64 │ │ 16 FP64 │ │ 16 FP64 │           │
│  │ 1 TC4   │ │ 1 TC4   │ │ 1 TC4   │ │ 1 TC4   │           │
│  │ 4 SFU   │ │ 4 SFU   │ │ 4 SFU   │ │ 4 SFU   │           │
│  │         │ │         │ │         │ │         │           │
│  │ Reg     │ │ Reg     │ │ Reg     │ │ Reg     │           │
│  │ 64 KB   │ │ 64 KB   │ │ 64 KB   │ │ 64 KB   │           │
│  └─────────┘ └─────────┘ └─────────┘ └─────────┘           │
│                                                             │
│  ┌─────────────────────────────────────────────────────┐   │
│  │ TMA (Tensor Memory Accelerator)                     │   │
│  └─────────────────────────────────────────────────────┘   │
│  ┌─────────────────────────────────────────────────────┐   │
│  │ L1 Data Cache + Shared Memory  (256 KB 总, 可分配) │   │
│  │  - 配置 1: SMEM 228 KB / L1 28 KB                   │   │
│  │  - 配置 2: SMEM 196 KB / L1 60 KB ... (多档)        │   │
│  └─────────────────────────────────────────────────────┘   │
│  ┌─────────────────────────────────────────────────────┐   │
│  │ Tex / L1 Instruction Cache / Constant Cache         │   │
│  └─────────────────────────────────────────────────────┘   │
└─────────────────────────────────────────────────────────────┘

来源：NVIDIA H100 Whitepaper (2022) 的 "H100 SM Architecture" 一节。数据为单 SM。

每个 SM 被分成 4 个 Sub-Core（也叫 partition），每个 Sub-Core 内部有：

• 1 个 Warp Scheduler：决定每周期发哪条指令到哪个 warp。
• 1 个 dispatch unit：把指令真正发到执行单元。
• 64 KB 寄存器文件（SM 总计 256 KB）。
• 32 个 FP32 cores + 16 个 INT32 cores（FP32 与 INT32 是各自独立的执行单元；每个 sub-core 的 warp scheduler 每周期只发射一条 warp 指令）+ 16 个 FP64 cores。整个 SM 合计 128 FP32 / 64 INT32 / 64 FP64——H100 SXM5 官方标称的 16896 个 CUDA core 就是 132 × 128。
• 1 个第四代 Tensor Core（矩阵乘的主力）。
• 4 个 SFU（Special Function Unit）：负责超越函数（exp、log、sqrt、rsqrt 等）。
• 8 个 LD/ST unit：处理 load / store 指令（每 SM 32 个）。

四个 Sub-Core 共享：

• 256 KB 的 L1 cache + Shared Memory（可配置比例）。
• TMA 单元（整个 SM 共用）。
• 指令 cache、constant cache。

下面把每个核心组件单独打开看。

2.2.1 Tensor Core 第四代：矩阵乘的主力

Tensor Core 是 NVIDIA 在 2017 年 Volta 架构（V100）首次引入的专用矩阵乘单元，到 Hopper 已经是第四代。每一代的关键升级如下：

代际	架构	年份	关键特性
1代	Volta (V100)	2017	FP16 输入 → FP32 累加，每个 Tensor Core 每周期一次 4×4×4 矩阵乘（64 次 FMA）
2代	Turing (T4)	2018	+ INT8 / INT4，矩阵尺寸保持 4×4×4
3代	Ampere (A100)	2020	+ BF16 / TF32 / FP64、2:4 结构化稀疏，warp 级 mma.sync 指令形状 m16n8k16，引入异步拷贝 cp.async
4代	Hopper (H100)	2022	+ FP8 (E4M3/E5M2)，WGMMA 异步 mma，warp-group 级操作
5代	Blackwell (B200)	2024	+ FP4 / FP6，Tensor Memory（TMEM），CTA Pair

数据来源：各架构 Whitepaper。

第四代 Tensor Core 最关键的两个特性是 FP8 和 WGMMA：

FP8 提供两种格式：

• E4M3（4 位指数 + 3 位尾数 + 1 位符号）：动态范围窄但精度高，适合前向激活。
• E5M2（5 位指数 + 2 位尾数 + 1 位符号）：动态范围宽但精度低，适合梯度。

FP8 把 FP16 的 Tensor Core 算力直接翻倍：

H100 SXM5 Tensor Core 峰值（稠密 / 2:4 稀疏）:
  FP64:    67 TFLOPs (无稀疏)
  TF32:    495 / 989 TFLOPs
  FP16:    989 / 1979 TFLOPs
  BF16:    989 / 1979 TFLOPs
  FP8:     1979 / 3958 TFLOPs
  INT8:    1979 / 3958 TOPS

来源：NVIDIA H100 Tensor Core GPU Datasheet / 产品规格页（官方表中带 * 的是含稀疏的数字，稠密值为其一半）。

FP8 是 LLM 推理性能跃迁的关键之一。FlashAttention v3 在 H100 上跑到接近 1.2 PFLOPs/s，靠的就是 FP8。代价是数值精度降低，需要配合 per-channel scale 或者 per-tensor scale 来维持模型质量——这是第 9 章会详细讲的内容。

WGMMA（Warp-Group MMA） 是 Hopper 引入的另一个关键创新。Ampere 及之前，Tensor Core 的 mma 指令是 warp-level 的：一个 warp 里 32 个线程协作发起一次矩阵乘，矩阵尺寸 16×8×16（M×N×K）。Hopper 把这个粒度提升到 warp-group（4 个 warp = 128 个线程），每条 WGMMA 指令操作 64×N×K 的矩阵，N 取 8 的倍数、从 8 一直到 256，K 是 16（FP16/BF16）或 32（FP8）。这套形状可以在 CUTLASS 的 CuTe 层逐条数出来：cutlass-4.7.0/include/cute/arch/mma_sm90_gmma.hpp:2052 是 wgmma.mma_async.sync.aligned.m64n256k16.f32.f16.f16，cutlass-4.7.0/include/cute/arch/mma_sm90_gmma.hpp:14797 是对应的 FP8 版本 m64n256k32.f32.e4m3.e4m3。

flowchart LR
  subgraph Ampere [Ampere · warp-level mma]
    AW[1 个 warp 32 线程] --> AM[mma.sync.m16n8k16<br/>矩阵 16×8×16]
  end
  subgraph Hopper [Hopper · warp-group mma]
    HW[4 个 warp 128 线程<br/>= 1 warp-group] --> HM[wgmma.mma_async.m64n128k16<br/>矩阵 64×128×16]
  end
  Ampere -.-> Note1[每条指令 16×8×16 = 2048 次乘加]
  Hopper -.-> Note2[每条指令 64×128×16 = 131072 次乘加]

WGMMA 的革命性在于两点：

1. 算力密度：以 m64n128k16 对 m16n8k16 计，单条指令计算量是 Ampere 的 64 倍。wgmma 由 warp-group 的 4 个 warp 各自发射，折到每个 warp 调度器也还有 16 倍——同样的指令调度能力，可以喂饱更多的算力。
2. 异步性：WGMMA 是异步指令——发出去之后不阻塞，warp 可以继续做别的事（比如发起下一次 TMA 拷贝），等真正需要结果时再用 wgmma.commit_group + wgmma.wait_group 同步。这让 Tensor Core 计算和数据加载可以真正流水起来——这是 FA3 性能跃迁的核心机制。

我们会在第 12 章和第 17 章把 WGMMA 用起来。

2.2.2 寄存器文件：极端宝贵的资源

每个 SM 的寄存器文件总共 256 KB，分到 4 个 Sub-Core，每个 64 KB。换算到具体能容纳多少个 32 位寄存器：

每 Sub-Core: 64 KB / 4 B = 16384 个 32-bit 寄存器
每 SM 总计:  16384 × 4 = 65536 个 32-bit 寄存器 (= 256 KB)

这个数字看似很大，但分摊到大量 active 线程上就紧巴巴了。Hopper 一个 SM 最多支持 2048 active threads（64 warps），所以要跑满 2048 个线程，每个线程最多只能拿到 32 个寄存器。

但实际上你写一个 LLM 算子的 kernel，编译器经常给每个线程分配 64-128 个寄存器（甚至更多）。这就意味着，如果每个线程要 128 寄存器，一个 SM 上能 active 的线程数就降到 65536/128 = 512，只能跑 16 个 warp——这会显著降低 occupancy（占用率）。

寄存器压力是 CUDA 性能调优中最常见的瓶颈之一。第 12 章手写 Tensor Core GEMM 时，会反复在"用更多寄存器存 tile → 计算更快"和"用更少寄存器 → 让更多 warp 活跃"之间权衡。第 17 章引入的 setmaxnreg 指令允许 warp 之间动态调整寄存器分配，是 Hopper 的一个杀手锏。

2.2.3 L1 Cache + Shared Memory：黄金 256 KB

Hopper 每个 SM 的 L1 cache 和 Shared Memory 物理上是同一块片上 SRAM，总容量 256 KB。CUDA 提供了 cudaFuncSetAttribute API 让程序员配置二者的比例：

配置	SMEM	L1	适用场景
SMEM 优先	228 KB	28 KB	LLM 算子
平衡	100 KB	156 KB	一般通用 kernel
L1 优先	0 KB	256 KB	不显式用 SMEM 的 kernel

注意比例不是任意的：H100 支持的 SMEM 容量档位是 0 / 8 / 16 / 32 / 64 / 100 / 132 / 164 / 196 / 228 KB，驱动会把请求向上取到最近的档位。也注意"默认"不等于 228 KB——不设 cudaFuncAttributePreferredSharedMemoryCarveout 时由驱动挑档位；单个 block 要用超过 48 KB（最多 227 KB，CUDA 每 block 预留 1 KB）的 SMEM，必须显式用 cudaFuncSetAttribute 把 cudaFuncAttributeMaxDynamicSharedMemorySize 提上去（静态 SMEM 的上限始终是 48 KB）。

绝大多数 LLM 算子都跑在"SMEM 优先"配置上。原因很直接：LLM kernel 的访存模式是程序员显式控制的（典型的：把 GEMM 的 tile 显式 stage 到 SMEM），不依赖 cache 自动命中；而SMEM 容量越大，能放下的 tile 越大，HBM 访问就越少。

228 KB SMEM 听起来不算大，但放在矩阵乘里，能放下的 tile 已经很可观：

FP16 矩阵, 一个 tile 占用 = M_tile × K_tile × 2 字节
 - M_tile=128, K_tile=64:    16 KB （一个 A tile）
 - M_tile=128, K_tile=128:   32 KB （加倍 K，access 更密集）
 - M_tile=256, K_tile=128:   64 KB （加倍 M，参与更多输出）

GEMM 一般同时需要 A tile 和 B tile, 还需要 double buffer:
 - A tile (128×64) × 2 buffer + B tile (64×128) × 2 buffer = 64 KB
 - 留下 228-64 = 164 KB 给其他需求 (e.g., FA2 的 K/V tile, output staging)

这是为什么 Hopper 把 SMEM 配置上限从 A100 的 164 KB 提升到 228 KB 是个大事——它直接拉宽了 GEMM/FA 的 tile 设计空间。

L1 + SMEM 的访问延迟在几十个时钟周期量级，HBM 则是数百个周期，相差一个数量级以上（NVIDIA 官方文档没有给出 H100 的具体延迟数字，第三方微基准测得的值随测法而异，这里只取量级）。所以一个 LLM 算子的核心设计哲学就是：把数据拉到 SMEM 一次，然后让 Tensor Core 反复用，最后再写回。

2.2.4 TMA：Hopper 的异步拷贝引擎

TMA（Tensor Memory Accelerator）是 Hopper 引入的全新硬件单元。它的功能是异步地把张量数据从 HBM 拷贝到 SMEM——听起来普通，但它解决了一个困扰 GPU 编程多年的痛点。

在 Ampere 上，把数据从 HBM 异步搬到 SMEM 的写法是：

// Ampere 时代的异步拷贝
for (int i = 0; i < TILE_SIZE; i += 4) {
    __pipeline_memcpy_async(&smem[tid * 4 + i], &gmem[gid * 4 + i], 16);
}
__pipeline_commit();

这段代码有几个问题：

1. 每个线程都要参与拷贝：32 个线程的协作开销大，需要计算地址、生成 mask、处理边界。
2. 抢发射槽和寄存器：cp.async 的数据搬运本身是异步的，但"32 个线程各自算地址、各自发一条指令"这部分是同步开销——它和算术指令抢同一个发射端口，还要占掉一批地址寄存器，间接压低 occupancy。
3. 仅支持 1D 连续访问：拷贝二维 tile 时需要 unroll 出 row 数量的 async copy 指令。
4. 没有 swizzle 支持：写 SMEM 时容易撞 bank conflict，需要程序员手动 swizzle。

TMA 把所有这些问题一次性解决：

// Hopper 时代的 TMA 拷贝
// cde::cp_async_bulk_tensor_2d_global_to_shared 是 libcu++ <cuda/barrier> 里的真函数
// （cde = cuda::device::experimental；CCCL 3.2 起标为 deprecated，推荐改用
// cuda::ptx::cp_async_bulk_tensor），它发出的正是 PTX
// `cp.async.bulk.tensor.2d.shared::cluster.global.tile.mbarrier::complete_tx::bytes`。
// mbar_wait 则是示意包装，真实写法是 barrier_arrive_tx + bar.wait（底层 mbarrier.try_wait），
// 第 17 章会给完整写法。
if (threadIdx.x == 0) {
    cde::cp_async_bulk_tensor_2d_global_to_shared(
        &smem_dst,
        &tma_descriptor,  // 预先在 host 上构建的 descriptor（CUtensorMap）
        x_offset, y_offset,
        mbar              // 完成时通知的 mbarrier（cuda::barrier，按引用传入）
    );
}
mbar_wait(mbar);

TMA 的关键能力：

1. 由专用硬件单元驱动：不占用 SIMT 算术 lane，SM 在 TMA 拷贝期间可以照常发射计算指令（包括 Tensor Core 指令）。
2. 单线程发起：一个 warp 中只需要一个线程发起 TMA 指令（其他 31 个线程闲置或做别的事），避免 32 线程协作开销。
3. 1D～5D 张量原生支持：通过预先构建的 TMA descriptor，硬件直接理解多维张量的 stride、padding、swizzle。
4. 内置 swizzle 模式：把 SMEM 写入按 NVIDIA 预定义的 swizzle 模式排布，自动避免 bank conflict。
5. mbarrier 完成通知：拷贝完成时通过 shared memory 里的 mbarrier 对象通知等待的线程，不需要轮询拷贝本身的进度。

flowchart LR
  subgraph Ampere2 [Ampere · 32 线程参与拷贝]
    AT[32 个线程算地址] --> ACA[发 32 条 cp.async]
    ACA --> AS[__pipeline_commit]
    AS --> AW[__pipeline_wait]
  end
  subgraph Hopper2 [Hopper · TMA 1 线程发起]
    HT[1 个线程发 cp.async.bulk.tensor]
    HT --> TMA[TMA 硬件单元直接拷]
    TMA --> MB[mbarrier.arrive]
    MB --> MW[mbarrier.wait]
  end
  Ampere2 -.-> Note3[算术单元被拷贝占用]
  Hopper2 -.-> Note4[算术单元同时可算 Tensor Core]

TMA 是 Hopper 上写出 SOTA kernel 的必经之路。FA3 论文报告它在 H100 上比 FA2 快 1.5–2.0×，而 TMA 与 WGMMA 的异步流水正是这次提升的主要来源。第 17 章我们会用 TMA 把 FA2 重写一遍，并按 FA3 论文口径对比性能。

2.3 Cluster：分布式共享内存

到这里我们已经讨论了 SM 内部的层次。但 Hopper 在 SM 之上又新增了一层抽象——Cluster（线程块簇）。

Hopper 之前，CUDA 编程模型是：

Grid (kernel 启动一次的所有线程)
└── Block (一组协作线程, 上限 1024 线程)
    └── Warp (32 线程的硬件调度单位)
        └── Thread

Block 内的线程可以通过 SMEM 协作，Block 之间则只能通过 HBM 协作——成本极高（数百个周期）。

Hopper 引入了 Cluster 这一层：一组 Block 组成一个 Cluster，Cluster 内的 Block 可以通过分布式共享内存（DSMEM）相互访问 SMEM。CUDA 保证可移植的 cluster 上限是 8 个 Block；H100 上还可以通过 cudaFuncAttributeNonPortableClusterSizeAllowed opt-in 到 16。

flowchart TB
  subgraph Grid [Grid]
    direction TB
    subgraph Cluster1 [Cluster 0 · 可移植上限 8 Block]
      B0[Block 0]
      B1[Block 1]
      B7[... Block 7]
      B0 -.直接访问 SMEM.-> B1
      B1 -.直接访问 SMEM.-> B7
    end
    subgraph Cluster2 [Cluster 1]
      C0[Block 8]
      C1[Block 9]
    end
  end
  Cluster1 -.通过 HBM 通信.-> Cluster2

Cluster 的硬件实现并不神秘：同一个 Cluster 的所有 Block 必须调度到同一个 GPC 内的 SMs 上。GPC 内部的 SMs 通过专用的 SM-to-SM Network 直接交换 SMEM 数据，无需走 L2/HBM。

Cluster 解决了什么真实问题？最经典的例子是 大尺寸 reduce——比如对一个 64K 元素的数组求和。在 Ampere 上，要么每个 block 算完 partial sum 后各发一次全局 atomic，要么用两次 kernel launch（一次每 block 求 partial sum，一次跨 block 求和）。在 Hopper 上，可以让 8 个 Block 组成一个 Cluster（可移植上限），第二阶段先在 Cluster 内通过 DSMEM 合并各 Block 的 partial sum，再由每个 Cluster 发一次 atomic，跨 Block 的 atomic 次数按 Cluster 大小成倍减少。

第 5 章会用 Cluster 实现 reduce 的最后一版——它的价值不在"快几倍"（带宽 bound 的 reduce 到那一步已经贴着 HBM 峰值了），而在把跨 block 的 atomic 次数再除以 Cluster 大小（Cluster=8 时减为 1/8）。

2.4 Warp Specialization：把"对称协作"变成"分工合作"

Hopper 上还有一个更深层次的范式转变——Warp Specialization。

Ampere 时代的 GEMM kernel，所有 warp 做的事是对称的：

Warp 0:  Load A[tile_a] -> Load B[tile_b] -> Compute -> Store
Warp 1:  Load A[tile_a] -> Load B[tile_b] -> Compute -> Store
Warp 2:  Load A[tile_a] -> Load B[tile_b] -> Compute -> Store
...

每个 warp 都干一样的活，只是处理不同的数据 tile。这个模式在过去十几年的 CUDA 编程中是默认的。

Hopper 上推荐的写法是专业化：

Warp Group 0 (Producer):  Load A -> Load B -> Load A -> Load B -> ... (一直拷贝)
Warp Group 1 (Consumer):  WGMMA -> WGMMA -> WGMMA -> WGMMA -> ... (一直算)
Warp Group 2 (Consumer):  WGMMA -> WGMMA -> WGMMA -> WGMMA -> ... (一直算)

把 warp 分成 Producer warp（专门发起 TMA 拷贝）和 Consumer warp（专门做 Tensor Core 计算），两类 warp 通过 mbarrier 同步。由于 WGMMA 和 setmaxnreg 都以 warp-group 为单位执行，实际分工的粒度是 warp-group（第 17 章是 1 个 producer warp-group + 2 个 consumer warp-group）。这种"流水线"写法的优势：

1. Producer 和 Consumer 异步并行：TMA 拷贝和 Tensor Core 计算物理上由不同硬件单元完成，可以真正并行。对称写法下，warp 在 load 时无法 compute，反之亦然。
2. 寄存器分配可以差异化：Producer warp 不做计算，不需要存中间结果，只需要少量寄存器；Consumer warp 需要存大量 fragment，可以拿到更多寄存器。Hopper 引入了 setmaxnreg PTX 指令，让 warp 在运行时把自己的寄存器配额还回去或拿更多。
3. Tensor Core 利用率拉高：理想情况下，Consumer warp 的时间都花在 WGMMA 上，不必停下来等数据。

sequenceDiagram
    participant P as Producer Warp
    participant M as mbarrier
    participant C as Consumer Warps
    P->>P: TMA Load Tile 0
    P->>M: arrive
    M->>C: signal
    P->>P: TMA Load Tile 1 (并行)
    C->>C: WGMMA Tile 0
    P->>M: arrive
    M->>C: signal
    P->>P: TMA Load Tile 2
    C->>C: WGMMA Tile 1 (并行)
    Note over P,C: 流水线持续, 算和拷贝重叠

Warp Specialization 是 Hopper 上写 SOTA kernel 的新范式。FA3 论文 [Shah et al., 2024] 的核心贡献之一就是把 attention 改写成 Warp Specialization 形式，让 H100 的 TMA 和 WGMMA 真正流水起来。

但 Warp Specialization 也有代价：同步成本。Producer 和 Consumer 之间通过 mbarrier 同步，每次 arrive/wait 都有不可忽略的开销。如果 tile 太小（比如 32×64 这种），同步开销可能超过收益。所以 Warp Specialization 适合大 tile + 长流水的场景——这是 GEMM 和 FA 的典型形态。第 17 章会有完整的工程化讨论。

2.5 其他值得提的特性

Hopper 还有几个特性虽然不在 LLM kernel 的核心路径上，但值得知道：

2.5.1 DPX 指令：动态规划加速

Hopper 引入了一组叫 DPX（Dynamic Programming X）的指令，专门加速动态规划类算法（Smith-Waterman、Needleman-Wunsch 等基因测序、路径规划算法）。这些指令对 LLM 训练/推理基本没用，但在生物信息、机器人路径规划领域有应用。

2.5.2 Transformer Engine：自动 FP8 转换

NVIDIA 白皮书把 Transformer Engine 描述为"软件 + Hopper 定制 Tensor Core 技术"的组合；落到代码上，它是 NVIDIA 开源的 Transformer Engine 库（提供 PyTorch、JAX 等框架的模块），用 FP8 计算的层由它自动监控数值范围（amax）并管理 scale，在 FP8 与 16 位之间自动做转换。这让 FP8 训练在精度上接近 BF16，同时吃到 FP8 Tensor Core 相对 BF16 翻倍的峰值算力。Megatron-LM 等训练框架都接入了它。

2.5.3 Confidential Computing：加密计算

按 NVIDIA 白皮书的说法，H100 是第一款支持 GPU 端机密计算的 GPU——可以让 GPU 内的数据和计算对 host CPU 都加密不可见。这对企业级 AI 服务（医疗、金融）有意义，但与 kernel 优化无关。

2.6 全景速记图

把这一章所有内容整合成一张 Hopper SM 全景速记图，方便后续章节查阅：

flowchart TB
  subgraph SM [Hopper SM]
    direction TB
    subgraph Top [4 个 Sub-Core, 总计]
      W[Warp Schedulers ×4 · 每周期发 4 条指令]
      C32[FP32 cores ×128]
      CINT[INT32 cores ×64]
      C64[FP64 cores ×64]
      TC[Tensor Core 4代 ×4 · WGMMA]
      SFU[SFU ×16 · exp/log/rsqrt]
      LSU[LD/ST Unit ×32]
      REG[寄存器文件 256 KB]
    end
    subgraph Mid [SM 共享]
      TMA[TMA 单元 · 异步拷贝]
      L1SMEM[L1+SMEM 256 KB · 可配比例]
      ICACHE[I-Cache / Const Cache]
    end
  end
  Top --> Mid
  L2[L2 50 MB · 整 GPU 共享]
  Mid --> L2
  HBM[HBM3 80 GB · 3.35 TB/s]
  L2 --> HBM

记一组关键数字（在后续章节会反复引用）：

Hopper H100 SXM5 关键数字
─────────────────────────────────────
SM 数量:                        132
Warp Schedulers / SM:           4
FP32 cores / SM:                128   (4 × 32)
Tensor Core / SM (4代):         4
寄存器 / SM:                    256 KB (65536 个 32-bit)
L1 + SMEM / SM:                 256 KB (SMEM 最多 228 KB)
L2 总:                          50 MB
HBM3 总:                        80 GB
HBM3 带宽:                      3.35 TB/s
Active threads / SM 上限:       2048   (64 warps)
Active threads / GPU 上限:      270k+
Tensor Core FP16 峰值 (稠密):   989 TFLOPs/s
Tensor Core FP8 峰值 (稠密):    1979 TFLOPs/s
SIMT FP32 峰值:                 67 TFLOPs/s

数据来源：NVIDIA H100 Datasheet 与 Hopper Whitepaper（L2 50 MB 是 H100 产品启用的容量，全尺寸 GH100 为 60 MB）。

2.7 这一章的小结与下一章

这一章我们建立了精确到具体硬件单元的 Hopper 心智地图：

1. GH100 die 上有 8 GPC × 9 TPC × 2 SM = 144 SM 物理，H100 SXM5 启用 132 SM。L2 cache 50 MB，HBM3 80 GB。
2. 每个 SM 分 4 个 Sub-Core，每个 Sub-Core 有独立的 warp scheduler、寄存器、SIMT cores 和一个第四代 Tensor Core。
3. 第四代 Tensor Core 引入 FP8 和 WGMMA：FP8 把稠密算力翻倍到 1979 TFLOPs，WGMMA 是异步指令让算和拷贝可以流水。
4. TMA 是 Hopper 的异步拷贝引擎：单线程发起、专用硬件、原生支持多维张量、内置 swizzle，让 SMEM 拷贝从"占用算术单元的同步操作"变成"完全后台运行的异步操作"。
5. Cluster 是 Block 之上的新一层抽象：同 GPC 内的一组 Block（可移植上限 8，H100 opt-in 到 16）可通过分布式共享内存协作，省去跨 Block 必须走 HBM 的痛点。
6. Warp Specialization 是 Hopper 的推荐编程范式：把 warp 分成 Producer/Consumer 流水，让 TMA 和 Tensor Core 真正异步并行。

第 3 章我们会换一个视角——从硬件视角切换到编程模型视角。读者已经知道硬件长什么样，现在我们看 CUDA 给程序员暴露了哪些抽象层次：host → device → grid → cluster → block → warp → thread 这七层之间的关系，以及它们各自对应到哪些 PTX/CUDA C++ API。

本章动手练习：

1. 跑 CUDA Samples 里的 deviceQuery（nvidia-smi -q 只给型号、显存、时钟等，不列 SM 数），对照本章数字，看看自己的 SM 数、每 SM 的 SMEM 容量、compute capability（9.0 即 Hopper、第四代 Tensor Core）。
2. 用 cudaDeviceGetAttribute(&val, cudaDevAttrMaxSharedMemoryPerBlockOptin, 0) 查询你 GPU 的 SMEM Optin 上限——这是每 block 的值：Hopper 是 227 KB、Ampere 是 163 KB（各自比「每 SM」的 228 / 164 少 1 KB，CUDA 每 block 预留）。
3. 阅读 NVIDIA 官方 CUDA Programming Guide 的 "Asynchronous Data Copies" 一节（含 TMA）与同目录下的 "Asynchronous Barriers"，对应到本章 TMA / mbarrier 的描述。