Cuda atomics et branches conditionnelles

Question

Je tente d'écrire une version CUDA d'un code serial dans le cadre de la mise en œuvre d'une condition aux limites périodiques dans un algorithme de dynamique moléculaire. L'idée est qu'une petite fraction de particules qui ont des positions hors de la boîte doivent être remis en utilisant l'un des deux ways, avec une limite sur le nombre de fois que j'utilise la première voie. Essentiellement, cela se résume à la MWE suivante. J'ai un tableau x[N] où N est grand, et le suivant serial code.

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

S'il vous plaît supposer que x[i] > 2.9 seulement pour une petite fraction (environ 12-15) des 30000 éléments de x[i].

Notez que la séquence de i n'est pas important, à savoir qu'il n'est pas nécessaire d'avoir le plus bas 10i à utiliser x[i] += 0.1, ce qui rend l'algorithme potentiellement parallélisables. Je pensais à la version CUDA suivante de la MWF, qui compile avec nvcc -arch sm_35 main.cu, où main.cu se lit comme

#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); 

} 

Bien sûr, ce n'est pas correct parce que la condition if sur (line a) utilise une variable qui est mise à jour en (line b), qui pourrait ne pas être courant. Ceci est quelque peu similaire à Cuda atomics change flag, cependant, je ne suis pas sûr si et comment l'utilisation de sections critiques aiderait.

Y at-il un moyen de s'assurer que count[0] est à jour lorsque chaque thread vérifie la condition if sur (line a), sans rendre le code trop série?

Answer 1

incrémenter Juste le compteur atomique à chaque fois, et utiliser ses return value dans votre test:

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

Si comme vous le dites environ 15 articles dépassent 2,9 et Nmax est d'environ 10, il y aura un petit nombre de " opérations atomiques supplémentaires, dont le surcoût est probablement minime (et je ne vois pas comment le faire plus efficacement, ce qui ne veut pas dire que ce n'est pas possible ...).