https://devtalk.nvidia.com/default/topic/463316/branch-divergent-branches/
避免分支之(一)
为了获得最好的性能,就需要避免同一个warp存在不同的执行路径。避免该问题的方法很多,比如这样一个情形,假设有两个分支,分支的决定条件是thread的唯一ID的奇偶性:
- __global__ void mathKernel1(float *c) {
- int tid = blockIdx.x * blockDim.x + threadIdx.x;
- float a, b;
- a = b = 0.0f;
- if (tid % 2 == 0) {
- a = 100.0f;
- } else {
- b = 200.0f;
- }
- c[tid] = a + b;
- }
一种方法是,将条件改为以warp大小为步调,然后取奇偶,如下:
- __global__ void mathKernel2(void) {
- int tid = blockIdx.x * blockDim.x + threadIdx.x;
- float a, b;
- a = b = 0.0f;
- if ((tid / warpSize) % 2 == 0) {
- a = 100.0f;
- } else {
- b = 200.0f;
- }
- c[tid] = a + b;
- }
$ nvprof --metrics branch_efficiency ./simpleDivergence
- $ nvcc -O3 -arch=sm_20 simpleDivergence.cu -o simpleDivergence
Branch Efficiency的定义如下:
我们也可以使用nvprof的inst_per_warp参数来查看每个warp上执行的指令数目的平均值。
- $ nvprof --metrics inst_per_warp ./reduceInteger
输出,原来的是新的kernel的两倍还多,因为原来的有许多不必要的操作也执行了:
- Neighbored Instructions per warp 295.562500
- NeighboredLess Instructions per warp 115.312500