第二课：GPU 程序是怎样被执行的

导语：下面继续讲 GPU 程序从启动到执行完成的全过程，重点理清 Kernel、Grid、Block、Warp、Thread 与 SM 的对应关系。

先看一个最简单的数组加法：

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

假设数组中有 1000 个元素，希望用 GPU 并行完成。

对应的 CUDA Kernel 可以写成：

__global__ void vector_add(
    const float* A,
    const float* B,
    float* C,
    int N)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;

    if (i < N) {
        C[i] = A[i] + B[i];
    }
}

CPU 启动这个 Kernel：

vector_add<<<4, 256>>>(A, B, C, 1000);

这里的：

<<<4, 256>>>

表示：

• 创建 4 个 Block；
• 每个 Block 有 256 个 Thread；
• 总共创建 $4\times256=1024$ 个 Thread。

---

一、Kernel 是什么

Kernel 是运行在 GPU 上的函数。

普通 CPU 函数：

void function();

CUDA Kernel：

__global__ void kernel();

Kernel 通常由 CPU 发起，然后交给 GPU 执行：

CPU
 │
 │ 启动 Kernel
 ▼
GPU
 └── 创建大量线程并执行

所以 GPU 一般不是完全独立工作的，而是由 CPU 负责：

• 分配数据；
• 启动任务；
• 管理程序流程；
• 处理输入输出。

GPU 负责：

• 大规模并行计算；
• 矩阵运算；
• 图像处理；
• 神经网络计算。

这种结构称为：

Host：CPU
Device：GPU

---

二、Thread 如何知道自己处理哪个数据

每个线程执行的代码是一样的：

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

区别在于，每个线程计算出的 i 不一样。

CUDA 为每个线程提供了几个内置变量：

threadIdx.x
blockIdx.x
blockDim.x
gridDim.x

它们的含义如下。

变量	含义
threadIdx.x	当前线程在 Block 内的编号
blockIdx.x	当前 Block 在 Grid 内的编号
blockDim.x	每个 Block 中的线程数
gridDim.x	Grid 中的 Block 数

全局线程编号通常这样计算：

$$i=blockIdx.x\times blockDim.x+threadIdx.x$$

例如每个 Block 有 256 个线程。

对于 Block 0：

blockIdx.x = 0
线程编号为 0～255

所以：

i = 0 × 256 + threadIdx.x

得到：

0～255

对于 Block 1：

blockIdx.x = 1

所以：

i = 1 × 256 + threadIdx.x

得到：

256～511

完整对应关系如下：

图示：二、Thread 如何知道自己处理哪个数据

注意：

每个 Block 中的 threadIdx.x 都会从 0 重新开始。

因此，必须结合 blockIdx.x 才能得到全局唯一编号。

---

三、为什么创建了 1024 个线程，却只处理1000个数据

前面启动了：

$$4\times256=1024$$

个线程，但数组长度只有 1000。

因此最后 24 个线程不能访问数组，否则会越界。

所以 Kernel 中通常会写：

if (i < N) {
    C[i] = A[i] + B[i];
}

线程 0～999 执行计算，线程 1000～1023 什么也不做。

这种写法很常见，因为 Block 的线程数通常取整为：

• 128；
• 256；
• 512。

Block 数量则按下面的公式计算：

$$N_{\text{blocks}}= \left\lceil \frac{N}{T_{\text{block}}} \right\rceil$$

对应代码：

int threads_per_block = 256;
int blocks =
    (N + threads_per_block - 1) / threads_per_block;

---

四、Grid、Block、Thread 是软件概念

程序员看到的层次是：

图示：四、Grid、Block、Thread 是软件概念

这里需要特别注意：

• Grid 是一次 Kernel 启动产生的全部线程；
• Block 是线程的分组；
• Thread 是程序员看到的最小执行单位。

但是这些都是编程模型中的概念。

GPU 真正的硬件结构是：

GPU
 ├── SM 0
 ├── SM 1
 ├── SM 2
 └── SM 3

