Cudaアトミックと条件分岐

Question

分子動力学アルゴリズムで周期的な境界条件を実装するために、serialコードのCUDAバージョンを記述しようとしています。考え方は、ボックスの外にある位置の粒子のほんの一部が、2つのうちの1つを使用して戻される必要があるということです。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]の30000個の要素のごく一部（約12〜15）のためにそのx[i] > 2.9を想定してください。 iの順序は重要ではない

注、すなわちアルゴリズムは、潜在的に並列化作り、10最低iはx[i] += 0.1を使用することが必要ではありません。私は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); 

} 

として読み込むnvcc -arch sm_35 main.cuでコンパイルMWE、以下のCUDAバージョンを考え、(line a)上if条件が(line b)で更新される変数を使用しているため、これは正しくありません、これは最新のものではない可能性があります。これはCuda atomics change flagと多少似ていますが、クリティカルセクションを使用するのがどのように役立つかどうか、またどのように役立つかはわかりません。

ifの条件が(line a)であることをすべてのスレッドがチェックすると、コードがあまりにもシリアルにならないように、count[0]が最新であることを確認する方法はありますか？

Answer 1

だけアトミックカウンターを毎回インクリメントし、そのテストでreturn value使用：あなたは約15アイテムが2.9を超え、Nmaxは10の周りで言うように、「少数のがあるでしょう場合

... 
    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 
    } 
... 

を余分な "原子操作、オーバーヘッドはおそらく最小限です（そして、私はより効率的にそれを行う方法を見ることができませんが、これは不可能ではありません...）。