为何两个CUDA归约内核gpuReduce与gpuReduceOpt存在两倍性能差距

你观察到的性能差异主要由分支发散导致，原理如下：
CUDA硬件以32个线程为一个warp调度执行，同一个warp内的所有线程必须同步执行相同的指令序列。如果同一warp内的线程分支判断结果不一致，就会触发分支发散：硬件会串行执行所有分支路径，未命中当前分支的线程会被屏蔽，执行效率直接下降。

我们可以对比两个内核的分支逻辑差异：

• 非优化版本gpuReduce的分支条件为 if((tid%(2*stride)) == 0)
  任意步长下，满足条件的线程都是间隔分布的：比如stride=1时，只有偶数tid的线程进入分支，同一warp内一半线程命中、一半不命中，分支发散直接让该步长的执行效率降为1/2；直到stride超过16之后，满足条件的线程才会按warp粒度连续分布，发散消失。大部分迭代步都存在分支发散开销。
• 优化版本gpuReduceOpt的分支条件为 if(index < blockDim.x)，其中index = 2*stride*tid
  任意步长下，满足条件的线程都是从tid=0开始连续分布的：比如stride=1时，前512个tid满足条件，刚好对应前16个完整warp全命中、后16个warp全不命中，同一warp内所有线程的分支结果完全一致，不存在任何分支发散开销。

另外模运算本身的计算延迟也高于乘法运算，但这部分占比很低，近两倍的性能差几乎全部来自分支发散的额外开销。

如果要进一步优化归约性能，你还可以尝试引入共享内存减少全局内存访问、后半段规约使用warp同步原语去掉__syncthreads()开销等操作。

内容的提问来源于stack exchange，提问作者hash join