Definiendo las estructuras Grid/Block para un kernel

Para llamar a ejecutar un kernel es necesario proveer una configuración de ejecución, esto es, las dimensiones del grid y del bloque con que se ejecutará dicho kernel. Esta información se proporciona mediante dos estructuras claves:

• Número de bloques en cada dimensión
• Threads por bloque en cada dimensión

Sintaxis de llamada a un kernel:

      nombreKernel <<< B, T >>> (arg1, arg2, ... argN);

donde:

• B es una estructura que define el número de bloques en el grid para cada dimensión (1D o 2D)
• T es una estructura que define el número de threads en un bloque para cada dimensión (1D, 2D o 3D).

Si se desea definir una estructura de 1D, se puede usar un entero para B y T en:

       nombreKernel <<< B, T >>> (arg1, arg2, ... argN);

donde:

• B es un entero que define un grid de dimensión 1D
• T es un entero que define un bloque 1D de ese tamaño.

por ejemplo:

       nombreKernel <<< 1, 100 >>> (arg1, arg2, ... argN);

dim3 en CUDA

dim3 es una estructura de datos utilizada comúnmente en CUDA para pasar al kernel la configuración de ejecución (es decir, las dimensiones del grid y del bloque).

Por ejemplo:

dim3 grid(512);                        //512 x 1 x 1

dim3 block(1024, 1024);              //1024 x 1024 x 1

kernelEjemplo <<grid, block >>>();

Puntos a tomar en cuenta de dim3:

• dim3 tiene 3 elementos x, y y z.
• dim3 es una estructura definida en %CUDA_INC_PATH%/vector_types.h
• dim3 en código C puede ser inicializado como dim3 grid = {512, 512, 1};
• dim3 en código C++ puede ser inicializado como dim3 grid(512,512,1);
• No es necesario proporcionar los 3 elementos. Si alguno de ellos falta, será inicializado por defecto con valor 1.

Identificación de un thread

CUDA identifica a cada thread de forma única. Para ello se apoya de ciertas variables incluidas en el driver:

• blockIdx.x, blockIdx.y, blockIdx.z
• Retornan el valor del ID del bloque en el eje x, y o z, respectivamente, del bloque que esté ejecutando el código dado.
• blockDim.x, blockDim.y, blockDim.z
• Retornan el valor de la dimensión del bloque, es decir, el número de threads en un bloque en el eje x, y o z.
• threadIdx.x, threadIdx.y, threadIdx.z
• Retornan el ID del thread en el eje x, y o z que está siendo ejecutado por el GPGPU en un bloque en particular.

De ésta forma, se puede expresar una colección de bloques, y una colección de threads dentro de un bloque, como un array 1D, 2D o 3D, lo cual puede ser de ayuda con información en 2D o 3D.

Ejemplo:

Suponer un grid de 1D (un grid se mapea directamente con una GPGPU), y una arquitectura de bloques de 1D con 4 bloques y cada uno con 8 threads.

Para identificar el ID Global del thread 26:

GlobalThreadID = blockIdx.x * blockDim.x + threadIdx.x
GlobalThreadID =         3       *         8         +      2
GlobalThreadID = 26       

Conceptos básicos en CUDA

• Una GPGPU corre de forma asíncrona al CPU del host
• Una GPGPU corre en un espacio de memoria separado del CPU del host
•  La velocidad de la memoria del host es 8-20 GB/s, mientras que la velocidad de la memoria en una GPGPU es aproximadamente 160-200 GB/s.
• Los programas en CUDA usan kernels
• Un kernel es una subrutina llamada desde el host que se ejecutara en la GPGPU.
• Un kernel no es una función, ya que no puede retornar un valor.
• Un kernel está definido con una especificador declarativo __global__ , el cual le indica al compilador que el kernel es llamable desde el procesador del host.
• Las llamadas a kernels son asíncronas
•  El host encola un kernel para su ejecución en el GPGPU y no lo espera a que finalice.
• En algún momento después el kernel corre en el GPGPU
• Para mayor eficiencia, se pueden crear pipelines para encolar kernels y mantener el GPGPU ocupado tanto sea posible.
• Más adelante, puede ser necesario determinar cuando el kernel o la pipeline ha sido completada, por lo que se utilizan comúnmente dos mecanismos de sincronización:
• cudaThreadSynchronize(), hace que el host se detenga y espere a que todos los kernels encolados se ejecuten
• Ejecutar una transferencia de datos con bloque mediante  cudaMemcpy(), ya que incluye dentro la llamada a cudaThreadSyncrhonize().
• La unidad básica de trabajo de una GPGPU es un thread. 
• Cada thread actúa como si tuviera su propio procesador con identidad y registros separados
• El hardware de la GPGPU (GPU thread scheduler) define cuando un grupo de threads que pueden correr concurrentemente. También tiene la habilidad para intercambiar entre threads de forma rápida y transparente desde el punto de vista del software.
• Existen algunas instrucciones que permiten que los threads se comuniquen a través de espacios de memoria compartida de CUDA.
• Una configuración de ejecución define el número de threads que ejecutará el kernel y el arreglo del grid computacional en 1D, 2D o 3D. Esta información se encierra en triples brackets "<<<" y ">>>" después del nombre del kernel y antes de la lista de parámetros entre paréntesis.
• La región más grande en una GPGPU es llamada memoria global. 
• La memoria global está sujeta a reglas de fusión que combinan múltiples transacciones de memoria en una sola carga en funciones que pueden obtener una transferencia más alta desde y hacia la memoria.
• En general, el mayor rendimiento de la memoria ocurre cuando se accesa de manera en trozos de 128 bytes consecutivos.
• Existen otras formas de memoria accesible para el programador de una GPGPU, como son constant, cache, shared, local, texture y register memory.
• Aunque el ancho de banda de la memoria global es de 160-200 GB/s, es lento comparado con el rendimiento de 1 teraflop que puede alcanzar. Por lo tanto es esencial hacer reuso de datos para alcanzar el máximo rendimiento.

