GPU 架构基础

在编写或调优 GPU kernel 之前，你需要先建立一个可用的心智模型，理解 GPU 实际上是如何运行代码的。 否则，诸如“提高 occupancy”或“减少 shared memory bank conflict”之类的建议，就只会变成一套需要死记硬背的规则。 你并不能真正理解它们何时适用、何时不适用。

本页会在对 kernel 工作有用的层次上介绍现代 GPU 的执行模型与内存系统。由于 CUDA 当前主导着 LLM 推理生态，文中的细节会偏向 NVIDIA 架构。不过，其中的核心概念同样广泛适用于 AMD GPU 以及其他加速器。

GPU 执行与内存示意图

了解工作如何组织、数据位于何处，从单个线程一直到全局内存。将鼠标悬停在任一区域即可高亮显示。

→→→

片上GPU 晶粒

SM 0核心计算单元

线程块 0

Warp 032 线程，锁步执行

0

1

2

3

4

5

6

7

8

9

10

11

12

13

14

15

16

17

18

19

20

21

22

23

24

25

26

27

28

29

30

31

Warp 132 线程

⋮ 最多 32 个 warp（最多 1024 线程）

线程块 1运行在同一个 SM 上

寄存器文件

256 KB · 每线程私有

共享内存

每个 block 一份 · 由程序员管理

L1 cache

每个 SM 一份 · 由硬件管理

Warp 调度器

每个周期发射已就绪的 warp

SM 1

线程块 2

寄存器 · 共享内存 · L1 · 调度器

SM 2

线程块 3

寄存器 · 共享内存 · L1 · 调度器

SM 3

线程块 4

寄存器 · 共享内存 · L1 · 调度器

最多 132 个 SM
(H100 SXM)

L2 Cache50 MB · 全部 SM 共享 · 由硬件管理

↕内存总线

片外HBM（全局内存 / VRAM）

80 GB · 3.35 TB/s (H100 SXM)

模型权重

KV cache

激活值

中间缓冲区

容量最大，但速度最慢。访问延迟约为 400+ 个周期。Kernel 优化的重点是尽量减少这里的往返访问。

GPU 如何运行代码

高层次的思路很简单。GPU 会让数千个轻量级 thread 同时处于 in-flight 状态，从而隐藏 latency， 并维持高吞吐量。

Thread、warp 和 thread block

当你启动一个 GPU kernel 时，会创建大量 thread。它们会按一种简单的层级结构组织起来，以帮助 GPU 高效执行。

• Thread：最小的执行单元。每个 thread 运行相同的 kernel 代码，但处理不同的数据。这称为 SIMT（Single Instruction, Multiple Threads）。

  你不会手动给每个 thread 分配任务。相反，每个 thread 通常会结合它在 block 中的位置 （threadIdx）以及 block 的位置（blockIdx），来计算自己应处理输入的哪一部分。

• Warp：由 32 个 thread（在 NVIDIA GPU 上）组成的一组执行单元，它们以 lockstep 方式执行指令。 warp 才是真正的调度单元。它一次只能执行一条指令。如果同一 warp 内的 thread 走了不同的代码路径 （branch divergence），那么 warp 就必须按顺序执行两条路径。这会降低效率，因为在任意时刻总会有一些 thread 处于非活动状态。

  不同的 warp 彼此独立，可以在同一时刻执行不同的指令。

• Thread block（或简称 block）：一组可以通过 shared memory 和 synchronization barrier 协作的 thread。一个 block 最多可以包含 1024 个 thread（32 个 warp）。其大小可通过 blockDim 获取。 一个 block 中的所有 thread 都会被调度到同一个 Streaming Multiprocessor（SM）上，但它们可能是分时执行的，而不是一次性同时执行。

Grid 与索引

启动 kernel 时，你需要定义一个由 block 组成的 grid：

kernel<<<gridDim, blockDim>>>();

• gridDim：grid 的维度，也就是每个维度上的 block 数量
• blockDim：block 的维度，也就是一个 block 在每个维度上的 thread 数量

GPU 会自动分配 ID：

• blockIdx：block 在 grid 中的位置
• threadIdx：thread 在 block 中的位置

每个 thread 将它们组合起来得到一个全局索引：

int i = blockIdx.x * blockDim.x + threadIdx.x;

这种索引机制使得 SIMT 模型在实践中可用。所有 thread 都运行相同的代码，但每个 thread 会根据自己在 grid 中的位置推导出不同的索引，并处理各自的数据。

Streaming Multiprocessor（SM）

SM 是 GPU 的核心计算单元。每个 SM 包含：

• 一组用于整数和浮点运算的算术执行单元（在 NVIDIA GPU 上是 CUDA core，在 AMD GPU 上是 execution unit）
• 用于加速矩阵运算的 tensor core（现代架构中提供）
• 每个周期选择就绪 warp 并发射指令的 warp scheduler
• register file、shared memory 和 L1 cache（在很多架构上，shared memory 与 L1 共享片上资源，并且可以配置。下文会展开说明）

