CUDA线程分歧和分支，示例

如何解决CUDA线程分歧和分支，示例

| 我有一些例子让我有些头疼： 我产生了线程分歧，但是我无法弄清楚首先计算哪个分支或哪个语句？ 第一个例子： 我有以下内核，我从1个块中的2个线程开始。 其中a [0] = 0，而a1 = 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 

我假设因为两个线程处于一个扭曲中，所以它们以锁步执行，并且（a）和（b）都同时读取a [0]和a1。 第二个例子： 与第一个完全相同，但是现在删除了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）参见图3，（a）对角线块     

解决方法

CUDA中的扭曲内没有分支执行顺序的优先规则-行为未定义。编译器，汇编器和JIT运行时可以自由地按自己的意愿对指令进行重新排序，并且绝对不能尝试依靠凭经验推断出的任何顺序，因为它可以更改（如您所知）。在这种情况下强制形式正确的唯一方法是使用原子内存访问操作，这将强制进行序列化。更好的是，寻找另一种算法。 在高斯-塞德尔案例中，正统方法是对矩阵或计算网格的图形分解中的每种颜色使用单独的内核启动。