从零学会GPU任务调度 - 第2节

2. GPU硬件架构入门

要理解 GPU 任务调度，得先搞明白它的硬件长什么样。

我个人习惯把 GPU 想象成一个超级工厂。工厂里有多个车间（GPC，图形处理簇），每个车间里有若干条生产线（SM，流多处理器）。每条生产线上有几十个工人（Core，CUDA 核心），他们同时干活，但必须步调一致——这就是 Warp（线程束）。

说白了，SM 是 GPU 最核心的计算单元。你写的 CUDA 程序，最终就是被拆成一个个小块，扔到 SM 上执行的。

核心关系链：

GPU → GPC（图形处理簇） → SM（流多处理器） → Warp（线程束） → Core（CUDA 核心）

GPU（图形处理器） GPC 0（图形处理簇） SM 0 Warp 调度器 32 个 Core SM 1 Warp 调度器 32 个 Core SM 2 Warp 调度器 32 个 Core SM 3 Warp 调度器 32 个 Core GPC 1（图形处理簇） SM 4 Warp 调度器 32 个 Core SM 5 Warp 调度器 32 个 Core SM 6 Warp 调度器 32 个 Core SM 7 Warp 调度器 32 个 Core 全局内存（Global Memory） 显存（VRAM） 所有 SM 共享 延迟高（几百周期）

二、SM（流多处理器）：GPU 的“迷你 CPU”

SM 是什么？你可以把它理解成一个“迷你 CPU”，但它和 CPU 完全不同。

CPU 擅长处理复杂、串行的任务，一个核心一次只能处理一个线程。而 SM 呢？它一次可以处理成百上千个线程。为什么？因为它内部有大量的计算单元（Core），还有一套高效的调度机制。

SM 的内部构成

一个典型的 SM 包含以下关键部件：

  • CUDA 核心（Core）：执行算术运算的工人。数量从几十到上百不等。
  • Warp 调度器：负责从就绪的 Warp 中选一个，发射指令到 Core 上执行。
  • 共享内存（Shared Memory）：SM 内部的高速缓存，所有线程都能访问，延迟极低。
  • 寄存器文件（Register File）：每个线程私有的“小本本”，存临时数据。
  • 加载/存储单元（LD/ST）：负责从显存中读写数据。
  • 特殊函数单元（SFU）：处理 sin、cos、sqrt 等复杂运算。

我的经验： 刚开始写 CUDA 时，我总以为 SM 越多越好。后来发现，如果共享内存用不好，SM 再多也白搭。有一次我写了一个矩阵乘法，因为共享内存 bank conflict 导致性能直接腰斩。嗯，从那以后我养成了先分析共享内存访问模式再动手的习惯。

三、Warp（线程束）：32 个线程的“小分队”

Warp 是 GPU 调度中最核心的概念。没有之一。

你想想看，一个 SM 里有几百个 Core，如果每个线程都独立调度，调度器得累死。所以 NVIDIA 想了个办法：把 32 个线程绑在一起，组成一个 Warp。这 32 个线程执行同一条指令，但处理不同的数据——这就是 SIMT（单指令多线程）模型。

Warp 的调度机制

SM 中的 Warp 调度器会维护一个就绪队列。当某个 Warp 因为等待数据（比如从显存读数据）而阻塞时，调度器会立刻切换到另一个就绪的 Warp 去执行。这个过程是硬件自动完成的，零开销。

关键数字：

一个 Warp = 32 个线程。这是硬件决定的，改不了。

为什么是 32？我猜是因为 32 个线程刚好能填满一个 SM 中一半的 Core，方便做双发射（dual-issue）。

Warp 发散（Warp Divergence）

这是新手最容易踩的坑。如果一个 Warp 里的线程走到了不同的分支（比如 if-else），会发生什么？

答案是：所有线程都会执行所有分支，只不过不符合条件的分支会被“屏蔽”掉。说白了，就是串行化了。

// 糟糕的写法：Warp 发散
if (threadIdx.x % 2 == 0) {
    // 偶数线程执行这里
    result = a + b;
} else {
    // 奇数线程执行这里
    result = a * b;
}

