cuda - CUDA中分支的概念(采取、不采取、发散)

Question

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

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

在我看来，当 tid %31 == 31 在扭曲中时，就会发生分歧，但什么是“未采取”？

Answer 1

来自 Nsight Visual Studio 版用户指南:

Not Taken / Taken Total: number of executed branch instructions with a uniform control flow decision; that is all active threads of a warp either take or not take the branch.

Diverged: Total number of executed branch instruction for which the conditional resulted in different outcomes across the threads of the warp. All code paths with at least one participating thread get executed sequentially. Lower numbers are better, however, check the Flow Control Efficiency to understand the impact of control flow on the device utilization.

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

#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 语句。令人惊讶的结果是您没有分歧。这可以通过查看内核函数的反汇编代码(针对 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 和两个 BRA)。两个 warp 都有 1 taken 分支，因为所有线程一致地碰到 EXIT 指令。第一个经纱有 2 未采用 分支，因为两个 BRA 路径在经 yarn 程中并未统一遵循。第二个 warp 有 1 not take 分支和 1 taken 分支，因为所有 warp 线程都统一遵循其中一个分支两个BRA。我想说，diverged* 再次等于零，因为两个分支中的指令完全相同，尽管在​​不同的操作数上执行。