CUDA


Hace algunos dias me tope con una tecnología de nVidia llamada CUDA. uhmmm ¿Para que sirve? básicamente para utilizar el procesamiento many-core de los GPU's.

Además los GPU's modernos nos permiten desarrollar programas utilizando vector programming, así como utilizar algunas "opciones extra" a la hora de trabajar con floats...

Bien, a esto le vamos a agregar que podemos desarrollar nuestros programas en lenjuague C, e incluso integrarlo en varios IDE's.

Hace un rato me puse a integrarlo en VC++ 2005 (Guía ) y todo funciona perfectamente (ea ojo, solo vamos a utilizar el editor, para compilar necesitamos utilizar el compilador de CUDA)

Veamos un pequeño ejemplo pasado en C++


#include "stdafx.h"
#include <stdio.h>
#include <stdlib.h>
#include <math.h>
#include <windows.h>

#define M_PI 3.14159265358979323846f

int _tmain(int argc, _TCHAR* argv[])
{
DWORD start = GetTickCount();
const int N = 200000;
float a[200000] = {0};

for (int i=0; i<N; i++)
a[i] = (float)i;


for (int i=0; i<N; i++)
a[i] =
sqrtf(sqrtf(cosf(a[i] / 1024 * (M_PI / 180.0f))) * sinf(1.0f)) *
sqrtf(sqrtf(cosf(a[i] / 1024 * (M_PI / 180.0f))) * sinf(1.0f)) *
sqrtf(sqrtf(cosf(a[i] / 1024 * (M_PI / 180.0f))) * sinf(1.0f));


DWORD end = GetTickCount();
printf("Processing time: %d %d (ms) \n" , end - start);

system("PAUSE");
return 0;
}


Sencillo, simplemente creamos un array y le aplicamos operaciones a lo bruto (a lo bruto).

Ahora nuestro ejemplo portado a CUDA sería algo así:


#include <tchar.h>
#include <stdio.h>
#include <cuda.h>
#include <cutil.h>

#define M_PI 3.14159265358979323846

__global__ void kernel(float *a, int N)
{
int idx = blockIdx.x * blockDim.x + threadIdx.x;
if (idx<N) a[idx] =
sqrtf(sqrtf(cosf(a[idx] / 1024 * (M_PI / 180.0f))) * sinf(1.0f)) *
sqrtf(sqrtf(cosf(a[idx] / 1024 * (M_PI / 180.0f))) * sinf(1.0f)) *
sqrtf(sqrtf(cosf(a[idx] / 1024 * (M_PI / 180.0f))) * sinf(1.0f));
}


int main(void)
{
unsigned int timer = 0;
CUT_SAFE_CALL(cutCreateTimer(&timer));
CUT_SAFE_CALL(cutStartTimer(timer));

float *a_h, *a_d;
const int N = 200000;

size_t size = N * sizeof(float);
a_h = (float *)malloc(size);


cudaMalloc((void **) &a_d, size);

for (int i=0; i < N; i++)
a_h[i] = (float)i;

cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);


int block_size = 4;
int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);


kernel <<< n_blocks, block_size >>> (a_d, N);


cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost);
CUT_SAFE_CALL(cutStopTimer(timer));


printf("Processing time: %f (ms) \n", cutGetTimerValue(timer));
CUT_SAFE_CALL(cutDeleteTimer(timer));

free(a_h);
cudaFree(a_d);
system("PAUSE");
}


Solo que esta pieza corre dos veces más rápido!! (AMD Athlon X2 vs GeForce 8800 GTX) y eso que aún estamos contando el tiempo de creación de datos que en realidad solo pasaría una vez en el programa.

¿En que es diferente el código?
El código en CUDA, tiene una función extra, esta función es nuestro kernel, en otras palabras la función que cargaremos al GPU; esto es elemental puesto que a todos los datos los vamos a procesar igual (como en los shaders), ademas todos los datos debemos de crearlos en el host, osea en la PC (ouch!! solo imaginen que la tarjeta de video esta aparte) y en el dispositivo (la tarjeta de video que suponemos esta aparte) y al hacer esto debemos de preocuparnos por su sincronización.

Este ejemplo es estupidamente sencillo, pero ahora imaginemos que podemos procesar partículas (jejeje voy a hacer eso esta semana), fluidos, física, u otro tipo de cosas donde el procesamiento es elemental, sobre todo cuando vamos a aplicar una misma función a una gran cantidad de datos (SIMD - Simple Input Multiple Data).

1 comentarios:

Niobé Jerez del Castillo dijo...

Genial, pero mira, en la llamada que haces al kernel, cómo sabes qué es lo que le tienes que pasar?? me refiero que cómo sabes que tienes que poner dentro de <<< >>> , muchas gracias! saludos

Publicar un comentario

About Me

Mi foto
carlos
Guadalajara, Jalisco, Mexico
Ver mi perfil completo

Twitter