CUDA 线程发散和分支，示例

Question

我有一些例子让我有些奇怪的头痛:我产生了线程分歧，但我无法弄清楚首先计算哪个分支或哪些语句？

第一个示例:
我有以下内核，它由 1 个 block 中的 2 个线程启动。其中 a[0]=0，且 1 =0。

__global__ void branchTest_kernel( float* a){

  int tx = threadIdx.x;

  if(tx==0){                   // or tx==1
     a[1] = a[0] + 1;  (a)
  }else if(tx==1){             // or tx==0
     a[0] = a[1] + 1;;         (b)
  }
}

输出

a[0] = 1  
a[1] = 1 

我假设因为两个线程位于一个 warp 中，所以它们以锁步方式执行，并且 (a) 和 (b) 都同时读取 a[0] 和 a 1 .

第二个示例:
与第一个完全相同，但是现在删除了 else if 部分:

__global__ void branchTest_kernel( float* a){

  int tx = threadIdx.x;

  if(tx==0){
     a[1] = a[0] + 1;  (a)
  }else{
     a[0] = a[1] + 1;  (b)
  }


} 

输出

a[0] = 1  
a[1] = 2 

是什么导致这种行为突然现在(b)是第一个，(a)是第二个......(可能是最内部的分支)有人可以解释一下分支的优先规则吗？或者哪里可以找到此类信息？

我在实现 Gauss-Seidel 求解器期间遇到了这个示例: Gauss Seidel See Figure 3, (a) diagonal block

Answer 1

CUDA 中的扭曲内的分支执行顺序没有优先规则 - 行为未定义。编译器、汇编器和 JIT 运行时可以自由地按照它们认为合适的方式重新排序指令，并且您绝对不能尝试并依赖您凭经验推断出的任何顺序，因为它可以更改(正如您所发现的)。在这种情况下强制形式正确性的唯一方法是使用原子内存访问操作，这将强制序列化。更好的是，寻找另一种算法。

在您的 Gauss-Seidel 案例中，正统方法是为矩阵或计算网格的图形分解中的每种颜色使用单独的内核启动。