CUDA中分支的概念（采取、不采取、发散）

Question

在 Nsight Visual Studio 中，我们将有一个图表来呈现“已采取”、“未采取”和“分歧”分支的统计信息。我对“不采取”和“分歧”之间的区别感到困惑。例如

kernel()
{
  if(tid % 32 != 31)
     {...}
  else
     {...}
}

在我看来，当tid %31 == 31处于扭曲状态时，就会发生分歧，但什么是“不采取”呢？

Answer 1

来自 Nsight Visual Studio 版用户指南：

Not Taken / Taken Total：具有统一控制流决策的已执行分支指令数；也就是说，warp 的所有活动线程要么采用分支，要么不采用分支。

Diverged：已执行的分支指令总数，其中条件导致扭曲线程中的不同结果。至少有一个参与线程的所有代码路径都会按顺序执行。数字越低越好，但是，请检查流量控制效率以了解控制流对设备利用率的影响。

现在，让我们考虑以下简单的代码，这可能是您当前在测试中考虑的内容：

#include<thrust\device_vector.h>

__global__ void test_divergence(int* d_output) {

    int tid = threadIdx.x;

    if(tid % 32 != 31)
        d_output[tid] = tid;
    else
        d_output[tid] = 30000;

}

void main() {

    const int N = 32;

    thrust::device_vector<int> d_vec(N,0);

    test_divergence<<<2,32>>>(thrust::raw_pointer_cast(d_vec.data()));

}

Nsight 生成的分支统计图如下所示。正如您所看到的，Taken等于100%，因为所有线程都会碰到该if语句。令人惊讶的结果是你没有Diverge。这可以通过查看内核函数的反汇编代码（针对 2.1 的计算能力进行编译）来解释：

MOV R1, c[0x1][0x100];
S2R R0, SR_TID.X;                      
SHR R2, R0, 0x1f;                      
IMAD.U32.U32.HI R2, R2, 0x20, R0;      
LOP.AND R2, R2, -0x20;
ISUB R2, R0, R2;
ISETP.EQ.AND P0, PT, R2, 0x1f, PT;
ISCADD R2, R0, c[0x0][0x20], 0x2;
SEL R0, R0, 0x7530, !P0;
ST [R2], R0;
EXIT;

正如您所看到的，编译器能够优化反汇编代码，以便除了指令导致的统一EXIT分支之外不存在分支，正如 Greg Smith 在下面的评论中指出的那样。

编辑：格雷格·史密斯评论之后的更复杂的例子

我现在正在考虑以下更复杂的示例

/**************************/
/* TEST DIVERGENCE KERNEL */
/**************************/
__global__ void testDivergence(float *a, float *b)
{
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    if (tid < 16) a[tid] = tid + 1;
    else b[tid] = tid + 2;
}

/********/
/* MAIN */
/********/
void main() {

    const int N = 64;

    float* d_a; cudaMalloc((void**)&d_a,N*sizeof(float));
    float* d_b; cudaMalloc((void**)&d_b,N*sizeof(float));

    testDivergence<<<2,32>>>(d_a, d_b);

}

这是分支统计图

这是反汇编代码

           MOV R1, c[0x1][0x100];                   
           S2R R0, SR_CTAID.X;                         R0 = blockIdx.x
           S2R R2, SR_TID.X;                           R0 = threadIdx.x
           IMAD R0, R0, c[0x0][0x8], R2;               R0 = threadIdx.x + blockIdx.x * blockDim.x
           ISETP.LT.AND P0, PT, R0, 0x10, PT;          Checks if R0 < 16 and puts the result in predicate register P0
/*0028*/   @P0 BRA.U 0x58;                                 If P0 = true, jumps to line 58
          @!P0 IADD R2, R0, 0x2;                           If P0 = false, R2 = R0 + 2
          @!P0 ISCADD R0, R0, c[0x0][0x24], 0x2;           If P0 = false, calculates address to store b[tid] in global memory
          @!P0 I2F.F32.S32 R2, R2;                                                "
          @!P0 ST [R0], R2;                                                       "
/*0050*/  @!P0 BRA.U 0x78;                                 If P0 = false, jumps to line 78
/*0058*/   @P0 IADD R2, R0, 0x1;                           R2 = R0 + 1
           @P0 ISCADD R0, R0, c[0x0][0x20], 0x2;
           @P0 I2F.F32.S32 R2, R2;
           @P0 ST [R0], R2;
/*0078*/       EXIT;

可以看到，现在BRA反汇编代码中有两条指令。从上图中，每个扭曲都会碰撞到3分支（一个用于EXIT，两个用于BRAs）。两个线程束都已1 采取分支，因为所有线程都一致地碰到EXIT指令。第一个经线没有2 产生分支，因为两个BRAs 路径没有均匀地穿过经线。第二经纱未1 发生分支并1 发生分支，因为所有经纱均一致地遵循两个BRA经纱之一。我想说，diverged* 再次等于零，因为两个分支中的指令完全相同，尽管在不同的操作数上执行。