Hacer mientras no trabaje dentro de Kernel CUDA

Ok, soy bastante nuevo en CUDA, y estoy algo perdido, realmente perdido.

Estoy tratando de calcular pi usando el método de Monte Carlo, y al final solo obtengo un agregado en lugar de 50.

No quiero “hacer tiempo” para llamar al kernel, ya que es demasiado lento. Mi problema es que mi código no se repite, se ejecuta solo una vez en el kernel.

Y también, me gustaría que todos los hilos accedan al mismo nitro y pi, de modo que cuando un hilo golpee los contadores, todos los demás se detendrían.

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

Hay una variedad de problemas con su código.

  1. Sugiero utilizar la comprobación de errores de cuda adecuada y ejecutar su código con cuda-memcheck para detectar cualquier error de tiempo de ejecución. He omitido la comprobación correcta de errores en mi código a continuación para abreviar la presentación, pero lo he ejecutado con cuda-memcheck para indicar que no hay errores de tiempo de ejecución.

  2. El uso de curand() probablemente no sea correcto (devuelve números enteros en un rango grande). Para que este código funcione correctamente, desea una cantidad de punto flotante entre 0 y 1. La llamada correcta para eso es curand_uniform() .

  3. Dado que desea que todos los subprocesos funcionen con los mismos valores, debe evitar que esos subprocesos se crucen entre sí. Una forma de hacerlo es usar actualizaciones atómicas de las variables en cuestión.

  4. No debería ser necesario volver a ejecutar curand_init en cada iteración. Una vez por hilo debe ser suficiente.

  5. No cudaMemcpy..Symbol operaciones __shared__ variables __shared__ . Por conveniencia, y para preservar algo que se asemeja a su código original, he elegido convertirlos a variables __device__ .

Aquí hay una versión modificada de su código que soluciona la mayoría de los problemas anteriores:

 $ 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 $ 

He modificado las iteraciones a un número mayor, pero puedes usar 50 si quieres para ITER_MAX .

Tenga en cuenta que hay muchas críticas que podrían ser dirigidas contra este código. Mi objective aquí, ya que es claramente un ejercicio de aprendizaje, es señalar cuál sería el número mínimo de cambios para obtener un código funcional, utilizando el algoritmo que se describe. Como solo un ejemplo, es posible que desee cambiar la configuración de inicio del kernel ( <<<1,32>>> ) a otros números más grandes, con el fin de utilizar más plenamente la GPU.