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)

Get Raw

0001/*

0002 * tester.c

0003 */

0004int main() {

0005  struct BigInteger* a = (struct BigInteger*)malloc(sizeof(struct BigInteger));

0006  char* s = "1234567890";

0007  int i = 0;

0008 

0009  newBI_testing(a, s, 0);

0010}

Get Raw

0001/*

0002 * BigInteger.cu

0003 */

0004#include "stdio.h"

0005#include "conio.h"

0006#include "stdlib.h" 

0007#include "string.h"

0008#include "BigInteger.h"

0009#include "limits.h"

0010 

0011#include "cuda_runtime.h"

0012#include "device_launch_parameters.h"

0013 

0014__device__ int CU_MAX_LENGTH = 4096;

0015__device__ int CU_RET = 0;

0016__device__ int BLOCK_SIZE = 256;

0017__device__ int NUM_BLOCKS = 16; 

0018 

0019/*

0020 * newBI_testing

0021 *

0022 * Preámbulo para pruebas

0023 */

0024void newBI_testing(void* dst, char* s, int sig) {

0025  //variables para CUDA

0026  struct BigInteger* cu_dst;

0027  char* cu_s;

0028  int* cu_sig;

0029 

0030  //allocamos memoria

0031  cudaMalloc((void**)&cu_dst, sizeof(struct BigInteger));

0032  cudaMalloc((void**)&cu_s, sizeof(char) * strlen(s));

0033  cudaMalloc((void**)&cu_sig, sizeof(int));

0034 

0035  //copiamos s a cu_s y sig a cu_sig

0036  cudaMemcpy(cu_s, s, sizeof(char)* strlen(s), cudaMemcpyHostToDevice);

0037  cudaMemcpy(cu_sig, &sig, sizeof(char)* strlen(s), cudaMemcpyHostToDevice);

0038 

0039  //llamamos a CUDA

0040  _newBI_testing<<<1, 1>>>(cu_dst, cu_s, cu_sig);

0041 

0042  //sincronizamos

0043  cudaDeviceSynchronize();

0044   

0045  //copiamos de vuelta

0046  cudaMemcpy(dst, cu_dst, sizeof(struct BigInteger), 

0047             cudaMemcpyDeviceToHost);

0048}

0049 

0050/*

0051 * _newBI_testing

0052 *

0053 * Lanzadera para _newBI

0054 */

0055__global__ static void _newBI_testing(void* dst, char* s, int* sig) {

0056  _newBI(dst, s, sig);

0057}

0058 

0059/*

0060 * Función newBI.

0061 *

0062 * Genera un nuevo dato BI a partir del string que recibe como entrada.

0063 * Se cargan en orden inverso para permitir el crecimiento de manera 

0064 * sencilla.

0065 */

0066__device__ void _newBI(void* dst, char* s, int* sig){

0067  //longitud del string

0068  int i = cu_strlen(s) - 1;

0069 

0070  //limpiamos el array

0071  clean<<<NUM_BLOCKS, BLOCK_SIZE>>>(dst);

0072 

0073  //validamos que  no sobrepase el límite establecido

0074  if (i > CU_MAX_LENGTH) {

0075    showError(1);

0076 

0077    return;

0078  }

0080  //semáforo para clean.

0081  //mientras tanto, se puede ir validando la longitud (no afecta a ret)

0082  cudaDeviceSynchronize();

0083 

0084  //recorremos el string y lo guardamos en integers

0085  newBI_fill<<<NUM_BLOCKS, BLOCK_SIZE >>>(dst, s, i);

0086 

0087  //ajustamos la longitud

0088  ((struct BigInteger*)dst)->count = i;

0089 

0090  //semáforo para newBI_fill

0091  //mientras tanto, vamos ajustando la longitud de ret

0092  cudaDeviceSynchronize();

0093 

0094  //validamos signo

0095  if (*sig == -1)

0096    ((struct BigInteger*)dst)->n[((struct BigInteger*)dst)->count] *= -1;

0097}

0099/*

0100 * Función _newBI_fill.

0101 *

0102 * Rellena a.n de manera paralela a partir del string s de manera inversa

0103 * usando len como punto medio

0104 */

0105__global__ static void newBI_fill(void* va, char* s, int len) {

0106  int index = blockIdx.x * blockDim.x + threadIdx.x;

0107  int stride = blockDim.x * gridDim.x;

0108  int i = 0;

0109  char c;

0110 

0111  for (i = index; i < len; i += stride) {

0112    if (i > len) {

0113      //puede que len no sea divisible por 16 

0114      return;

0115    }

0116 

0117    //vamos capturando los caracteres de manera inversa

0118    c = (int)(s[len - i] - 48);

0119 

0120    if (c >= 0 && c <= 9) 

0121      ((struct BigInteger*)va)->n[i] = c;

0122    else {

0123      showError(3);

0125      return;

0126    }

0127  }

0128}

0129 

0130/*

0131 * Función clean.

0132 *

0133 * Limpia la estructura

0134 */

0135__global__ static void clean(void* va){

0136  int index = blockIdx.x * blockDim.x + threadIdx.x;

0137  int stride = blockDim.x * gridDim.x;

0138  int i = 0;

0139 

0140  for (i = index; i < CU_MAX_LENGTH; i += stride) 

0141    ((struct BigInteger*)va)->n[i] = 0;

0142}

0143 

0144/*

0145 * Función showError.

0146 *

0147 * Muestra un error en base al índice que se le pasa

0148 */

0149__device__ static void showError(int k) {

0150  if (k == 1)

0151    printf("Error. Limite alcanzado");

0152  else if (k == 3)

0153    printf("Error. Datos erróneos en newBI");

0154    /*etc..*/

0155 

0156  printf("\n");

0157 

0158  //modificamos la variable de retorno para poder validarla

0159  CU_RET = k;

0160}

0161 

0162/*

0163 * Función cu_strlen.

0164 *

0165 * Sinónimo de strlen C

0166 */

0167__device__ static int cu_strlen(char* s) {

0168  int ret = 0;

0169 

0170  while (*s++ != '\0') 

0171    ++ret;

0172 

0173  return ret;

0174}