CUDA之Branch/Divergent branches详解

https://devtalk.nvidia.com/default/topic/463316/branch-divergent-branches/

避免分支之（一）

为了获得最好的性能，就需要避免同一个warp存在不同的执行路径。避免该问题的方法很多，比如这样一个情形，假设有两个分支，分支的决定条件是thread的唯一ID的奇偶性：

[cpp] view plain copy

1. __global__ void mathKernel1(float *c) {  
2.     int tid = blockIdx.x * blockDim.x + threadIdx.x;  
3.     float a, b;  
4.     a = b = 0.0f;  
5.     if (tid % 2 == 0) {  
6.         a = 100.0f;  
7.     } else {  
8.         b = 200.0f;  
9.     }  
10.     c[tid] = a + b;  
11. }                          

一种方法是，将条件改为以warp大小为步调，然后取奇偶，如下：

[cpp] view plain copy

1. __global__ void mathKernel2(void) {  
2.     int tid = blockIdx.x * blockDim.x + threadIdx.x;  
3.     float a, b;  
4.     a = b = 0.0f;  
5.     if ((tid / warpSize) % 2 == 0) {  
6.         a = 100.0f;  
7.     } else {  
8.         b = 200.0f;  
9.     }  
10.     c[tid] = a + b;  
11. } 
    

$ nvprof --metrics branch_efficiency ./simpleDivergence

1. $ nvcc -O3 -arch=sm_20 simpleDivergence.cu -o simpleDivergence

Branch Efficiency的定义如下：

我们也可以使用nvprof的inst_per_warp参数来查看每个warp上执行的指令数目的平均值。

[cpp] view plain copy

1. $ nvprof --metrics inst_per_warp ./reduceInteger  

输出，原来的是新的kernel的两倍还多，因为原来的有许多不必要的操作也执行了：

[cpp] view plain copy

1. Neighbored Instructions per warp 295.562500  
2. NeighboredLess Instructions per warp 115.312500