第 2 章 Hopper 微架构地图

"Architecture is what enables algorithms. To understand modern GPU algorithms, you must understand the hardware they are written for." ——NVIDIA 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: 50 MB (分两块, 各 25 MB, NUMA-like)           │
│  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, Figure 4 / Table 2.

注意 132 这个数字：H100 SXM5 实际启用 132 个 SM，但 die 上物理存在 144 个。其余 12 个被禁用是为了良品率——制造一颗 814 mm² 的 GPU die，几乎不可能全部 144 个 SM 都没缺陷，禁用一部分能让芯片整体良品率从 ~10% 提升到 ~70%。这是高端芯片的常见做法（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 的拓扑（同一 GPC 内的 SM 共享 L2 的同一段更快），但 CUDA 编程模型里没有 GPC/TPC 这两层抽象。真正影响代码的是从 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   │           │
│  │         │ │         │ │         │ │         │           │
│  │ 16 FP32 │ │ 16 FP32 │ │ 16 FP32 │ │ 16 FP32 │           │
│  │ 16 INT  │ │ 16 INT  │ │ 16 INT  │ │ 16 INT  │           │
│  │ 8 FP64  │ │ 8 FP64  │ │ 8 FP64  │ │ 8 FP64  │           │
│  │ 1 TC4   │ │ 1 TC4   │ │ 1 TC4   │ │ 1 TC4   │           │
│  │ 1 SFU   │ │ 1 SFU   │ │ 1 SFU   │ │ 1 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 192 KB / L1 64 KB ... (多种)        │   │
│  └─────────────────────────────────────────────────────┘   │
│  ┌─────────────────────────────────────────────────────┐   │
│  │ Tex / L1 Instruction Cache / Constant Cache         │   │
│  └─────────────────────────────────────────────────────┘   │
└─────────────────────────────────────────────────────────────┘

来源：NVIDIA Hopper Whitepaper, 2022, Figure 7. 数据为单 SM。

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

• 1 个 Warp Scheduler：决定每周期发哪条指令到哪个 warp。
• 1 个 dispatch unit：把指令真正发到执行单元。
• 64 KB 寄存器文件（SM 总计 256 KB）。
• 16 个 FP32 cores + 16 个 INT32 cores（注意：FP32 和 INT32 共享发射端口，不能同周期发射）+ 8 个 FP64 cores。
• 1 个第四代 Tensor Core（占据每个 sub-core 大量晶体管，是矩阵乘的主力）。
• 1 个 SFU（Special Function Unit）：负责超越函数（exp、log、sqrt、rsqrt 等）。
• 1 个 LD/ST unit：处理 load / store 指令。

四个 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 已经是第四代。每一代的关键升级如下：

数据来源：各架构 Whitepaper。

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

FP8 提供两种格式：

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

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

H100 SXM5 Tensor Core 峰值（带稀疏不带稀疏）:
  FP64:    67 TFLOPs (无稀疏)
  TF32:    495 / 989 TFLOPs (无稀疏 / 2:4 稀疏)
  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/16/24/...256，K 是 16（FP16）或 32（FP8）。

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 FLOPs]
  Hopper -.-> Note2[每条指令算 64×128×16 = 131072 FLOPs]

WGMMA 的革命性在于两点：

1. 算力密度：单条指令计算量是 Ampere 的 64 倍。这意味着 Tensor Core 的"输入指令带宽"压力降低 64 倍——同样的指令调度能力，可以喂饱更多的算力。
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 总计:  64K × 4 = 65536 个 32-bit 寄存器

这个数字看似很大，但分摊到大量 active 线程上就紧巴巴了。Hopper 一个 SM 最多支持 2048 active threads（64 warps），所以每个线程理论最多能拿到 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 让程序员配置二者的比例：

绝大多数 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 配置范围从 Ampere 的 164 KB 提升到 228 KB 是个大事——它直接拉宽了 GEMM/FA 的 tile 设计空间。

L1 + SMEM 的访问延迟是 20-30 cycles，对比 HBM 的 ~800 cycles，是 30-40 倍的差距。所以一个 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. 占用算术单元：拷贝指令和算术指令共享同一个发射端口。拷贝时不能算。
3. 仅支持 1D 连续访问：拷贝二维 tile 时需要 unroll 出 row 数量的 async copy 指令。
4. 没有 swizzle 支持：写 SMEM 时容易撞 bank conflict，需要程序员手动 swizzle。

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

// Hopper 时代的 TMA 拷贝
if (threadIdx.x == 0) {
    cp_async_bulk_tensor_2d_global_to_shared(
        &smem_dst,
        &tma_descriptor,  // 预先在 host 上构建的 descriptor
        x_offset, y_offset,
        &mbar             // 完成时通知的 mbarrier
    );
}
mbar_wait(&mbar);

