"COCINANDO" CON TARJETAS GRÁFICAS (GPU): enero 2014

viernes, 31 de enero de 2014

Grids Bidimensionales

Esta entrada mostrará un código donde el kernel será lanzado con bloques de 2D, y un grid de threads 2D para realizar operaciones sobre una matriz de datos. Las operaciones ejemplifican la conversión de los índices del thread y del bloque en 1D, esto debido a que el manejo de los datos en la GPGPU es 1D.

#include <stdlib.h>
#include <stdio.h>

__global__ void kernel(int *array) {

// se obtienen los índices x y y para cada thread.
int index_x = blockIdx.x * blockDim.x + threadIdx.x;
int index_y = blockIdx.y * blockDim.y + threadIdx.y;

// el array viene en forma 1D (a pesar de ser matriz), por lo
// que hay que covertir de 2D a un índice 1D

int grid_width = gridDim.x * blockDim.x;
int index = index_y * grid_width + index_x;

// convierte el índice de bloque de 2D en un índice 1D.
int result = blockIdx.y * gridDim.x + blockIdx.x;

// escribe el resultado
array[index] = result;
}
int main(void) {
int num_elements_x = 16;
int num_elements_y = 16;

int num_bytes = num_elements_x * num_elements_y * sizeof(int);

int *device_array = 0;
int *host_array = 0;

// reserva memoria en host
host_array = (int*) malloc(num_bytes);
// reserva memoria en GPU
cudaMalloc((void**) &device_array, num_bytes);

// crea bloques bidimensionales de threads de 4 x 4
dim3 block_size;
block_size.x = 4;
block_size.y = 4;

// configura un grid bidimensional de 4 x 4
dim3 grid_size;
grid_size.x = num_elements_x / block_size.x;
grid_size.y = num_elements_y / block_size.y;

// se envían el arreglo de bloques y el grid de threads por bloque
kernel<<<grid_size, block_size>>>(device_array);

// copia al host
cudaMemcpy(host_array, device_array, num_bytes, cudaMemcpyDeviceToHost);

// impresión
for (int row = 0; row < num_elements_y; ++row) {
for (int col = 0; col < num_elements_x; ++col) {
printf("%2d ", host_array[row * num_elements_x + col]);
}
printf("\n");
}
printf("\n");

// liberar memoria
free(host_array);
cudaFree(device_array);
}

Salida del programa:

0 0 0 0 1 1 1 1 2 2 2 2 3 3 3 3

4 4 4 4 5 5 5 5 6 6 6 6 7 7 7 7

8 8 8 8 9 9 9 9 10 10 10 10 11 11 11 11

12 12 12 12 13 13 13 13 14 14 14 14 15 15 15 15

miércoles, 22 de enero de 2014

Suma de vectores II

El principal propósito de trabajar con una GPGPU es el de potenciar el cálculo con cantidades inmensa de datos. Hemos descrito anteriormente un post donde se realizaba la suma de vectores funcionando perfectamente, sin embargo a continuación se muestra un código que funciona para una cantidad de elementos mucho mayor del vector de números:

__global__ void sumaVector(long *v1, long *v2, long *v3, long N) {

int threadId = blockIdx.x * blockDim.x + threadIdx.x;
while (threadId < N) {
v3[threadId] = v1[threadId] + v2[threadId];
threadId += blockDim.x * gridDim.x;
}
}

int main(int argc, char** argv) {
long N = 9000000; // 9 millones

long *v1, *v2, *v3;
v1 = (long *) malloc(N * sizeof(long));
v2 = (long *) malloc(N * sizeof(long));
v3 = (long *) malloc(N * sizeof(long));

for (long i = 0; i < N; i++) {
// datos de prueba
v1[i] = 10;
v2[i] = 11;

}

printf("%ld", v1[1000001]); //OK

long *dv1, *dv2, *dv3;

cudaMalloc((void**) &dv1, N * sizeof(long));
cudaMalloc((void**) &dv2, N * sizeof(long));
cudaMalloc((void**) &dv3, N * sizeof(long));

// copiando memoria a la GPGPU
cudaMemcpy(dv1, v1, N * sizeof(long), cudaMemcpyHostToDevice);
cudaMemcpy(dv2, v2, N * sizeof(long), cudaMemcpyHostToDevice);

// número de bloques
int B = 1024;
int T = 1024;

// Llamando a ejecutar el kernel
sumaVector<<<B, T>>>(dv1, dv2, dv3, N);

// copiando el resultado a la memoria Host
cudaMemcpy(v3, dv3, N * sizeof(long), cudaMemcpyDeviceToHost);

cudaFree(dv1);
cudaFree(dv2);
cudaFree(dv3);

printf("%ld", v3[900001]); //OK

return (EXIT_SUCCESS);
}

Características importantes a tomar en cuenta

Este cálculo se realiza en dimensión 1D.
Se define la llamada del kernel para 1024 bloques y cada bloque con 1024 threads. Por lo tanto la GPGPU ejecutará 1,048,576 threads.
El vector tiene un tamaño de 9 millones de elementos (se usó valores long para disponer en cualquier caso de números grandes).
Debido a que el número de elementos excede el número de threads en el kernel, la codificación de éste incluye un ciclo que indica que se deberá estar realizando hasta que se terminen de calcular las 9 millones de veces.