第4章 GPU调度基础

4、GPU调度基础：任务调度单元（GPC/TPC）、Warp调度机制、Occupancy（占用率）概念

好，咱们今天聊点硬核的。GPU调度，说白了就是「怎么把活儿分下去，让每个计算单元都不闲着」。我刚开始接触GPU底层时，总觉得调度就是个排队问题。后来调优调得多了才发现——这里面的门道，比你想象的要深得多。

先看一张图，这是NVIDIA GPU里最核心的硬件划分逻辑。

这张图我画了很多遍才满意。你看，从GPU到GPC，再到TPC，最后到SM，每一层都有明确的职责。

GPC（图形处理簇）：每个GPC包含光栅引擎和多边形引擎。说白了，它负责把3D图形「拍扁」成2D像素。我做过一个项目，发现GPC数量直接决定了三角形吞吐量——你想想看，GPC越多，能同时处理的三角形就越多。
TPC（纹理处理簇）：每个TPC下面挂着SM和纹理单元。纹理采样、过滤这些活儿，都在这里干。我记得有一次调纹理性能，死活上不去，后来发现是TPC里的纹理缓存命中率太低——嗯，这又是另一个故事了。
SM（流式多处理器）：这才是真正的「干活主力」。每个SM里有多个CUDA核心、共享内存、寄存器文件。Warp调度就在SM里完成。

核心要点：GPC负责图形前端，TPC负责纹理，SM负责计算。这三层配合好了，GPU才能跑满。

好，现在活儿分到了SM里。SM怎么执行？答案是——以Warp为单位。

一个Warp是32个线程。为什么是32？不是16，也不是64？我个人的理解是：NVIDIA在设计时发现，32个线程刚好能填满一个SM里的CUDA核心数，而且调度开销最小。你想想看，如果Warp太大，线程间同步成本高；太小，调度器忙不过来。

Warp调度有几个关键点：

零开销上下文切换：SM里维护了多个Warp的上下文。当一个Warp在等数据时，调度器立刻切到另一个Warp。切换成本几乎为零——因为所有状态都保存在寄存器里。
指令发射：每个时钟周期，调度器从就绪的Warp里选一个，发射一条指令。如果Warp数量够多，调度器总能找到活干。
分支发散：这是个大坑。Warp里的32个线程如果走不同的分支，那性能直接打折扣。我遇到过最夸张的情况——一个Warp里只有1个线程在干活，其他31个都在等。那性能，惨不忍睹。

避坑指南：我曾经在一个图像处理项目里，因为if-else分支太多，导致Warp发散严重。后来我把数据重新排列，让同一Warp里的线程走相同分支，性能提升了3倍。记住——尽量让Warp里的线程「步调一致」。

Occupancy，中文叫「占用率」。说白了，就是SM里同时活跃的Warp数量，除以SM能支持的最大Warp数量。

举个例子：一个SM最多能放64个Warp，你现在放了48个，那Occupancy就是75%。

为什么Occupancy重要？因为高Occupancy意味着调度器有更多选择。当一个Warp卡在内存访问上时，调度器可以立刻切换到另一个Warp。Occupancy低了，调度器就「没米下锅」了。

影响Occupancy的因素有三个：

注意：高Occupancy不等于高性能。有时候为了降低Occupancy，你可以给每个线程分配更多寄存器，反而能减少寄存器溢出，提升性能。这需要权衡——我一般先用NVIDIA的Occupancy Calculator算一下，再微调。

咱们来点实际的。假设一个SM有65536个寄存器，最大支持64个Warp，每个Warp32个线程。你的kernel每个线程用了40个寄存器，线程块大小是256。

算一下：

每个线程块需要的寄存器 = 256 × 40 = 10240
每个SM能放的线程块数 = floor(65536 / 10240) = 6
每个SM能放的Warp数 = 6 × (256 / 32) = 48
Occupancy = 48 / 64 = 75%

嗯，75%还不错。但如果每个线程用64个寄存器呢？

每个线程块需要的寄存器 = 256 × 64 = 16384
每个SM能放的线程块数 = floor(65536 / 16384) = 4
每个SM能放的Warp数 = 4 × 8 = 32
Occupancy = 32 / 64 = 50%

你看，寄存器用量翻倍，Occupancy直接掉到50%。这就是为什么我总说——寄存器是稀缺资源，省着点用。

我的习惯：写kernel时，先用--ptxas-options=-v编译，看看寄存器用量。如果超过32，我会考虑优化代码，或者用__launch_bounds__限制线程块大小。

好了，这一章的内容就这些。GPC/TPC是硬件骨架，Warp调度是执行引擎，Occupancy是性能标尺。这三样东西搞明白了，GPU调度你就入门了。