4、执行配置优化：Grid与Block大小选择、Occupancy计算与优化、Warp调度与分支发散
好，咱们进入第四章。这一章，说白了就是教你「怎么把核函数喂饱」。
你写了一个完美的核函数，算法对，数据也对。但一跑起来，GPU利用率只有30%。为什么？因为你没告诉GPU「该用多少人、怎么分工」。这就是执行配置的活儿。
我个人习惯，调优的第一步永远是看执行配置。配置不对，后面全是白费。
4.1 Grid与Block大小选择：别拍脑袋定
很多人上来就写 dim3 block(256, 1, 1)，然后 Grid 随便算一下。嗯，这其实是个坑。
Block 大小不是随便选的。它直接影响两个东西：每个 SM 能驻留多少线程，以及线程间怎么协作。
核心原则：
- Block 大小最好是 32 的倍数。因为 warp 是 32 线程一组。不是 32 的倍数，最后一个 warp 就浪费了。
- 不要小于 64，也不要大于 512。太小了 SM 利用率上不去；太大了寄存器压力大，容易 spill。
- Grid 大小要足够大，至少是 SM 数量的几倍。这样才能隐藏延迟。
我在项目中遇到过，有人把 block 设成 128，结果 occupancy 只有 33%。改成 256 后，直接翻倍。为什么？因为 128 让每个 SM 能放的 block 数量受限了。
一个常用的经验公式：
// 对于大多数通用计算
dim3 block(256, 1, 1); // 或者 128, 512
// 对于矩阵乘法等访存密集型
dim3 block(16, 16, 1); // 2D tile
// 对于归约等计算密集型
dim3 block(512, 1, 1); // 更多线程做计算你想想看，如果你的数据量是 10000，用 block=256，grid=40。每个 SM 分到几个 block？如果只有 4 个 SM，每个 SM 分 10 个 block，还行。但如果只有 2 个 SM，每个分 20 个，那调度器就忙不过来了。
我的小技巧：用 cudaOccupancyMaxPotentialBlockSize 这个 API 自动算。它会根据你的核函数资源占用，给出最优的 block 大小。省心。
4.2 Occupancy计算与优化：别被数字骗了
Occupancy，直译就是「占用率」。它表示每个 SM 上活跃的 warp 数量与最大 warp 数量的比值。
很多人以为 occupancy 越高越好。其实不一定。我见过 occupancy 100% 但性能很差的例子。
为什么？因为 occupancy 只是「能放多少」，不是「实际跑多快」。如果每个线程用了太多寄存器，虽然 occupancy 低，但每个线程跑得快，整体反而更快。
注意：Occupancy 是手段，不是目的。目的是隐藏延迟。只要延迟能被隐藏，occupancy 低一点也没关系。
Occupancy 的计算公式：
Occupancy = (活跃 warp 数) / (SM 最大 warp 数)
影响因素：
- 每个线程的寄存器数
- 每个 block 的共享内存大小
- block 大小
- SM 硬件限制（最大线程数、最大 block 数）举个例子，假设一个 SM 最多 64 个 warp（2048 线程），每个线程用 32 个寄存器。SM 总共 65536 个寄存器。那么最多能放 65536/32 = 2048 个线程，刚好 64 个 warp。occupancy 100%。
但如果每个线程用 64 个寄存器，那只能放 1024 个线程，occupancy 50%。
我曾经优化一个图像处理核函数，一开始 occupancy 只有 25%。我以为是寄存器太多，拼命减少寄存器。结果 occupancy 上去了，性能反而下降了。为什么？因为我把循环展开取消了，指令数变多了。后来我保持 occupancy 在 50%，但优化了指令级并行，性能反而提升了 30%。
Occupancy 优化策略：
- 用
__launch_bounds__限制最大线程数，让编译器少用寄存器。 - 减少共享内存使用，每个 block 少用一点，就能放更多 block。
- 调整 block 大小，找到 occupancy 和指令级并行的平衡点。
用 NVIDIA 的 occupancy calculator 工具，或者直接在代码里用 API 查：
int minGridSize, blockSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, myKernel, 0, 0);
// blockSize 就是推荐值4.3 Warp调度与分支发散：别让线程吵架
Warp 是 GPU 执行的基本单位。32 个线程一组，一起执行同一条指令。
问题来了：如果这 32 个线程走不同的分支呢？比如 if-else，一半走 if，一半走 else。
这就是分支发散。GPU 的做法是：先执行 if 分支，把走 else 的线程屏蔽掉；再执行 else 分支，把走 if 的线程屏蔽掉。相当于串行化了。
分支发散的影响：
- 同一个 warp 内，分支越多，性能越差。
- 不同 warp 之间，分支不影响。因为 warp 是独立调度的。
- 最坏情况：32 个线程走 32 个不同分支，性能降到 1/32。
我遇到过最典型的例子：一个粒子模拟程序，每个粒子根据位置做不同处理。一开始用 if-else 判断区域，结果性能惨不忍睹。后来我改成数据重排，把相同区域的粒子放到一起，让同一个 warp 内的粒子走相同分支。性能提升了 5 倍。
怎么避免分支发散？
- 用三元运算符代替 if-else。三元运算符会被编译成选择指令，不会发散。
- 数据重排，让同一个 warp 的数据性质相同。
- 用查找表代替复杂分支。
- 把分支提到 warp 级别，用
__any()或__all()判断整个 warp 是否走同一分支。
举个例子：
// 坏：分支发散
if (data[tid] > 0) {
result[tid] = data[tid] * 2;
} else {
result[tid] = data[tid] + 1;
}
// 好：用三元运算符
result[tid] = (data[tid] > 0) ? (data[tid] * 2) : (data[tid] + 1);你想想看，三元运算符在硬件层面就是一个 sel 指令，所有线程同时执行，只是选择不同的结果。没有分支，没有发散。
避坑指南：我曾经在循环里用了 break，结果不同线程在不同迭代退出，导致 warp 内严重发散。后来改成用 continue 加条件判断，虽然多算了几个无用迭代，但整体快了 2 倍。有时候「多做一点」比「发散」更好。
4.4 本章小结：配置调优三板斧
好，咱们捋一下：
- Grid/Block 大小：block 选 128-512，32 的倍数，grid 要够大。
- Occupancy：别盲目追求 100%，找到平衡点。用工具算，别拍脑袋。
- 分支发散：尽量让同一个 warp 走同一条路。用三元运算符、数据重排、查找表。
这三板斧用好了，你的核函数性能至少能翻一倍。我见过太多人在这上面栽跟头，所以这一章我讲得特别细。嗯，下一章咱们聊聊共享内存和 bank conflict，那又是一个大坑。
想要更多资料，联系可可