TMA 的关键能力：

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

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 快 ~2x，TMA 贡献了至少一半的提升。第 17 章我们会用 TMA 把 FA2 重写一遍，对比同一份 kernel 用/不用 TMA 的性能差距。

2.3 Cluster：分布式共享内存

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

Ampere 之前，CUDA 编程模型是：

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

Block 内的线程可以通过 SMEM 协作，Block 之间则只能通过 HBM 协作——成本极高（800+ cycles）。

Hopper 引入了 Cluster 这一层：最多 16 个 Block 组成一个 Cluster，Cluster 内的 Block 可以通过分布式共享内存（DSMEM）相互访问 SMEM。

flowchart TB
  subgraph Grid [Grid]
    direction TB
    subgraph Cluster1 [Cluster 0 · 最多 16 Block]
      B0[Block 0]
      B1[Block 1]
      B15[... Block 15]
      B0 -.直接访问 SMEM.-> B1
      B1 -.直接访问 SMEM.-> B15
    end
    subgraph Cluster2 [Cluster 1]
      C0[Block 16]
      C1[Block 17]
    end
  end
  Cluster1 -.通过 HBM 通信.-> Cluster2

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

Cluster 解决了什么真实问题？最经典的例子是 大尺寸 reduce——比如对一个 64K 元素的数组求和。在 Ampere 上，要么用 atomic（全局原子操作慢），要么用两次 kernel launch（一次每 block 求 partial sum，一次跨 block 求和）。在 Hopper 上，可以让 16 个 Block 组成一个 Cluster，第二阶段直接在 Cluster 内通过 DSMEM 求 partial sum，省去一次 kernel launch 的开销。

第 5 章会用 Cluster 实现一个比传统 reduce 快 3 倍的版本。

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 0 (Producer):  Load A -> Load B -> Load A -> Load B -> ... (一直拷贝)
Warp 1 (Consumer):  WGMMA -> WGMMA -> WGMMA -> WGMMA -> ... (一直算)
Warp 2 (Consumer):  WGMMA -> WGMMA -> WGMMA -> WGMMA -> ... (一直算)
Warp 3 (Consumer):  WGMMA -> WGMMA -> WGMMA -> WGMMA -> ... (一直算)

把 warp 分成 Producer warp（专门发起 TMA 拷贝）和 Consumer warp（专门做 Tensor Core 计算），两类 warp 通过 mbarrier 同步。这种"流水线"写法的优势：

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，没有任何 idle 周期。

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 有几十 cycle 开销。如果 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 转换

Transformer Engine 是 NVIDIA 提供的一个软件库（不是硬件特性），它会自动监控 LLM 训练过程中各层的数值范围，动态决定哪些层用 FP8、哪些层用 BF16，并自动管理 scale。这让 FP8 训练在精度上接近 BF16，速度上提升 2x。Megatron-LM、DeepSpeed、PyTorch 都集成了它。

2.5.3 Confidential Computing：加密计算

Hopper 是第一代支持 GPU 端机密计算的 NVIDIA 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 ×64]
      CINT[INT32 cores ×64]
      C64[FP64 cores ×32]
      TC[Tensor Core 4代 ×4 · WGMMA]
      SFU[SFU ×4 · exp/log/rsqrt]
      LSU[LD/ST Unit ×4]
      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:                64    (4 × 16)
Tensor Core / SM (4代):         4
寄存器 / SM:                    256 KB (65536 个 32-bit)
L1 + SMEM / SM:                 256 KB (228 KB 配 SMEM 时)
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。

2.7 这一章的小结与下一章

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

1. GH100 die 上有 8 GPC × 9 TPC × 2 SM = 144 SM 物理，启用 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 内的 16 个 Block 可通过分布式共享内存协作，省去跨 Block 必须走 HBM 的痛点。
6. Warp Specialization 是 Hopper 的推荐编程范式：把 warp 分成 Producer/Consumer 流水，让 TMA 和 Tensor Core 真正异步并行。

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

本章动手练习：

1. 跑 nvidia-smi -q 输出你卡的详细信息，对照本章数字，看看自己的 SM 数、SMEM 配置、Tensor Core 代际。
2. 用 cudaDeviceGetAttribute(&val, cudaDevAttrMaxSharedMemoryPerBlockOptin, 0) 查询你 GPU 的 SMEM Optin 上限。Hopper 应该是 228 KB，Ampere 是 164 KB。
3. 阅读 NVIDIA 官方 CUDA C++ Programming Guide 7.30 节 "Asynchronous SIMT Programming Model"，对应到本章 TMA / mbarrier 的描述。