Publicado el

Suma concurrente de vectores con NVidia CUDA

CUDA es una arquitectura hardware y software para la computación en la GPU. Consiste en utilizar el procesador gráfico o GPU como coprocesador junto con la CPU.

La GPU es un procesador especializado cuyas capacidades de cómputo se resumen en lo siguiente:

  • Múltiples núcleos dirigidos por un gran ancho de banda de memoria.
  • Gran paralelismo.
  • Optimización para cálculos en coma flotante.
  • Bajo coste.
  • Ingente cantidad de transistores.
    • Procesamiento de datos.
    • Almacenamiento en caché.
    • Control del flujo.

Estos son generalmente los pasos de cualquier aplicación CUDA:

  1. Copiar los datos a procesar desde la memoria de la CPU a la memoria de la GPU.
  2. Invocar el Kernel en la GPU.
  3. El Kernel ejecuta iteraciones del bucle en paralelo en cada núcleo.
  4. Copiar el resultado desde la memoria de la GPU a la memoria de la CPU.

El siguiente programa utiliza esta arquitectura para realizar de forma concurrente la suma de dos vectores a y b, y almacenar el resultado en un tercer vector c. Cabe destacar que este es el ejemplo más sencillo que se puede encontrar del uso de CUDA.

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71

#include 
#include 
#include   
 
using namespace std;
 
#define	TAM	8 // Tamanho del vector
 
// Declaración del kernel
__global__ void kernelSuma(int* a, int* b, int* c);
 
int main(int argc, char *argv[]) {
	int *a_h, *b_h, *c_h;
	int *a_d, *b_d, *c_d;
 
	// Reservamos memoria para los vectores host
	a_h = new int[TAM];
	b_h = new int[TAM];
	c_h = new int[TAM];
 
	// Reservamos memoria en la GPU
	cudaMalloc((void **)&a_d, TAM*sizeof(int));
	cudaMalloc((void **)&b_d, TAM*sizeof(int));
	cudaMalloc((void **)&c_d, TAM*sizeof(int));
 
        // Se configura un tamanio de bloque de 4 hilos
	int tamanio_bloque = 4;
	int num_bloques = TAM/tamanio_bloque + (TAM%tamanio_bloque == 0 ? 0:1);
 
        // Inicialización de los vectores en memoria host
        for(int i=0;i<TAM;i++) {
               a_h[i] = i;
               b_h[i] = i*2;
               }
 
 
     	// Copiamos los datos desde la memoria host a la memoria device
	cudaMemcpy(a_d, a_h, TAM*sizeof(int), cudaMemcpyHostToDevice);
     	cudaMemcpy(b_d, b_h, TAM*sizeof(int), cudaMemcpyHostToDevice);
 
	kernelSuma<<<tamanio_bloque, num_bloques>>>(a_d, b_d, c_d);
 
	// Copia los resultados del vector c desde la memoria de la GPU
        // a la memoria del host
	cudaMemcpy(c_h, c_d, TAM*sizeof(int), cudaMemcpyDeviceToHost);
 
        // Imprime los resultados
	for(int i = 0; i < TAM; i++)
		cout << c_h[i] << " ";
 
	cout << endl;
 
	// Libera la memoria
 
        // Tanto en el host...
	free(a_h);
	free(b_h);
	free(c_h);
        // ...como en la device.
	cudaFree(a_d);
	cudaFree(b_d);
	cudaFree(c_d);
 
} 
 
// Definición del kernel
__global__ void kernelSuma(int *a, int *b, int *c) {
        // i = Número de bloque * dimensión del bloque + hilo dentro del bloque
	int i = blockIdx.x * blockDim.x + threadIdx.x;
	c[i] = a[i]+b[i];
}