因此必须建立软件到硬件的映射。

---

五、Block 如何分配给 SM

假设 GPU 有 4 个 SM，Kernel 创建了 10 个 Block。

开始时，GPU 可能这样分配：

SM 0 ← Block 0
SM 1 ← Block 1
SM 2 ← Block 2
SM 3 ← Block 3

当某个 Block 完成后，新的 Block 会继续进入：

SM 0：Block 0 完成 → 执行 Block 4
SM 1：Block 1 完成 → 执行 Block 5

直到全部 Block 执行完成。

实际情况中，一个 SM 通常可以同时容纳多个 Block，例如：

SM 0
 ├── Block 0
 ├── Block 4
 └── Block 8

能同时放多少个 Block，取决于资源是否足够，例如：

• 寄存器；
• 共享内存；
• Warp 数量；
• Block 数量上限；
• 每个 Block 的线程数量。

因此不能简单理解为：

一个 SM 只能执行一个 Block

更准确的说法是：

一个 Block 只能被分配到一个 SM，但一个 SM 可以同时驻留多个 Block。

---

六、为什么 Block 不能跨 SM 执行

同一个 Block 中的线程可以：

• 共享 Shared Memory；
• 通过同步指令相互等待；
• 协同处理同一块数据。

例如：

__shared__ float buffer[256];

这个共享内存只属于当前 Block。

如果一个 Block 的线程被拆到不同 SM，那么：

• 线程无法共享同一块低延迟存储器；
• 同步控制会非常复杂；
• 硬件实现成本会大幅增加。

所以 GPU 规定：

一个 Block 完整地分配给一个 SM

例如：

Block 0 的全部线程 → SM 2

Block 0 不会出现：

一半线程在 SM 1
另一半线程在 SM 2

而且一个 Block 开始在某个 SM 上执行后，通常会一直在这个 SM 上运行到结束。

---

七、Block 进入 SM 后，会被拆成 Warp

假设一个 Block 有 256 个线程。

NVIDIA GPU 中，一个 Warp 有 32 个线程，因此：

$$256\div32=8$$

这个 Block 会被拆成 8 个 Warp：

图示：七、Block 进入 SM 后，会被拆成 Warp

因此：

• Thread 是程序员定义的；
• Warp 是硬件自动形成的；
• 程序员一般不需要手动创建 Warp。

Warp 数量为：

$$N_{\text{warps}}= \left\lceil \frac{T_{\text{block}}}{32} \right\rceil$$

如果一个 Block 有 100 个线程：

$$\left\lceil\frac{100}{32}\right\rceil=4$$

会形成：

Warp 0：32 个有效线程
Warp 1：32 个有效线程
Warp 2：32 个有效线程
Warp 3：只有 4 个有效线程

最后一个 Warp 中其余 28 个位置空闲，但硬件仍然需要以一个完整 Warp 的形式调度。

因此 Block 的线程数量通常设为 32 的整数倍，例如：

128、256、512

---

八、Warp 才是硬件真正调度的单位

虽然程序员看到的是 Thread，但 Warp Scheduler 调度的是 Warp。

假设 SM 中有这些 Warp：

Warp 0：准备执行
Warp 1：等待显存
Warp 2：准备执行
Warp 3：等待前一条指令完成
Warp 4：准备执行

调度器会选择一个已经准备好的 Warp：

周期 1：选择 Warp 0
周期 2：选择 Warp 2
周期 3：选择 Warp 4

每次选择一个 Warp，并向执行单元发射指令。

因此不能理解为：

GPU 调度 Thread 7
GPU 调度 Thread 18

更准确的说法是：

GPU 调度包含 Thread 0～31 的 Warp 0

---

九、一个 Warp 怎样执行指令

假设 Warp 中有 32 个线程，代码为：

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

从程序角度看，每个线程独立执行：

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 中的线程共同执行相同指令：

指令 1：加载 A[i]
指令 2：加载 B[i]
指令 3：执行加法
指令 4：写回 C[i]

在执行加法时：

