GPU任务调度与CUDA并发执行实战 - 第2节
一、从宏观到微观：GPU 的“三明治”结构
要理解 GPU 任务调度，得先搞明白它的硬件长什么样。
我个人习惯把 GPU 想象成一个超级工厂。工厂里有多个车间（GPC，图形处理簇），每个车间里有若干条生产线（SM，流多处理器）。每条生产线上有几十个工人（Core，CUDA 核心），他们同时干活，但必须步调一致——这就是 Warp（线程束）。
说白了，SM 是 GPU 最核心的计算单元。你写的 CUDA 程序，最终就是被拆成一个个小块，扔到 SM 上执行的。
核心关系链：
GPU → GPC（图形处理簇） → SM（流多处理器） → Warp（线程束） → Core（CUDA 核心）
二、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 上执行的。
- CPU 端启动 Kernel：你写了一个 CUDA kernel 函数，CPU 调用它，并指定线程网格（Grid）和线程块（Block）的大小。
- 线程块分配到 SM：GPU 的全局调度器把线程块分配给各个 SM。一个 SM 可以同时处理多个线程块，但受限于共享内存和寄存器资源。
- 线程块拆成 Warp：SM 把线程块中的线程分成若干个 Warp，每个 Warp 32 个线程。比如一个 256 线程的线程块，会被分成 8 个 Warp。
- Warp 调度器发射指令：Warp 调度器从就绪的 Warp 中选一个，把指令发射到 Core 上执行。如果某个 Warp 在等数据，调度器就切到另一个 Warp。
- 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 调度的基石。搞懂了它们，后面的内容就好理解了。