4、线程与线程束：Warp（线程束）的概念、线程调度机制、分支发散对性能的影响
4.1 线程束到底是什么？
好，咱们直接切入正题。线程束（Warp），是GPU执行的基本单位。
你写代码时，觉得是成千上万个线程在并行跑。但硬件层面，GPU没那么奢侈。它把32个线程捆成一捆，这就是一个Warp。这32个线程，共享同一个指令，但处理不同的数据。
说白了，就是SIMD（单指令多数据）的变种，NVIDIA管它叫SIMT（单指令多线程）。
我个人习惯把Warp想象成一辆公交车。车上32个座位，每个座位坐一个线程。司机（指令调度器）喊一声“向左转”，所有乘客都得向左转。但每个乘客手里拿的地图（数据）可能不一样。
核心要点：Warp是GPU调度和执行的最小单位。不是单个线程，是32个线程的集合。
4.2 线程调度机制：零开销切换
CPU切换线程，开销很大。要保存上下文、恢复寄存器，折腾半天。
GPU不一样。它手里攥着大量的Warp，比如一个SM（流多处理器）里可能同时有64个Warp在“待命”。
调度器的工作很简单：
- 当前Warp在等数据？好，换下一个。
- 下一个Warp指令准备好了？直接发射。
切换Warp，不需要保存任何东西。为什么？因为每个Warp都有自己的寄存器文件，物理上就是独立的。调度器只需要换一个指针，指向下一个Warp的指令地址就行。
我在项目中遇到过一个问题：有同事觉得线程数越多越好，结果把每个线程的寄存器用量撑爆了。导致SM里能同时容纳的Warp数量骤降，调度器没得切换了，性能直接腰斩。嗯，这里要注意，寄存器是稀缺资源，别乱用。
避坑指南：我曾经在优化一个矩阵乘法的kernel时，发现占用率（occupancy）上不去。查了半天，发现是每个线程用了太多局部变量，编译器被迫用了local memory（局部内存），速度慢得离谱。后来我强制限制了每个线程的寄存器数量（用__launch_bounds__），虽然单个线程慢了，但整体吞吐量翻了一倍。
4.3 分支发散：性能杀手
现在聊聊最让人头疼的问题——分支发散。
你想想看，一个Warp里的32个线程，执行同一条指令。但如果代码里有if-else呢？
if (threadId % 2 == 0) {
// 偶数线程走这里
result = a + b;
} else {
// 奇数线程走这里
result = a * b;
}硬件怎么处理？
它会先让所有线程执行if分支，那些不满足条件的线程，就“静默”掉（相当于NOP操作）。然后再让所有线程执行else分支，这次轮到满足条件的线程静默。
结果就是：一个Warp本来一个周期能干完的活，现在要两个周期。如果分支嵌套三层，那就得四个周期。你想想看，这效率损失有多大。
警告：分支发散不是所有情况都致命。关键是看发散发生在哪个粒度。
- Warp内发散：同一个Warp里的线程走不同分支。这是最糟糕的，性能直接打折扣。
- Warp间发散：不同Warp走不同分支。这完全没问题，因为不同Warp本来就是独立调度的。
4.4 如何避免分支发散？
我总结了几个实战技巧：
- 数据预处理：把数据按条件排序。比如所有偶数线程处理的数据放一起，奇数线程的放一起。这样同一个Warp里的线程大概率走同一个分支。
- 用三元运算符代替if-else：有些情况下，编译器能把三元运算符优化成无分支的指令（比如
sel指令）。 - 重构算法：有时候换个思路，就能避免分支。比如用查表法代替条件判断。
我记得有一次优化一个粒子系统的kernel。每个粒子要根据它的类型执行不同的物理计算。一开始用switch-case，性能惨不忍睹。后来我把粒子按类型排序，然后每个Warp只处理一种类型的粒子。分支发散消失了，性能提升了3倍多。
4.5 总结一下
| 概念 | 说明 | 性能影响 |
|---|---|---|
| Warp | 32个线程的集合，共享指令 | 调度基本单位 |
| 零开销切换 | Warp间切换不保存上下文 | 隐藏内存延迟的关键 |
| 分支发散 | Warp内线程走不同分支 | 导致部分线程静默，降低吞吐 |
| 避免发散 | 数据排序、算法重构 | 可提升数倍性能 |
最后说一句：理解Warp，是理解GPU性能的起点。你写kernel时，脑子里要时刻想着“我这个Warp里的32个线程，它们的行为一致吗？”如果不一致，那就得小心了。
小技巧：NVIDIA的Nsight Compute工具里有个“Branch Efficiency”指标。如果这个值低于90%，说明你的分支发散比较严重，该优化了。
好了，这一章就到这里。下一章我们会聊聊共享内存和同步机制，那又是另一个有意思的话题。