所有活跃线程：执行加法
每个线程：使用自己的寄存器和数据

这就是 SIMT：

$$\text{Single Instruction, Multiple Threads}$$

单条指令同时作用于多个线程。

---

十、Thread 是独立的，但又不是完全独立的

从编程模型看，每个 Thread 都有：

• 自己的线程编号；
• 自己的寄存器；
• 自己的程序状态；
• 自己的数据；
• 自己的逻辑执行路径。

因此看起来每个线程像一个独立的小程序。

但硬件执行时，同一个 Warp 中的线程共享：

• 指令发射；
• 程序计数器控制；
• 调度机会；
• 部分执行资源。

因此 Warp 中线程的独立性是有限的。

这会导致一个重要问题：分支发散。

---

十一、什么是分支发散

考虑下面的代码：

if (threadIdx.x % 2 == 0) {
    C[i] = A[i] + B[i];
} else {
    C[i] = A[i] - B[i];
}

Warp 中：

Thread 0：执行加法
Thread 1：执行减法
Thread 2：执行加法
Thread 3：执行减法
...

但是 Warp 通常只能在某一时刻发射一条共同指令，因此无法让一半线程做加法，另一半线程同时做减法。

硬件通常会分两次执行。

第一次：

执行 if 路径
偶数线程：激活
奇数线程：关闭

第二次：

执行 else 路径
偶数线程：关闭
奇数线程：激活

示意如下：

第一次执行加法：
Thread 0  活跃
Thread 1  屏蔽
Thread 2  活跃
Thread 3  屏蔽

第二次执行减法：
Thread 0  屏蔽
Thread 1  活跃
Thread 2  屏蔽
Thread 3  活跃

这叫做：

Warp Divergence，Warp 分支发散。

原本可以一次完成的工作，现在需要分两次执行，因此性能下降。

---

十二、并不是所有 if 都会导致严重分支发散

例如：

if (blockIdx.x == 0) {
    ...
} else {
    ...
}

如果整个 Warp 内的线程条件相同，那么 Warp 不会分裂。

例如 Warp 0 中所有线程都满足：

blockIdx.x == 0

则全部线程走同一个分支。

真正影响性能的是：

同一个 Warp 内的线程走不同分支。

所以需要区分：

不同 Block 走不同分支：通常没关系
不同 Warp 走不同分支：通常没关系
同一 Warp 内线程走不同分支：可能降低性能

---

十三、SM 中大致有什么

暂时可以把一个 SM 理解为：

图示：十三、SM 中大致有什么

各部分作用如下。

Warp Scheduler

选择下一个可以执行的 Warp。

Dispatch Unit

把指令发送给对应执行单元。

CUDA Core

主要执行普通浮点运算，例如：

FP32 加法
FP32 乘法
FMA

Integer Unit

执行整数运算，例如：

整数加法
地址计算
位运算

Load/Store Unit

负责：

加载数据
存储数据
地址访问

Special Function Unit

执行一些特殊运算，例如：

sin
cos
exp
倒数
平方根近似

Tensor Core

执行矩阵乘加运算，主要用于：

• 深度学习；
• 矩阵计算；
• Transformer；
• 卷积。

Register File

保存每个线程的局部数据，是速度最快的存储资源之一。

Shared Memory

同一个 Block 中的线程共享的片上存储器。

---

十四、一个 Kernel 的完整执行流程

可以总结为以下过程。

第一步：CPU 准备数据

CPU 内存
A、B、C

然后把数据复制到 GPU 显存。

CPU Memory
     │
     ▼
GPU Global Memory

第二步：CPU 启动 Kernel

vector_add<<<4, 256>>>(A, B, C, N);

GPU 创建：

1 个 Grid
4 个 Block
1024 个 Thread

第三步：GPU 把 Block 分配给 SM

Block 0 → SM 0
Block 1 → SM 1
Block 2 → SM 2
Block 3 → SM 3

第四步：Block 被拆成 Warp

每个 Block 有 256 个线程：

256 Thread → 8 Warp

