避免翘曲发散

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 所建议的，我相信对您的代码进行相当简单的转换可能会减少您对 warp divergence 的担忧：

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 的负载（创建已经提到的 "conditional load"）。这种预测最小化了负载本身的发散成本。此转换下的其余代码是非发散的（除了 tid >= N，这应该无关紧要）。

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

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

避免翘曲发散

Avoid warp divergence

c

cuda

gpu-warp