Cuda 原子和条件分支

Question

我正在尝试编写 CUDA 版本的 serial 代码，作为在分子动力学算法中实现周期性边界条件的一部分。这个想法是，有一小部分位置开箱即用的粒子需要用两个 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，使得算法潜在可并行化。我想到了以下 CUDA 版本的 MWE，它使用 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] 是最新的，而不会使代码过于串行？

Answer 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，则会有少量 "extra" 原子操作，其开销可能是最小的（我看不出如何更有效地做到这一点，这并不是说这是不可能的......）。

Cuda 原子和条件分支

Cuda atomics and conditional branches

c++

algorithm

cuda

atomic

compare-and-swap