第五步：Warp Scheduler 选择 Warp

Warp 0 准备好 → 发射
Warp 1 等待内存 → 暂停
Warp 2 准备好 → 发射

第六步：执行指令

加载 A
加载 B
执行加法
写回 C

第七步：所有 Thread 完成

一个 Block 中的全部线程完成后，这个 Block 结束。

第八步：全部 Block 完成

Grid 执行结束，Kernel 完成。

CPU 可以把结果从 GPU 显存复制回来。

---

十五、必须区分的几个概念

概念	类型	作用
Kernel	软件	在 GPU 上运行的函数
Grid	软件	一次 Kernel 的所有 Block
Block	软件	可以协作的一组线程
Thread	软件	程序员看到的执行单元
Warp	硬件调度概念	32 个线程组成的调度单位
SM	硬件	执行 Warp 的处理器模块
CUDA Core	硬件	执行普通算术运算的单元

它们之间的关系可以画成：

图示：十五、必须区分的几个概念

---

十六、一个常见误区

错误理解：

一个 Thread 对应一个 CUDA Core

这通常不成立。

线程数量可能有几万个甚至几十万个，但 GPU 不可能拥有同样数量的物理运算单元。

更准确的理解是：

大量 Thread
    ↓
组成大量 Warp
    ↓
Warp 被分时调度
    ↓
复用有限数量的执行单元

例如一个 SM 中可能驻留上千个线程，但只有较少数量的物理执行单元。

这些线程不是全部同时在 ALU 上运行，而是由 Warp Scheduler 分批发射。

---

十七、为什么需要大量 Warp

假设 Warp 0 执行一条显存加载指令：

x = A[i];

数据可能需要较长时间才能返回。

Warp 0 进入等待状态：

Warp 0：等待数据

调度器立即执行其他 Warp：

Warp 1：执行加法
Warp 2：执行乘法
Warp 3：访问共享内存

等 Warp 0 的数据返回后，再继续执行 Warp 0。

这就是 GPU 隐藏延迟的核心：

不是让一次访存变得特别快
而是让访存等待时间被其他线程的计算覆盖

因此 GPU 希望一个 SM 内同时有足够多的 Warp。

---

十八、Occupancy 的初步概念

Occupancy 通常表示：

$$\text{Occupancy}= \frac{W_{\text{resident}}} {W_{\text{max}}}$$

例如一个 SM 最多支持 64 个 Warp，当前驻留 48 个 Warp：

$$Occupancy=\frac{48}{64}=75\%$$

较高的 Occupancy 通常意味着有更多 Warp 可以用于隐藏延迟。

但 Occupancy 并不是越高性能就一定越好，因为性能还受以下因素影响：

• 显存带宽；
• 缓存命中率；
• 指令依赖；
• 分支发散；
• 算力利用率；
• 访存是否合并。

目前只需要记住：

一个 SM 中同时驻留的线程和 Warp 越多，通常越容易隐藏延迟，但资源使用过多也可能降低并发度。

---

本节核心结论

软件层次：

Kernel → Grid → Block → Thread

硬件执行层次：

GPU → SM → Warp → Thread

关键对应关系：

Block 分配给 SM
Block 被拆成 Warp
Warp 由调度器发射
线程在执行单元上完成计算

还要记住：

1. 一个 Block 只能属于一个 SM。
2. 一个 SM 可以同时驻留多个 Block。
3. 一个 NVIDIA Warp 通常包含 32 个 Thread。
4. Warp 是硬件调度单位，而不是单个 Thread。
5. 大量 Warp 可以帮助隐藏访存延迟。
6. 同一 Warp 内线程走不同分支，会产生分支发散。
7. Thread 数量远大于物理 CUDA Core 数量，线程会分批复用执行单元。

下一部分最适合学习 GPU 的存储层次：

寄存器
Shared Memory
L1 Cache
L2 Cache
Global Memory
Constant Memory
Texture Memory

理解存储层次后，才能真正理解 GPU 为什么有时很快、有时却很慢。