c - 避免翘曲发散

我有一个布尔1D数组T[N]控制移位值，如下所示：
**a：指向全局内存中n*n矩阵的指针数组
我想让每个矩阵a子结构一个移位*恒等式来获得：

a=a-shift*eye(n)

我有：

__device__ bool T[N];
__device__ float shift1[N];
__device__ float shift2[N];
__device__ float* a[N];

移位值由T控制
如果T[i]==true=>shift=shift1
否则移位=移位2；

int tid=threadIdx.x;

      if(tid < N){

              if(T[tid]){

               for (int i=0;i<n;i++){

                   a[tid][i*n+i]=a[tid][i*n+i]-shift1[tid];
               }

            }
        else {

          for (int i=0;i<n;i++){

                   a[tid][i*n+i]=a[tid][i*n+i]-shift2[tid];
               }
            }
        }
      __syncthreads();

这将导致扭曲发散并减慢我的代码。有没有技巧可以避免上述循环的扭曲发散？

最佳答案

正如@AnastasiyaAsadullayeva所建议的，我相信对代码进行一个相当简单的转换可以减少您对warp发散的担心：

int tid=threadIdx.x;
  float myshift;
  if (T[tid]) myshift = shift1[tid];
  else myshift = shift2[tid];
  if(tid < N){
           for (int i=0;i<n;i++){

               a[tid][i*n+i]=a[tid][i*n+i]-myshift;
           }

        }
  __syncthreads();

编译器将预测myshift的加载（创建前面提到的“条件加载”）。这种预测使负载本身的发散成本最小化。此转换下的其余代码是非发散的（除了tid >= N，这应该不需要担心）。
同样，如前所述，编译器可能已经观察并完成了整个转换。这是可能的，但是如果不运行一个实际的完整测试用例（您还没有提供）就无法确认。
一个更好的方法是以一种对您来说似乎很自然的方式编写代码，然后让编译器处理它。此时，您可以使用分析器和分析驱动的优化来确定代码中的扭曲散度是否实际是性能问题（分析器有度量和其他方法来评估扭曲散度并在代码中指示其严重性）