Blackwell GPU 执行模型详解

GPU Execution Model — Blackwell GPU 执行模型详解

来源: Modern GPU Programming For MLSys (MLC Community)
章节: Part I, Chapter 1 — GPU Execution Model
目标架构: NVIDIA Blackwell
整理时间: 2026-06-24

📌 核心摘要

本章系统介绍了现代 GPU（以 Blackwell 架构为代表）的执行模型，涵盖三个核心维度：

1. 线程层级（Thread Hierarchy）：从 Thread 到 Grid 的嵌套分组结构
2. 内存空间（Memory Spaces）：GMEM → SMEM → TMEM → Register 的数据通路
3. 计算引擎（Compute Engines）：CUDA Cores 与 Tensor Cores 的分工协作

核心思想：内核是一个流水线，数据在这些内存空间之间流动，并在独立的计算与数据传输引擎之间交接工作。反复的目标是让这些引擎同时保持忙碌。

一、执行层级（The Execution Hierarchy）

GPU 不将数千个线程呈现为扁平池，而是分组为嵌套层级结构，每一层级的存在都是为了在特定尺度上使协作成本最低。

1.1 线程层级结构

1.2 各层级详解

1.3 Blackwell 的关键变化

Blackwell 的关键操作不是由同一组线程发起的，每种操作有其天然的粒度。
| 操作 | 发起粒度 | 描述 |
|------|----------|------|
| TMA 拷贝 | 单线程 | 一个线程发出命令，硬件引擎搬运整个 Tile |
| TMEM ↔ 寄存器 | Warpgroup | 4 个 Warp 协作，每个 Warp 移动 TMEM Tile 的自己的切片 |
| tcgen05 MMA | 选举线程 | 一个选举出的线程提交，Tensor Core 执行 |
| Cluster MMA | 2-CTA | 跨两个 CTA 的协作 MMA |
Scope（作用域）：运行某个操作的线程集合，是本书三大核心设计要素（Scope / Layout / Dispatch）中的第一个。

二、内存空间（Memory Spaces）

不存在同时兼具大容量和高速度的单一内存。物理规律决定了容量与速度的权衡，GPU 因此提供多种内存，每种在权衡曲线上处于不同位置。

2.1 内存层级总览

2.2 数据通路

GMEM → SMEM → [计算] → 寄存器 → SMEM → GMEM
↑
TMEM（存放累加器）

2.3 TMEM 详解（Blackwell 新增）

TMEM 是 Blackwell 独有的内存空间，在之前架构上没有对应物。
为什么需要 TMEM？

• 之前的 GPU 将大型 MMA 累加器存放在寄存器中，与线程的其他值争夺稀缺资源
• 寄存器是固定的每线程资源，随着 MMA Tile 增大，累加器 Fragment 也增大，开始挤占线程需要的其他值
• 更大的 Tile 对 Tensor Core 吞吐量有利，但把整个累加器保持在寄存器中使大 Tile 更难使用
  TMEM 设计：
• CTA 作用域的 2D Scratchpad：128 Lane × 最多 512 列（32-bit）
• tcgen05.mma 将累加器写入 TMEM 而非寄存器
• 内核需在 Epilogue 阶段将 TMEM 显式读回寄存器
  两个重要后果：

1. TMEM 读取是显式的、Warpgroup 分布式的——由 4 个 Warp 协作完成
2. TMEM 必须显式分配和释放（不像寄存器由编译器自动分配）

三、Cluster 的分布式共享内存（DSMEM）

3.1 问题

• 单个 CTA 运行在一个 SM 上，使用该 SM 的 SMEM
• 但单个 CTA 的 SMEM 预算是有限的
• 大 Tile 常常需要比单个 Block 能提供的更多的操作数存储或复用

3.2 解决方案：Thread Block Cluster

Hopper 引入了 Thread Block Cluster：一组 CTA 比独立 Block 更紧密地协作：

• 可以彼此同步
• 可以读写彼此的共享内存（DSMEM）
  Blackwell 保留了 Cluster 并增强了它：
• 动态调度（Cluster Launch Control）
• 2-CTA 协作 MMA

3.3 DSMEM 工作机制

