Ne fonctionne pas dans le kernel CUDA

Ok, je suis assez nouveau dans CUDA, et je suis un peu perdu, vraiment perdu.

J’essaie de calculer pi en utilisant la méthode de Monte Carlo, et à la fin, je ne reçois qu’un ajout au lieu de 50.

Je ne veux pas “faire le temps” pour appeler le kernel, car il est trop lent. Mon problème est que mon code ne tourne pas en boucle, il ne s’exécute qu’une fois dans le kernel.

Et aussi, j’aimerais que tous les threads accèdent au même niter et pi, ainsi, quand un thread atteindrait les compteurs, tous les autres s’arrêteraient.

#define SEED 35791246 __shared__ int niter; __shared__ double pi; __global__ void calcularPi(){ double x; double y; int count; double z; count = 0; niter = 0; //keep looping do{ niter = niter + 1; //Generate random number curandState state; curand_init(SEED,(int)niter, 0, &state); x = curand(&state); y = curand(&state); z = x*x+y*y; if (z<=1) count++; pi =(double)count/niter*4; }while(niter < 50); } int main(void){ float tempoTotal; //Start timer clock_t t; t = clock(); //call kernel calcularPi<<>>(); //wait while kernel finish cudaDeviceSynchronize(); typeof(pi) piFinal; cudaMemcpyFromSymbol(&piFinal, "pi", sizeof(piFinal),0, cudaMemcpyDeviceToHost); typeof(niter) niterFinal; cudaMemcpyFromSymbol(&niterFinal, "niter", sizeof(niterFinal),0, cudaMemcpyDeviceToHost); //Ends timer t = clock() - t; tempoTotal = ((double)t)/CLOCKS_PER_SEC; printf("Pi: %g \n", piFinal); printf("Adds: %d \n", niterFinal); printf("Total time: %f \n", tempoTotal); } 

Votre code pose divers problèmes.

  1. Je suggère d’utiliser la vérification d’erreur cuda appropriée et d’exécuter votre code avec cuda-memcheck pour détecter les erreurs d’exécution. J’ai omis de vérifier les erreurs dans le code ci-dessous pour la brièveté de la présentation, mais je l’ai exécuté avec cuda-memcheck pour ne pas indiquer d’erreur d’exécution.

  2. Votre utilisation de curand() n’est probablement pas correcte (elle renvoie des nombres entiers sur une large plage). Pour que ce code fonctionne correctement, vous voulez une quantité à virgule flottante comprise entre 0 et 1. L’ appel correct pour cela est curand_uniform() .

  3. Puisque vous voulez que tous les threads fonctionnent sur les mêmes valeurs, vous devez empêcher ces threads de se chevaucher. Une solution consiste à utiliser les mises à jour atomiques des variables en question.

  4. Il ne devrait pas être nécessaire de réexécuter curand_init à chaque itération. Une fois par fil devrait être suffisant.

  5. Nous n’utilisons pas d’opérations __shared__ sur les variables __shared__ . Par souci de commodité et pour conserver quelque chose qui ressemble à votre code d’origine, j’ai choisi de convertir ces variables en variables __device__ .

Voici une version modifiée de votre code qui résout la plupart des problèmes ci-dessus:

 $ cat t978.cu #include  #include  #include  #define ITER_MAX 5000 #define SEED 35791246 __device__ int niter; __device__ int count; __global__ void calcularPi(){ double x; double y; double z; int lcount; curandState state; curand_init(SEED,threadIdx.x, 0, &state); //keep looping do{ lcount = atomicAdd(&niter, 1); //Generate random number x = curand_uniform(&state); y = curand_uniform(&state); z = x*x+y*y; if (z<=1) atomicAdd(&count, 1); }while(lcount < ITER_MAX); } int main(void){ float tempoTotal; //Start timer clock_t t; t = clock(); int count_final = 0; int niter_final = 0; cudaMemcpyToSymbol(niter, &niter_final, sizeof(int)); cudaMemcpyToSymbol(count, &count_final, sizeof(int)); //call kernel calcularPi<<<1,32>>>(); //wait while kernel finish cudaDeviceSynchronize(); cudaMemcpyFromSymbol(&count_final, count, sizeof(int)); cudaMemcpyFromSymbol(&niter_final, niter, sizeof(int)); //Ends timer double pi = count_final/(double)niter_final*4; t = clock() - t; tempoTotal = ((double)t)/CLOCKS_PER_SEC; printf("Pi: %g \n", pi); printf("Adds: %d \n", niter_final); printf("Total time: %f \n", tempoTotal); } $ nvcc -o t978 t978.cu -lcurand $ cuda-memcheck ./t978 ========= CUDA-MEMCHECK Pi: 3.12083 Adds: 5032 Total time: 0.558463 ========= ERROR SUMMARY: 0 errors $ 

J’ai modifié les itérations en un nombre plus élevé, mais vous pouvez utiliser 50 si vous le souhaitez pour ITER_MAX .

Notez que de nombreuses critiques pourraient être adressées à ce code. Mon objective ici, dans la mesure où il s’agit clairement d’un exercice d’apprentissage, est de préciser le nombre minimal de modifications à apporter pour obtenir un code fonctionnel, en utilisant l’algorithme que vous avez décrit. À titre d’exemple, vous voudrez peut-être modifier la configuration de lancement du kernel ( <<<1,32>>> ) par un nombre plus important, afin d’utiliser davantage le processeur graphique.