Cuda原子和條件分支

我試圖編寫CUDA版本的serial代碼作爲在分子動力學算法中實現週期性邊界條件的一部分。這個想法是，有一小部分位置在框外的粒子需要使用兩個ways中的一個來重新使用，並限制我使用第一種方式的次數。Cuda原子和條件分支

本質上，它歸結爲以下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最低位使用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有些類似，但是，我不確定是否以及如何使用關鍵部分會有所幫助。

有沒有辦法確保count[0]是最新的，當每個線程檢查(line a)上的if條件時，不會使代碼太串行？

來源

2017-08-30 h k

我認爲@Claude的答案是一個很好的答案，如果你可以容忍'count [0]'總是遞增的，那麼它就非常簡單和整齊。要重現序列代碼的確切行爲（'count'遞增，直到達到'Nmax'，然後停止），應該可以使用圍繞'atomicCAS'構建的自定義原子。然而，我懷疑在性能方面它會比克勞德的答案昂貴。 –

每次只需遞增原子計數器，並在測試使用其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，會有少量的「額外的「原子操作，其開銷可能很小（我看不出如何更有效地做到這一點，這並不是說這是不可能的）。

來源

2017-08-30 00:52:48 Claude

Cuda原子和條件分支

回答

相關問題