// 好的写法：避免 Warp 发散
// 让同一个 Warp 内的线程走同一个分支
int lane = threadIdx.x & 0x1F;  // 取低 5 位，判断在 Warp 内的位置
if (lane < 16) {
    result = a + b;
} else {
    result = a * b;
}

我曾经踩过的坑： 有一次我写了一个粒子模拟程序，每个粒子根据速度决定是否分裂。结果因为 Warp 发散，性能从 200fps 掉到了 30fps。排查了半天才发现是 if-else 导致的。后来我重新组织了数据，把要分裂的粒子和不分裂的分开处理，性能才恢复。

四、Core（CUDA 核心）：最基础的“工人”

Core 就是 CUDA 核心，也叫 SP（流处理器）。它是 GPU 中最基本的计算单元，负责执行整数和浮点数的加减乘除等运算。

但注意：Core 不是 CPU 核心。CPU 核心有复杂的控制逻辑、分支预测、乱序执行等。而 Core 非常“傻”，它只做一件事：执行指令。控制逻辑全在 SM 的 Warp 调度器里。

Core 的规格

GPU 架构 每个 SM 的 Core 数 Core 频率 支持的数据类型
Fermi (2010) 32 ~1.4 GHz FP32, INT32
Kepler (2012) 192 ~1.0 GHz FP32, INT32
Maxwell (2014) 128 ~1.2 GHz FP32, INT32
Pascal (2016) 64 ~1.5 GHz FP32, INT32, FP16
Turing (2018) 64 ~1.7 GHz FP32, INT32, FP16, Tensor
Ampere (2020) 128 ~1.8 GHz FP32, INT32, FP16, BF16, Tensor

一个常见的误解： 很多人以为 Core 越多性能越好。其实不然。比如 Kepler 有 192 个 Core，但实际性能不如 Maxwell 的 128 个 Core。为什么？因为 Kepler 的 Core 频率低，而且调度效率差。所以看 GPU 性能不能只看 Core 数量，还要看架构、频率、内存带宽等。

五、三者如何协同工作？

现在我们把三个概念串起来，看看一个 CUDA 程序是怎么在 GPU 上执行的。

  1. CPU 端启动 Kernel：你写了一个 CUDA kernel 函数，CPU 调用它，并指定线程网格（Grid）和线程块（Block）的大小。
  2. 线程块分配到 SM：GPU 的全局调度器把线程块分配给各个 SM。一个 SM 可以同时处理多个线程块，但受限于共享内存和寄存器资源。
  3. 线程块拆成 Warp：SM 把线程块中的线程分成若干个 Warp，每个 Warp 32 个线程。比如一个 256 线程的线程块，会被分成 8 个 Warp。
  4. Warp 调度器发射指令：Warp 调度器从就绪的 Warp 中选一个，把指令发射到 Core 上执行。如果某个 Warp 在等数据，调度器就切到另一个 Warp。
  5. Core 执行计算：Core 执行指令，结果写回寄存器或共享内存。

一句话总结：

SM 是车间，Warp 是班组，Core 是工人。班组（Warp）里的 32 个工人（Core）同时干一样的活，但处理不同的零件（数据）。

六、避坑指南：新手最容易犯的错

  • 误区一： 以为线程越多越好。实际上，如果线程太多导致寄存器溢出（spill），性能会暴跌。
  • 误区二： 忽略 Warp 发散。同一个 Warp 内的线程尽量走同一个分支。
  • 误区三： 共享内存用成全局内存。共享内存延迟只有几十个周期，全局内存延迟几百个周期，差了一个数量级。
  • 误区四： 不关注 occupancy（占用率）。occupancy 是 SM 上活跃 Warp 数量与最大 Warp 数量的比值。太低的话，调度器没有足够的 Warp 来隐藏延迟。

我曾经犯过的错： 刚学 CUDA 时，我写了一个程序，把所有的数据都放在全局内存里，结果性能还不如 CPU。后来用 Nsight 分析发现，occupancy 只有 12%，大部分时间 SM 都在等数据。加了共享内存后，occupancy 提升到 75%，性能翻了 6 倍。

好了，这一章就到这里。SM、Warp、Core 这三个概念是 GPU 调度的基石。搞懂了它们，后面的内容就好理解了。

← 上一章 📖 返回目录 下一章 →