CUDA线程分歧和分支,例子

Question

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

第一个例子:
我有以下内核,我从1个块中的2个线程开始.用[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)同时读取[0]和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参见图3,(a)对角线块

Answer 1

CUDA中的warp中的分支执行顺序没有优先规则 - 行为未定义.编译器,汇编器和JIT运行时可以根据需要自由重新排序指令,并且绝对不能尝试依赖于经验推断的任何顺序,因为它可以更改(如您所知).在这种情况下强制形式正确性的唯一方法是使用原子内存访问操作,这将强制序列化.更好的是,寻找另一种算法.

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