BigInteger - Codificación de la arquitectura: Creación de BigInteger

 Buenas!

Otro fin de semana que llega, cada vez hace más calor... (menos mal que ya termina el verano :-) )

Pues aquí seguimos, con la buena racha. De momento, he conseguido rehacer todo el setup de CUDA, y de momento he conseguido avanzar (no prometo un avance semanal, así que tengo alguna entrada guardada en la recámara!).

Vamos a entrar de lleno en la codificación paralela con CUDA. En concreto, empezamos con la creación de BigInteger.

• ¿No es la entrada que buscabas? Entra en el índice
• ¿Buscas el código fuente? Aquí lo tienes

¡Vamos allá!

Vamos a empezar con una pequeña introducción de conceptos y demás idiosincrácia própia de la arquitectura CUDA.

• Preámbulo: parte del código C que sirve para ajustar el contexto de CUDA. Dentro del preámbulo, realizamos las siguientes operaciones
• Crear variables para CUDA
• Copiar memoria de CPU (código C) a GPU (código CUDA)
• Llamar a la función CUDA principal

Hay que destacar la diferencia entre CPU (o código "host") y GPU (o código "device"). No es posible llamar a código "device" desde código "host", por lo que existe un tipo de función híbrido "global", que permite ejecutar código "host" y "device" de manera conjunta.

También, las funciones "device" no pueden usar funciones nativas de "host" (por ejemplo strlen), mientras que las funciones "host" no pueden usar funciones nativas de "device"; las funciones "global" pueden mezclar ambas tecnologías.

Por tanto, lo habitual es:

• Creamos variables "host" en una función puramente "host"
• Llamamos a una función "global", pasando las variables. Esta función será el preámbulo.
• La función "global" llama a las funciones "deivice"

Pensaréis que esto es todo, ¿verdad? ¡Pues no!

¿Recordáis toda la lata que os dimos con los clústeres? Pues bien, esots clústeres se definen en las llamadas a las funciones "global"; y para complicar más la cosa, una función "device" no puede llamar a una función "global" (a no ser que hagamos algunos cambios en la configuración del compilador), ya que esto es una característica avanzada llamada "palalelismo dinámico" (y que es lo que hemos diseñado).

Bien, ahora sí, estamos en disposición de ver algo de código (ya era hora)

