一、执行配置优化：Grid/Block大小选择、Occupancy计算与调优、Warp发散问题
好，咱们直接进入正题。核函数写好了，性能却上不去？八成是执行配置没调对。我见过太多人，花了两周写核函数，结果Grid/Block大小随便填了个256，就跑来说CUDA不行。嗯，其实不是CUDA不行，是你没伺候好它。
1.1 Grid/Block大小选择：别拍脑袋
先问个问题：你平时怎么选Block大小？是不是128、256、512轮着试？我以前也这么干，后来发现这其实是个数学问题。
Block大小不是随便定的。它受两个硬约束：
- 每个Block最多1024个线程（老架构是512或768）
- 每个线程块占用的共享内存和寄存器不能超标
我个人习惯，先算一个理论值。假设每个线程处理一个数据元素，数据总量是N。那么：
// 推荐做法：先定Block大小，再算Grid大小
int blockSize = 256; // 通常是32的倍数，128~512之间
int gridSize = (N + blockSize - 1) / blockSize;
// 启动核函数
myKernel<<<gridSize, blockSize>>>(d_data, N);为什么选256？因为它是32的倍数，而且大多数GPU上，256线程的Block能很好地平衡占用率和资源消耗。我在项目中遇到过，用128时SM利用率不够，用512时寄存器压力太大，256刚刚好。
1.2 Occupancy计算与调优：别被数字骗了
Occupancy，说白了就是SM上活跃的Warp数量占理论最大值的比例。很多人追求100% Occupancy，其实没必要。我见过一个案例，Occupancy从100%降到67%，性能反而提升了30%。为什么？因为每个线程能用的寄存器多了，减少了全局内存访问。
Occupancy的计算公式：
Occupancy = 实际活跃Warp数 / 最大支持Warp数但真正影响性能的，是每个SM上同时运行的Warp数。这个值受四个因素限制：
- 每个SM的最大线程数（比如1024）
- 每个SM的最大Block数（比如32）
- 每个线程的寄存器用量
- 每个Block的共享内存用量
我建议你用NVIDIA的cudaOccupancyMaxPotentialBlockSize函数来自动计算：
int minGridSize, blockSize;
cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize,
myKernel, 0, 0);
// blockSize就是理论最优值
// minGridSize是达到最大Occupancy的最小Grid大小我曾经调一个矩阵乘法核函数，Occupancy从100%降到75%，每个线程的寄存器从32个增加到48个，结果全局内存访问减少了40%，整体性能提升了22%。嗯，这就是典型的“降占用率提性能”。
1.3 Warp发散问题：32条线程的默契
Warp是GPU执行的基本单位，32条线程一起跑。如果它们走不同的分支，就出问题了——这叫Warp发散。
看个例子：
// 糟糕的写法：Warp内分支发散
if (threadIdx.x % 2 == 0) {
// 偶数线程走这里
result = a + b;
} else {
// 奇数线程走这里
result = a * b;
}
// 好的写法：用三元运算符或数学变换
result = (threadIdx.x % 2 == 0) ? (a + b) : (a * b);
// 或者更彻底：让所有线程做同样的操作
result = a + (threadIdx.x % 2 == 0 ? b : (a * (b - 1)));为什么三元运算符好？因为编译器有时能把它优化成无分支代码。但if-else基本没戏。
Warp发散的本质是：32条线程中，只要有一条走了不同分支，其他31条就得等着。你想想看，这多浪费。
怎么避免Warp发散？三个原则：
- 分支条件尽量基于Warp内不变的值，比如blockIdx.x、blockDim.x这些
- 如果必须用数据做分支，先对数据排序，让相似的数据聚在一起
- 能用数学变换就别用分支，比如用min/max代替if-else
举个例子，处理稀疏数据时：
// 先对数据排序，让非零元素集中
// 然后每个Warp处理连续的一段
// 这样Warp内要么全是非零，要么全是零
// 分支只发生在Warp边界上说白了，Warp发散就是让32条线程“步调一致”。你让它们各走各的，性能就崩了。
知识体系总览
下面这张图，是我做优化时脑子里一直挂着的框架：
这三个方面是联动的。你改了Block大小，Occupancy就变了；Occupancy变了，每个线程的寄存器数就变了；寄存器数变了，你可能需要重新组织数据来避免Warp发散。说白了，这是一个系统工程。
最后说一句：别迷信工具给出的最优配置。我习惯先用cudaOccupancyMaxPotentialBlockSize算个起点，然后手动微调。有时候差32个线程，性能就差20%。嗯，这就是GPU优化的魅力——细节决定成败。