避免翘曲发散

Question

我有布尔一维数组T[N] 控制移位值，如下所示：

**a: 指向全局内存中n*n 矩阵的指针数组 我想为每个矩阵a 构造一个 shift*Identity 来获得：

a=a-shift*eye(n)

我有：

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

shift的值由T控制 如果 T[i]==true => shift=shift1 否则 shift=shift2;

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();

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

Answer 1

正如@AnastasiyaAsadullayeva 所建议的那样，我相信对您的代码进行相当简单的转换可能会减少您对扭曲发散的担忧：

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 除外，这应该是无关紧要的）。

同样，如前所述，编译器可能已经观察到并完成了整个转换。这是可能的，但如果不运行您未提供的实际完整测试用例，则无法确认。

更好的方法是以您认为自然的方式编写代码，然后让编译器处理它。此时，您可以使用分析器和分析驱动的优化来确定翘曲发散是否实际上是您代码中的性能问题（分析器有指标和其他方法来评估翘曲发散并在您的代码中指示其严重性。）