• CTA 可以寻址和访问对等 CTA 的 SMEM
• 线程可以在对等 CTA 的 SMEM 中指定位置，直接将 Tile 从自己的 SMEM 批量拷贝到对等的 SMEM
• 字节落地后触发完成 Barrier
• 2-CTA Cluster GEMM 正是基于此机制，在两个 CTA 之间共享操作数 Tile，无需通过全局内存中转

四、计算引擎：CUDA Cores 与 Tensor Cores

SM 提供两种不同的数学引擎，而非一种。两者之间的分工塑造了几乎每个内核的编写方式。

4.1 两种引擎对比

4.2 为什么这个分工很重要

Tensor Cores 的算术吞吐量比 CUDA Cores 高约 10 倍以上（FLOP/s）。密集线性代数（GEMM、卷积、Attention）只有在 Tensor Cores 上运行时才能达到峰值性能。
获取性能的关键在于：保持 Tensor Cores 持续有数据可消费。

4.3 代际演进

4.4 Cluster 对引擎的扩展
2-CTA 协作 MMA：两个 CTA 各自贡献 SMEM 操作数到单个更大的 Tensor Core MMA Tile
TMA 多播：数据移动引擎的一次加载可将同一 GMEM Tile 交付给多个 CTA，消除独立加载产生的冗余全局流量

五、GEMM 数据流水线（The GEMM Data Pipeline）

5.1 三阶段流水线

以通用矩阵乘法（GEMM）为例，单个 GEMM Tile 流经三个阶段：

┌─────────────────────────────────────────────────────────┐
│ GEMM 三阶段流水线 │
│ │
│ ┌──────────┐ ┌──────────┐ ┌──────────┐ │
│ │ Load │ → │ Compute │ → │ Epilogue │ │
│ │ │ │ │ │ │ │
│ │ TMA 拷贝 │ │ tcgen05 │ │ TMEM → │ │
│ │ GMEM → │ │ MMA │ │ 寄存器 → │ │
│ │ SMEM │ │ SMEM → │ │ GMEM │ │
│ │ │ │ TMEM │ │ │ │
│ └──────────┘ └──────────┘ └──────────┘ │
│ 单线程发起 选举线程发起 Warpgroup 协作 │
│ 字节计数Barrier 完成Barrier TMA Store │
└─────────────────────────────────────────────────────────┘

阶段 1：Load

1. TMA 拷贝将 A 或 B 操作数 Tile 从 GMEM 流式传输到 SMEM
   一个线程发出拷贝，提前记录预期到达的字节数
   TMA 引擎报告进度，完成 Barrier 在所有预期字节交付后才翻转

阶段 2：Compute

tcgen05 MMA 从 SMEM 读取操作数 Tile，将乘积累加到 TMEM Tile
一个选举出的线程提交，完成后发出 Barrier 信号

阶段 3：Epilogue

Warpgroup 将 TMEM 累加器读回寄存器
将结果转换为目标输出数据类型
存储到 GMEM（通常通过 SMEM 中转并发出 TMA Store）

5.2 慢内核 vs 快内核

写成这样三个阶段看起来是严格串行的，但慢内核和快内核的全部区别在于重叠（Overlap）。
慢内核（Naive）：

加载 → 等待 → 计算 → 等待 → 存储

每个引擎在等待前一个引擎完成时空闲。

快内核（Pipeline）：

时间轴：
TMA:    [Tile k-1 Load] [Tile k Load]   [Tile k+1 Load]
MMA:          [Tile k-1 Compute] [Tile k Compute]   [Tile k+1 Compute]
Epilogue:         [Tile k-2 Writeback] [Tile k-1 Writeback] [Tile k Writeback]

• Tensor Core 在计算 Tile k 时
• TMA 引擎已经在获取 Tile k+1
• Epilogue 正在排出 Tile k-1
• 三个引擎同时保持忙碌

5.3 安全交接的关键

让三个异步引擎安全交接工作，正是 Barrier 和 Phase 模型（Async Coordination: mbarriers）的职责。Part III 的 GEMM 优化阶梯正是建立在此基础之上。

六、核心设计要素

贯穿全书的三个核心设计要素：