On-chip 指的是物理上位于 GPU 硅片内部、紧邻计算单元的位置。Off-chip 指的是位于 GPU 芯片之外， 需要跨越内存接口（连线、控制器）进行访问。

现代数据中心 GPU 拥有许多 SM。例如，NVIDIA H100 SXM 有 132 个 SM，A100 有 108 个。 GPU 的总吞吐量取决于你的 kernel 能否让这些 SM 始终在做有用的工作。

只要还有足够的 register、shared memory 和 warp slot，每个 SM 就可以并发驻留并执行多个 thread block。

占用率（Occupancy）

Occupancy 用来衡量一个 SM 上活动 warp 数量相对于其可支持最大值的比例。更高的 occupancy 会给 warp scheduler 更多可选 warp，从而有助于隐藏内存 latency。当一个 warp 在等待数据（例如来自 HBM） 时，SM 可以执行另一个 warp 的指令。

不过，更高的 occupancy 并不总是更好。一个 kernel 如果每个 thread 使用更多 register， 或者每个 block 使用更多 shared memory，那么 occupancy 会更低，但它仍然可能跑得更快， 因为每次内存访问时每个 thread 做了更多有用工作。FlashAttention 就是一个很好的例子： 它会有意使用大量 shared memory，将数据保留在片上，用牺牲 occupancy 的方式换取大幅减少 HBM 访问。

Occupancy 是一种隐藏 latency 的工具，而不是一个应该盲目最大化的性能指标。先 profile，再判断 occupancy 是否真的是你的瓶颈。

内存层次结构

GPU 拥有很深的 memory hierarchy。理解它至关重要，因为大多数 kernel 瓶颈都与内存有关，而不是计算本身。

寄存器（Register）

Register 是 GPU 上最快的存储。每个 thread 都拥有逻辑上私有的 register，这意味着其他 thread 无法访问它们。 但从物理上看，这些 register 来自 SM 上一个大型的 register file（例如 H100 上每个 SM 为 256 KB）， 它由所有 resident thread 共享，并在它们之间进行划分。

访问 register 的成本远低于访问 shared memory 或 HBM，因为数据已经在片上，并且可被计算单元直接使用。

Register 是有限资源。每个 thread 使用的 register 越多，SM 能并发运行的 thread（以及 warp）就越少。 这是 occupancy 权衡的一个主要驱动因素。

Shared memory（SMEM）与 L1 cache

现代 GPU 在每个 SM 内都提供两类高速片上内存：shared memory（SMEM）和 L1 cache。它们通常使用同一块物理 SRAM， 但扮演的角色不同。

L1 cache 由硬件管理。当 thread 从 global memory（HBM）读取数据时，GPU 可能会自动把数据存入 L1。 如果之后再次访问相同数据，就可以更快地直接从 L1 提供（cache hit），而不必回到 HBM。这个过程对程序员是透明的： 你不会显式把数据加载到 L1，也无法控制哪些数据会留在里面。因此，L1 最适合被视为一种 best-effort 优化， 当内存访问模式存在时间局部性或空间局部性时，它可以提升性能。

相比之下，shared memory 由程序员管理。它是一块小型、显式分配的内存空间，由同一个 thread block 中的所有 thread 共享。 thread 可以直接读写 shared memory，并使用同步机制（例如 __syncthreads()）来协调访问。这使 shared memory 成为 thread 间协作时一个可预测、可控的工作区。

shared memory 的主要目的，是减少代价高昂的 HBM 访问。一种常见模式是先把数据从 global memory 加载到 shared memory 一次， 然后让多个 thread 多次复用。由于 shared memory 位于片上，且速度远快于 HBM，这通常能显著提升性能。这个模式在许多高性能 kernel 中都很常见，例如矩阵乘法和 attention 机制。

shared memory 被组织为多个 bank（通常是 32 个）。当一个 warp 中的多个 thread 同时访问同一个 bank 时，就会发生 bank conflict， 这些访问将被串行化。避免 bank conflict 是 kernel 调优中的常见微优化。

下面是一个对比：

	Shared memory	L1 cache
管理方式	程序员	硬件
作用范围	Thread block	每个 SM（SM 上的所有 block）
持续时间	block 生命周期	由硬件驱逐
Bank conflict	是，可能发生	否（由硬件处理）
最适合	可提前规划的复用	不规则或不可预测的访问

一个常见问题是：既然已经有了 register，为什么还需要 shared memory（SMEM）和 L1 cache？

Register 的确是 GPU 上最快的存储，但单靠它们还不够。

1. Register 是私有的。每个 thread 都有自己的 register，因此数据无法共享。如果多个 thread 需要同一份数据， 它们就必须从 global memory 重新加载。L1 cache 和 shared memory 则支持 thread 之间的数据复用。
2. Register 数量有限。每个 thread 能拿到的 register 数量很少。较大的数据（例如权重、tile）装不下， 因此必须来自内存。L1 可以自动缓存这些数据，而 shared memory 则允许你显式地暂存并复用它们。
3. **Thread 之间无法协调。**Thread 不能通过 register 相互通信。shared memory 提供了一个用于协作与同步的共享工作区。

L2 缓存