[code lan=gral] /* * tester.c */ int main() { struct BigInteger* a = (struct BigInteger*)malloc(sizeof(struct BigInteger)); char* s = "1234567890"; int i = 0; newBI_testing(a, s, 0); } [/code] [code lan=gral] /* * BigInteger.cu */ #include "stdio.h" #include "conio.h" #include "stdlib.h" #include "string.h" #include "BigInteger.h" #include "limits.h" #include "cuda_runtime.h" #include "device_launch_parameters.h" __device__ int CU_MAX_LENGTH = 4096; __device__ int CU_RET = 0; __device__ int BLOCK_SIZE = 256; __device__ int NUM_BLOCKS = 16; /* * newBI_testing * * Preámbulo para pruebas */ void newBI_testing(void* dst, char* s, int sig) { //variables para CUDA struct BigInteger* cu_dst; char* cu_s; int* cu_sig; //allocamos memoria cudaMalloc((void**)&cu_dst, sizeof(struct BigInteger)); cudaMalloc((void**)&cu_s, sizeof(char) * strlen(s)); cudaMalloc((void**)&cu_sig, sizeof(int)); //copiamos s a cu_s y sig a cu_sig cudaMemcpy(cu_s, s, sizeof(char)* strlen(s), cudaMemcpyHostToDevice); cudaMemcpy(cu_sig, &sig, sizeof(char)* strlen(s), cudaMemcpyHostToDevice); //llamamos a CUDA _newBI_testing<<<1, 1>>>(cu_dst, cu_s, cu_sig); //sincronizamos cudaDeviceSynchronize(); //copiamos de vuelta cudaMemcpy(dst, cu_dst, sizeof(struct BigInteger), cudaMemcpyDeviceToHost); } /* * _newBI_testing * * Lanzadera para _newBI */ __global__ static void _newBI_testing(void* dst, char* s, int* sig) { _newBI(dst, s, sig); } /* * Función newBI. * * Genera un nuevo dato BI a partir del string que recibe como entrada. * Se cargan en orden inverso para permitir el crecimiento de manera * sencilla. */ __device__ void _newBI(void* dst, char* s, int* sig){ //longitud del string int i = cu_strlen(s) - 1; //limpiamos el array clean<<<NUM_BLOCKS, BLOCK_SIZE>>>(dst); //validamos que no sobrepase el límite establecido if (i > CU_MAX_LENGTH) { showError(1); return; } //semáforo para clean. //mientras tanto, se puede ir validando la longitud (no afecta a ret) cudaDeviceSynchronize(); //recorremos el string y lo guardamos en integers newBI_fill<<<NUM_BLOCKS, BLOCK_SIZE >>>(dst, s, i); //ajustamos la longitud ((struct BigInteger*)dst)->count = i; //semáforo para newBI_fill //mientras tanto, vamos ajustando la longitud de ret cudaDeviceSynchronize(); //validamos signo if (*sig == -1) ((struct BigInteger*)dst)->n[((struct BigInteger*)dst)->count] *= -1; } /* * Función _newBI_fill. * * Rellena a.n de manera paralela a partir del string s de manera inversa * usando len como punto medio */ __global__ static void newBI_fill(void* va, char* s, int len) { int index = blockIdx.x * blockDim.x + threadIdx.x; int stride = blockDim.x * gridDim.x; int i = 0; char c; for (i = index; i < len; i += stride) { if (i > len) { //puede que len no sea divisible por 16 return; } //vamos capturando los caracteres de manera inversa c = (int)(s[len - i] - 48); if (c >= 0 && c <= 9) ((struct BigInteger*)va)->n[i] = c; else { showError(3); return; } } } /* * Función clean. * * Limpia la estructura */ __global__ static void clean(void* va){ int index = blockIdx.x * blockDim.x + threadIdx.x; int stride = blockDim.x * gridDim.x; int i = 0; for (i = index; i < CU_MAX_LENGTH; i += stride) ((struct BigInteger*)va)->n[i] = 0; } /* * Función showError. * * Muestra un error en base al índice que se le pasa */ __device__ static void showError(int k) { if (k == 1) printf("Error. Limite alcanzado"); else if (k == 3) printf("Error. Datos erróneos en newBI"); /*etc..*/ printf("\n"); //modificamos la variable de retorno para poder validarla CU_RET = k; } /* * Función cu_strlen. * * Sinónimo de strlen C */ __device__ static int cu_strlen(char* s) { int ret = 0; while (*s++ != '\0') ++ret; return ret; } [/code]

Vamos a analizar el código, por partes...

Lo que vemos en "tester.c" es lo que sería una función "host", código C de toda la vida.

En "BigInteger.cu" tenemos el preámbulo ("newBI_testing"), donde realizamos las acciones previas, y llamamos a la función "global" (podemos ver que lleva el identificador "__global__").

Finalmente, la función "global", llama a la función "device", que hace todos los cálculos.

Podemos ver que la llamada a "_newBI_testing" está precedida por <<<1, 1>>>. Estos parámetros adicionales indican la configuración del clúster. En este caso, <<<1, 1>>> le hace saber al compilador que queremos abrir un clúster de un hilo (aunque se reservarán 16).

Más tarde, en "_newBI", usamos <<<NUM_BLOCKS, BLOCK_SIZE>>>, que son variables que equivalen a <<<16, 256>>> (es decir, abrimos 256 clústeres de 16 hilos).

Podemos ver cómo el código es muy similar al de "newBI" convencional, con algunas excepciones:

• cudaDeviceSynchronize(): hace la función de semáforo, y nos sirve para esperar que todos los hilos finalicen
• newBI_fill(): hemos movido el bucle de llenado a una función a parte, para poder paralelizarla
• Aparecen las variables "index" y "stride": son variables que se usan en este contexto para poder direccionar correcamente los hilos. Aquí encontraréis una pequeña guía que lo explica mil veces mejor de lo que yo podría explicarlo.
• cu_strlen(): hemos tenido que crear un sinónimo de la función "strlen" para poder usar dentro de CUDA.

Comentar que, en la versión final, las funciones de "testing" no estarían, y sería responsabilidad del desarrollador realizar todos los preámbulos y operaciones previas para llamar a BigInteger dentro de un contexto CUDA.

Por último, un fun fact: se puede usar BigInteger.cu para crear un BigInteger, y luego operar con él desde BigInteger.c, ya que se usa la misma estructura de datos!

Con esto, os dejo, con suerte, hasta la semana que viene.

¡Hasta la próxima!