0

我有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 控制 if T[i]==true => shift=shift1 else 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();

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

4

1 回答 1

2

正如@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（创建已经提到的“条件负载”）。这种预测将负载本身的发散成本降至最低。此转换下的此代码的其余部分是非发散的（除了 where tid >= N，这应该是无关紧要的）。

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

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

于 2015-09-14T20:11:11.653 回答