Uno de los caballos de batalla más empleado en sistemas de optimización, minería de datos, sistemas expertos probabilísticos y simulación, es la utilización de números aleatorios, los números aleatorios nos permiten generar sobre la marcha poblaciones dentro de un espacio muestral controlado, que nos van a permitir, aplicando ciertas heurísticas, realizar inferencias y/o llegar a resultados. Normalmente estos métodos están clasificados dentro de lo que se conoce como "métodos Montecarlo" o basados en Montecarlo.
Por la intensa aplicación de estos métodos, se hace necesario disponer de mecanismos para la generación de números aleatorios que sea eficiente y con poco costo de procesamiento.
Existen varias soluciones para abordar el problema de la generación de números pseudoaleatorios con CUDA. Así, la librería CURAND provee al desarrollador de un importante surtido de funciones de alto nivel para la obtención de números aleatorios.
La librería CURAND contiene dos APIs: una para uso en código host, y la otra para uso en código de dispositivo.
Host API: para utilizar este API es necesario incluir curand.h en la lista de librerías en nuestro programa. Esta librería oculta al programador toda la problemática relacionada con el estado del generador de números aleatorios y su progresión de estado dentro del kernel de dispositivo. El funcionamiento básico es el siguiente:
1º) Se realiza la inicialización del estado del generador, para ello se utiliza la función curandCreateGenerator().
2º) Se especifican las opciones del generador, indicando además la semilla a utilizar. CURAND ofrece 2 tipos de generadores: CURAND_RNG_XORWOW, es un generador de números pseudoaleatorios que ha sido implementado usando el algoritmo XORWOW (32) introducido por George Marsaglia es una implementación de un generador del tipo xorshif también propuesto por él. Es un algoritmo muy rápido, calcula el siguiente elemento de la secuencia aplicando repetidamente un OR exclusivo sobre el resultado de aplicar un desplazamiento de bit al propio número. El código en lenguaje C publicado por Marsaglia es el siguiente:
unsigned long xor128(){
static unsigned long x=123456789,y=362436069,z=521288629,w=88675123;
unsigned long t;
t=(xˆ(x<<11));x=y;y=z;z=w;
return( w=(wˆ(w>>19))ˆ(tˆ(t>>8)) ) }
El Segundo tipo de generador que ofrece CURAND es un tipo de generador números quasialeatorios: CURAND_RNG_SOBOL32, que permite generar secuencias de 32-bit y hasta 20.000 dimensiones. Es un algoritmo algo más complejo ideado por I.M. Sobol en 1967, que permite crear secuencias uniformes de baja discrepancia utilizando muchas dimensiones.
Una vez seleccionado el tipo de RNG realizaremos la llamada a la función de la librería: curandSet-PseudoRandomGeneratorSeed() indicando la semilla y tipo de generador deseado.
3º) Reservar el espacio de memoria en el dispositivo.
4º) Invocar una función de generación de las disponibles pasando como parámetro el puntero a la memoria del dispositivo donde se generará la secuencia. Las distintas funciones de generación son:
- curandSetPseudoRandomGeneratorSeed(), permite generar números pseudoaleatorios y quasialeatorios cada elemento es de 32 bits, sin signo, donde todos sus bits son aleatorios.
- curandGenerateUniform(), genera una secuencia uniforme-mente distribuida de números en coma flotante comprendidos entre (0.0, 1.0].
- curandGenerateNormal(), permite obtener una secuencia de distribución normal, dados los valores de la media y la desviación estándar.
Para obtener secuencias con valores utilizando doble precisión se cuenta con las funciones homologas a las dos últimas: curandGenerateUniformDouble(), para la distribución uniforme y curandGenerateNormaDouble() para la normal.
5º) Una vez terminados los procesos deseados, liberamos los recursos solicitados mediante el uso de la instrucción curandDestroyGenerator().
Device API: para utilizar esta modalidad, deberemos incluir la librería curand_kernel.h en nuestro programa. Como es evidente, las funciones aquí definidas se invocarán en los kernels de dispositivo o en funciones de dispositivo, diseñadas para ser ejecutadas por los kernel. En general, el uso de esta API permite una mejor optimización del código y ofrece mejores tiempos de respuesta. La librería permite, al igual que su versión para host, generar secuencias de números pseudoaleatorios y quasialeatorios.
Las funciones principales son las siguientes: curand_init(), para iniciar la secuencia, la función curand(), que permite obtener un número pseudoaleatorio entero sin signo. curand_uniform(), obtiene un número en coma flotante de una distribución uniforme en el intervalo (0.0, 1.0]. curand_normal(), obtiene un número en coma flotante de una distribución normal en el intervalo (0.0, 1.0] utilizando un valor de 0 para la media y de 1 para la desviación estándar.
También se dispone de versiones de estas dos últimas funciones para obtener secuencias con números de doble precisión: curand_uniform_double() y curand_normal_double(). Adicionalmente contamos con dos funciones que permiten obtener pares de elementos utilizando la estructura float2 y double2, que componen de un campo x y otro y, para obtener secuencias de distribución normal: curand_normal2() y curand_normal2_double() respectiva-mente.
El valor que se pasa a estas funciones es el estado del generador, que puede ser de dos tipos, dependiendo del generador que deseemos utilizar: curandState si deseamos utilizar un RNG pseudoaleatorio, o curandStateSobol32 si deseamos obtener valores de la secuencia quasialeatorios. Para el primer caso, en el que se trabaja con números pseudoaleatorios hay que pasar una semilla a la función, mientras que en el caso de quasialeatorios hay que pasar como argumento los vectores de dirección y el desplazamiento. Finalmente cabe destacar una función para pasar por alto elementos de la secuencia, que es lo mismo que llamar varias veces a la función curand() pero mucho más rápido: skipahead(), en esta función indicamos el número de elementos a descartar y el estado.
Para la inicialización del generador de números aleatorios en CUDA, utilizamos la función curand_init(). Como particularidad hay que utilizar una estructura de datos que mantenga el estado del generador y lo actualice durante el proceso de generación de números. Necesitamos mantener y pasar esa estructura en cada kernel que quiera generar números aleatorios:
curandState* devStates;
cudaMalloc ( &devStates, Elementos*sizeof( curandState ) );
/* inicializacion de numeros aletorios */
randomize<<< ( Elementos + ( hilos – 1 )) / hilos, hilos >>> ( devStates, time(NULL) );
En el código anterior, en primer lugar se declara la estructura de estados de tipo curandState ya que se van a utilizar números pseudoaleatorios.
En la siguiente sentencia se reserva la memoria que necesita el dispositivo para actualizar los estados a cada hilo del kernel. Finalmente la última instrucción consiste en una llamada al kernel reandomize (la hemos llamado igual que la función existente en muchos lenguajes de alto nivel para la inicialización de la generación de números aleatorios). Como semilla se emplea el valor del tiempo actual del sistema, preparando así los estados para la obtención de números aleatorios.
__global__ void randomize( curandState *state, unsigned long semilla )
{
int id = threadIdx.x + blockIdx.x * blockDim.x;
curand_init ( semilla, id, 0, &state[id] );
}
El kernel randomize se encarga de iniciar el estado para cada hilo al llamar a curand_init. Separamos esta tarea en un kernel independiente con el objetivo de aumentar el rendimiento global del algoritmo. Por un lado para reducir el uso de la pila de memoria por hilo, ya que la generación de los estados de números aleatorios, en el peor de los casos puede llegar a consumir hasta 16KB. Por otro lado al separar esta tarea de la obtención de números aleatorios curand() permitimos al kernel, que está manejando la obtención, disponer de más velocidad y aumentar sus posibilidades de concurrencia al requerir menos recursos del procesador. Esto permite manejar más bloques por procesador.
Para obtener números aleatorios dentro del kernel utilizaremos la función curand_uniform( &localState ), que como hemos indicando anteriormente obtienen un elemento de la secuencia (en nuestro caso pseudoaleatoria) de elementos en coma flotante dentro del intervalo (0.0, 1.0].
Yo por el momento cuento con una modesta GeForce GTX 660, las pruebas que he realizado en principio no me han dejado muy impresionado. Supongo que CUDA 5 permitirá exprimir a la perfección esta arquitectura.
