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 

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

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

 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 

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

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

 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 

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

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

 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 

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

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

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