Cuda 原子和条件分支答案

【问题标题】：Cuda atomics and conditional branchesCuda 原子和条件分支
【发布时间】：2017-08-30 08:52:48
【问题描述】：

我正在尝试编写serial 代码的CUDA 版本，作为在分子动力学算法中实现周期性边界条件的一部分。这个想法是，需要使用两个ways 之一将一小部分具有开箱即用位置的粒子放回原处，并限制我使用第一种方式的次数。

基本上，它归结为以下 MWE。我有一个数组x[N]，其中N 很大，下面是serial 代码。

#include <cstdlib>

int main()
{
  int N =30000;
  double x[30000];
  int Nmax = 10, count = 0;

  for(int i = 0; i < N; i++)
    x[i] = 1.0*(rand()%3);

  for(int i = 0; i < N; i++)
   {
      if(x[i] > 2.9)
        {
          if(count < Nmax)
            {
              x[i] += 0.1; //first way
              count++;
            }
          else
            x[i] -= 0.2; //second way
        }
    }
}

请假设 x[i] > 2.9 仅用于 x[i] 的 30000 个元素中的一小部分（大约 12-15 个）。

请注意，i 的顺序并不重要，即，不必让10 最低的i 使用x[i] += 0.1，从而使算法具有潜在的可并行性。我想到了 MWE 的以下 CUDA 版本，它与 nvcc -arch sm_35 main.cu 编译，其中 main.cu 读作

#include <cstdlib>

__global__ void PeriodicCondition(double *x, int *N, int *Nmax, int *count)
{
  int i = threadIdx.x+blockIdx.x*blockDim.x;
  if(i < N[0])
    {
      if(x[i] > 2.9)
        {
           if(count[0] < Nmax[0]) //===============(line a)
             {
               x[i] += 0.1; //first way
               atomicAdd(&count[0],1); //========(line b)
             }
           else
             x[i] -= 0.2; //second way
        }
    }
}

int main()
{
  int N = 30000;
  double x[30000];
  int Nmax = 10, count = 0;

  srand(128512);
  for(int i = 0; i < N; i++)
    x[i] = 1.0*(rand()%3);

  double *xD;
  cudaMalloc( (void**) &xD, N*sizeof(double) );
  cudaMemcpy( xD, &x, N*sizeof(double),cudaMemcpyHostToDevice );

  int *countD;
  cudaMalloc( (void**) &countD, sizeof(int) );
  cudaMemcpy( countD, &count, sizeof(int),cudaMemcpyHostToDevice );

  int *ND;
  cudaMalloc( (void**) &ND, sizeof(int) );
  cudaMemcpy( ND, &N, sizeof(int),cudaMemcpyHostToDevice );

  int *NmaxD;
  cudaMalloc( (void**) &NmaxD, sizeof(int) );
  cudaMemcpy( NmaxD, &Nmax, sizeof(int),cudaMemcpyHostToDevice );

  PeriodicCondition<<<938,32>>>(xD, ND, NmaxD, countD);

  cudaFree(NmaxD);
  cudaFree(ND);
  cudaFree(countD);
  cudaFree(xD);

}

当然，这是不正确的，因为(line a) 上的if 条件使用了在(line b) 中更新的变量，该变量可能不是最新的。这有点类似于Cuda atomics change flag，但是，我不确定使用临界区是否以及如何提供帮助。

当每个线程检查(line a) 上的if 条件时，有没有办法确保count[0] 是最新的，而不会使代码过于串行？

【问题讨论】：

我认为@Claude 的答案是一个很好的答案，如果您可以容忍count[0] 始终递增，那么它会非常简单和整洁。要重现您的序列代码的确切行为（count 递增直到达到Nmax，然后停止），应该可以使用围绕atomicCAS 构建的自定义原子。然而，我怀疑在性能方面它会比 Claude 的答案更昂贵。

标签： c++ algorithm cuda atomic compare-and-swap

【解决方案1】：

每次只增加原子计数器，并在您的测试中使用它的return value：

...
  if(x[i] > 2.9)
    {
       int oldCount = atomicAdd(&count[0],1);
       if(oldCount < Nmax[0])
         x[i] += 0.1; //first way
       else
         x[i] -= 0.2; //second way
    }
...

如果您说大约 15 个项目超过 2.9 并且 Nmax 大约为 10，则会有少量“额外”原子操作，其开销可能很小（我看不出如何做更多有效，这并不是说不可能......）。

【讨论】：