Generar números aleatorios en CUDA

La generación de números aleatorios (RNG) tiene diversas aplicaciones en simulaciones computacionales, algoritmos evolutivos, método de Monte-Carlo, entre otros y por lo tanto, serán de importancia para el cálculo en GPGPUs.

En estos problemas, podemos distinguir:

1. Los Números Aleatorios "Verdaderos" (True Random Number): Los más complicados de generar, se basan en métodos no determinísticos, generalmente en fenómenos físicos (por ejemplo, radioactivos, atmósfera) que se espera tengan resultados aleatorios. La generación de éstos números no es periódica, es decir, que no se repite la secuencia de números generados. El sitio Random.org proporciona servicios gratuitos de generación de números aleatorios de éste tipo.
2. Los Pseudo Números Aleatorios: Usan algoritmos computacionales capaces de producir secuencias largas de números aparentemente aleatorios, los cuales están determinados por un valor inicial al que se le denomina semilla (seed). En estos números la secuencia eventualmente se repite.

Obviamente es imposible generar números aleatorios verdaderos en una computadora determinística. La función de aleatoriedad aplica algún tipo de transformación sobre otro número, determinando una sucesión que "parece" aleatoria. De cualquier forma, si dos generaciones de números empezaran a partir de la misma semilla, el resultado sería el mismo.

En la programación en C/C++, se hace gran uso de la función rand() para generar estos tipos de números pseudoaleatorios. Las semilla generalmente más usada es el tiempo, lo cual genera resultados aceptables.

Sin embargo,al generar números aleatorios en una GPGPU puede volverse algo complicado. La solución más sencilla e ingenua es crear todos los números aleatorios necesarios en el host y colocarlos en la memoria global de la GPGPU (pre-generación). La desventaja está en el bandwith necesario para transferir dichos números a la memoria del dispositivo.

Por lo tanto es más eficiente generar los números aleatorios directamente en la memoria del dispositivo en un kernel exclusivamente dedicado a ello. Para ello se puede aprovechar la paralelización en la generación de dichos números, obviamente tomando en cuenta que es necesario partir de diferentes semillas, si no, el número generado sería el mismo. La biblioteca NVIDIA CURAND hace más fácil la creación de números dentro del kernel del dispositivo. Dichos números estarán almacenados en la memoria local, y disponibles para el cálculo que se requiera.

Los números generados por CURAND son pseudoaleatorios y/o cuasialeatorios. Una secuencia de pseudoaleatorios satisface la mayoría de las propiedades estadísticas de una secuencia de números verdaderamente aleatorios, sin embargo, es generada por un algoritmo determinista. Una secuencia cuasialeatoria de puntos n-dimensionales es determinada por un algoritmo determinista diseñado para llenar el espacio n-dimensional.

A continuación se muestra un código de ejemplo, genera un vector de números flotantes en el dispositivo. Para efectos de muestra, se copian al host e imprime.

#include <stdio.h>
#include <curand_kernel.h>
#include <time.h>

__global__ void setup_kernel(curandState * state, unsigned long seed) {
int id = threadIdx.x;

/* cada thread tiene la misma semilla, y un diferente número
* de secuencia
*/
curand_init(seed, id, 0, &state[id]);

}

__global__ void generate(curandState* globalState, float *result) {
int ind = threadIdx.x;

// copiar estado a la memoria local para mayor eficiencia
curandState localState = globalState[ind];

// generar número pseudoaleatorio
float r = curand_uniform(&localState);

//copiar state de regreso a memoria global
globalState[ind] = localState;

//almacenar resultados
result[ind] = r;
}

int main(int argc, char** argv) {
int N = 30;

curandState* devStates;
float *devResults;
float hostResults[N];

// reservando espacio para los states PRNG en el device
cudaMalloc(&devStates, N * sizeof(curandState));

// reservando espacio para el vector de resultados en device
cudaMalloc((void**) &devResults, N * sizeof(float));

dim3 tpb(N, 1, 1);

// setup semillas
setup_kernel<<<1, tpb>>>(devStates, time(0));

// generar números aleatorios
generate<<<1, tpb>>>(devStates, devResults);

cudaMemcpy(hostResults, devResults, N * sizeof(float),
cudaMemcpyDeviceToHost);

cudaFree(devStates);
cudaFree(devResults);

for (int i = 0; i < N; i++) {
printf("%f\n", hostResults[i]);
}

return 0;
}


Link: CUDA Curand