L2 cache 是最大的片上内存，由 GPU 上所有 SM 共享，并在访问 HBM 之前充当最后一道高速缓存。H100 的 L2 cache 大约为 50 MB，A100 大约为 40 MB。

L2 cache 由硬件管理。你不会显式把数据加载进去。相反，它会自动缓存最近的 global memory 访问。 对于那些在 thread block 之间存在一定数据复用的 workload，L2 可以显著减少 HBM traffic。

L1 cache 和 shared memory 主要在单个 SM 内发挥作用，而 L2 的存在是为了捕获跨 SM、跨 thread block 的复用。 那些放不进片上内存、或者会被多个 block 访问的数据，仍然可以通过 L2 复用，而不必反复从 HBM 获取。

在 LLM 推理中，L2 cache 可以帮助处理：

• 被多个 attention head 访问的 KV cache 条目
• 在 batch 元素之间复用的权重 tile
• 较小的查找表或元数据

不过，当 working set 非常大时（这在 LLM 推理中很常见），L2 hit rate 会下降，而 HBM 带宽会成为真正的约束。

HBM（High Bandwidth Memory）

HBM 是 GPU 的主内存（通常也称为 VRAM 或 global memory）。它存储模型权重、KV cache、activation 以及其他所有大型数据结构。 现代数据中心 GPU 使用 HBM2e 或 HBM3：

GPU	HBM 容量	HBM 带宽
A100 SXM	80 GB	2.0 TB/s
H100 SXM	80 GB	3.35 TB/s
H200 SXM	141 GB	4.8 TB/s

HBM 容量很大，但相对于片上内存来说速度较慢。一次 HBM 访问通常要花费数百个 cycle。 这就是为什么 kernel 优化会聚焦于尽量减少 HBM 与计算单元之间的数据传输。

---

memory hierarchy 构成了一个金字塔。

层级	大小（H100）	带宽	作用范围	延迟	管理方式
Registers	每个 SM 256 KB	最高	每个 thread	~1 cycle	编译器
Shared memory / L1	每个 SM 最多 228 KB	有效约 ~20 TB/s	每个 block（SMEM）/ SM（L1）	~20-30 cycles	程序员（SMEM）/ 硬件（L1）
L2 cache	50 MB	~12 TB/s	所有 SM	~200 cycles	硬件
HBM	80 GB	3.35 TB/s	全局	~400+ cycles	程序员

所有 kernel 优化技术，不管是 tiling、fusion 还是 data layout 调整，本质上都在做同一件事： 把数据访问尽可能往这座金字塔的上层移动，也就是尽量把高频使用的数据留在 register 或 shared memory 中， 而不是反复从 HBM 读取。

Tensor Core

从 Volta 架构（2017）开始，NVIDIA GPU 引入了 tensor core：这是一类专为矩阵乘加运算设计的专用硬件单元。 它们作用于较小的矩阵 tile（例如 16×16 或 8×8，取决于数据类型），在矩阵计算上的吞吐量远高于标准 CUDA core。

对于 LLM 推理而言，tensor core 很重要，因为：

• Transformer 中的核心计算就是矩阵乘法（projection、attention score、feed-forward layer）
• tensor core 支持混合精度格式（FP16、BF16、FP8、INT8），可以减少 memory traffic 并提高吞吐量
• 在现代 GPU 上，tensor core 吞吐量与标准 core 吞吐量之间的差距非常大。在 H100 级别 GPU 上，FP16 和 BF16 tensor core 的吞吐量显著高于标准 FP32 吞吐量

不过，要高效使用 tensor core，你的数据必须具备正确的格式和 layout。tile 未对齐、数据类型错误、 或者内存访问模式不理想，都可能阻止编译器或库将你的计算映射到 tensor core 上。

这也是为什么会有 cuBLAS 和 Triton 这类框架存在。它们负责处理把操作映射到 tensor core 上的复杂性， 从而避免你在每个 kernel 中都手工管理这些细节。

常见问题

一个 warp 中的所有 thread 会不会连数据都做完全一样的事情？

它们执行的是相同的操作，但处理的不是相同的数据。如果用的是相同数据，那么并行就毫无意义。例如，对于这个 kernel：

C[i] = A[i] + B[i];

你会启动 32 个 thread（一个 warp）。实际情况如下：

• Thread 0 计算 C[0] = A[0] + B[0]
• Thread 1 计算 C[1] = A[1] + B[1]
• ...
• Thread 31 计算 C[31] = A[31] + B[31]

在一个 warp 内，相同的指令是这里的 add，但数据不同。每个 thread 都使用不同的索引（i）。

CUDA core 和 tensor core 有什么区别？

CUDA core 是用于标准浮点和整数运算的通用执行单元。tensor core 是面向小 tile 上矩阵乘加运算的专用单元 （例如 4×4 或 16×16，取决于架构和数据类型）。对于 Transformer layer 这类矩阵密集型 workload， tensor core 提供的吞吐量显著高于 CUDA core。

延伸资源

• CUDA C++ Programming Guide
• NVIDIA H100 Tensor Core GPU Architecture Whitepaper
• 究竟什么是“CUDA”？