NVIDIA Tesla C2075

./deviceQuery de NVIDIA Tesla C2075


[cuda@scamlab release]$ ./deviceQuery
[deviceQuery] starting...

./deviceQuery Starting...

 CUDA Device Query (Runtime API) version (CUDART static linking)

Found 1 CUDA Capable device(s)

Device 0: "Tesla C2075"
  CUDA Driver Version / Runtime Version          4.2 / 4.2
  CUDA Capability Major/Minor version number:    2.0
  Total amount of global memory:                 5375 MBytes (5636292608 bytes)
  (14) Multiprocessors x ( 32) CUDA Cores/MP:    448 CUDA Cores
  GPU Clock rate:                                1147 MHz (1.15 GHz)
  Memory Clock rate:                             1566 Mhz
  Memory Bus Width:                              384-bit
  L2 Cache Size:                                 786432 bytes
  Max Texture Dimension Size (x,y,z)             1D=(65536), 2D=(65536,65535), 3D=(2048,2048,2048)
  Max Layered Texture Size (dim) x layers        1D=(16384) x 2048, 2D=(16384,16384) x 2048
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total number of registers available per block: 32768
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  1536
  Maximum number of threads per block:           1024
  Maximum sizes of each dimension of a block:    1024 x 1024 x 64
  Maximum sizes of each dimension of a grid:     65535 x 65535 x 65535
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and execution:                 Yes with 2 copy engine(s)
  Run time limit on kernels:                     Yes
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Concurrent kernel execution:                   Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support enabled:                Yes
  Device is using TCC driver mode:               No
  Device supports Unified Addressing (UVA):      Yes
  Device PCI Bus ID / PCI location ID:           3 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

CUDA Hello World

Programa Hello World.

Imprime "Hello" a partir de un vector, y realiza una suma de vectores para crear el mensaje "World!".



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

const int N = 16;
const int blocksize = 16;


/**
* Suma vector 'b' al 'a'
*/

__global__ void hello(char *a, int *b) {
    a[threadIdx.x] += b[threadIdx.x];
}

/**
*
*/
int main(int argc, char** argv) {

    char a[N] = "Hello \0\0\0\0\0\0";
    int b[N] = {15, 10, 6, 0, -11, 1, 0, 0, 0, 0, 0, 0, 0, 0, 0, 0};

    char *ad; // d for device variables
    int *bd;
    const int csize = N * sizeof (char);
    const int isize = N * sizeof (int);

    printf("%s", a);

    cudaMalloc((void**) &ad, csize);
    cudaMalloc((void**) &bd, isize);
    cudaMemcpy(ad, a, csize, cudaMemcpyHostToDevice);
    cudaMemcpy(bd, b, isize, cudaMemcpyHostToDevice);

    dim3 dimBlock(blocksize, 1);
    dim3 dimGrid(1, 1);
    hello << <dimGrid, dimBlock >> >(ad, bd);
    cudaMemcpy(a, ad, csize, cudaMemcpyDeviceToHost);
    cudaFree(ad);
    cudaFree(bd);

    printf("%s\n", a);
    return (EXIT_SUCCESS);
}

¿Cómo compilar CUDA desde consola?

Instrucciones para compilar un programa en CUDA desde consola en Linux

1. Crear archivo "file.cu"
2. Compilar
• nvcc -c file.cu
3. Linkear
• g++ -o nameExecutable -L/usr/local/cuda/lib64/ -lcudart -lcuda file.o
4. Ejecutar
• ./nameExecutable