3、核函数执行模型：线程块与线程束、指令流水线、寄存器压力来源
好，咱们进入第三章。这一章我打算聊聊GPU核函数到底是怎么跑起来的。说白了，你写的那段CUDA代码，到了硬件层面是怎么被拆解、调度、执行的。很多同学写核函数时只关注算法逻辑，忽略了底层的执行模型，结果性能惨不忍睹。我自己就踩过不少坑，今天把这些经验掰开揉碎了讲给你听。
3.1 线程块与线程束：GPU的“组织单元”
先说说线程块（Thread Block）和线程束（Warp）。这两个概念是理解GPU执行模型的基础。
线程块是你编程时直接面对的“工作组”。你把一堆线程打包成一个块，块内的线程可以共享显存、同步执行。但硬件并不直接执行线程块——它把线程块进一步拆成更小的单元，也就是线程束。
线程束是啥？说白了，就是32个线程组成的一个“执行小队”。GPU的SM（流多处理器）每次取指令，都是针对一个线程束来取的。也就是说，32个线程在同一时刻执行同一条指令，只是处理的数据不同。这就是SIMT（单指令多线程）的精髓。
关键点：线程束是硬件调度的基本单位。你写的线程块大小最好是32的整数倍，否则会有“浪费”的线程槽位。我见过有人设线程块大小为128，性能还行；设成100，性能直接掉一截——因为最后4个线程凑不满一个线程束，硬件还得等。
我记得有一次优化一个图像处理核函数，线程块大小设成了96。结果性能比128差了不少。后来一查，96不是32的整数倍，导致线程束内部分线程闲置。改成128后，吞吐量直接提升了15%。嗯，这个教训挺深刻的。
3.2 指令流水线：隐藏延迟的艺术
接下来聊聊指令流水线。GPU的指令执行不是“取指→译码→执行→写回”这么简单，它背后有一套复杂的流水线机制，目的只有一个：隐藏延迟。
GPU的延迟主要来自两方面：
- 访存延迟：从全局显存读数据，动辄几百个周期
- 计算延迟：浮点运算、整数运算等，几个到几十个周期
怎么隐藏？靠的是线程束切换。当一个线程束在等数据时，SM立刻切换到另一个就绪的线程束去执行。只要你有足够多的线程束在“排队”，SM就能一直忙，不会闲着。
你想想看，这就像餐厅里的厨师。一个订单在等食材，厨师不会干等着，而是立刻做下一个订单。只要订单够多，厨师就能一直炒菜。GPU也是这个道理。
我的建议：每个SM上活跃的线程束数量至少要有4-6个，才能有效隐藏访存延迟。如果线程束太少，SM就会频繁“空转”，性能上不去。我一般会保证每个SM至少有1024个线程（32个线程束）在跑。
指令流水线还有一个细节：指令级并行。现代GPU可以在一个周期内发射多条独立的指令。如果你的代码中连续几条指令没有数据依赖，它们可以并行执行。这需要编译器帮你做指令调度，但你的代码结构也会影响调度效果。
举个例子：
// 不好的写法：连续依赖
float a = b + c;
float d = a + e; // 依赖上一条结果
float f = d + g; // 又依赖上一条
// 好的写法：独立计算
float a = b + c;
float x = y + z; // 与上一条无关
float m = n + p; // 与前面都无关
第二种写法，编译器可以把三条加法指令并行发射，流水线利用率更高。我刚开始写核函数时没注意这个，后来用nvprof一看，指令吞吐量只有理论值的60%。调整代码结构后，提升到了85%。
3.3 寄存器压力来源：性能的隐形杀手
最后聊聊寄存器压力。这是很多同学容易忽略的点，但往往是性能瓶颈的根源。
GPU的寄存器是每个SM私有的，数量有限。比如NVIDIA的A100，每个SM有65536个32位寄存器。这些寄存器要分配给所有活跃的线程束。每个线程能分到多少，取决于你核函数里用了多少局部变量。
寄存器压力，说白了就是你的核函数用了太多寄存器，导致每个SM能同时容纳的线程束数量减少。线程束少了，隐藏延迟的能力就下降了，性能自然就差了。
我曾经踩过的坑：有一次写一个复杂的物理模拟核函数，里面用了大量临时变量。编译后一看，每个线程用了64个寄存器。结果每个SM只能放1024个线程（64个寄存器/线程 × 1024线程 = 65536寄存器），线程束数量直接减半。性能比预期低了30%。
寄存器压力的来源主要有三个：
- 局部变量过多：每个float、int变量都占用一个寄存器。数组、结构体也会占用多个。
- 编译器优化不足：有时候编译器为了简化代码，会把中间结果都存到寄存器里，而不是复用。
- 循环展开过度：手动展开循环会引入大量临时变量，寄存器压力飙升。
怎么缓解？我有几个经验：
- 减少局部变量：能用表达式的地方就别用临时变量。比如
float r = a + b; c = r * d;可以写成c = (a + b) * d;，省一个寄存器。 - 使用__launch_bounds__：这个编译器提示可以告诉编译器你期望的线程块大小，编译器会据此调整寄存器分配策略。
- 检查编译报告：用
--ptxas-options=-v编译，看看每个线程用了多少寄存器。如果超过32，就要警惕了。
我记得有一次优化一个矩阵乘法核函数，每个线程用了48个寄存器。我通过合并临时变量、减少循环展开，把寄存器数降到了24个。结果每个SM的线程束数量翻倍，性能提升了40%。嗯，这个优化效果非常明显。
一个小技巧：如果你发现寄存器压力大，可以试试用共享显存（shared memory）来存储一些中间结果。虽然共享显存访问比寄存器慢，但能释放寄存器，让更多线程束并行执行。有时候整体性能反而更好。
好了，这一章的内容就这些。线程块和线程束是GPU的“组织单元”，指令流水线是隐藏延迟的关键，寄存器压力则是你优化时最需要关注的指标。下一章我会讲寄存器分配的具体策略，包括怎么用编译器选项、怎么手动调整代码来降低寄存器压力。到时候见。