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

Empezando a "cocinar"

GPU. Graphics Processing Unit

De forma simple, una unidad de procesamiento gráfico es un chip de procesamiento dedicado para realizar operaciones de construcción de imágenes para salida de video en una computadora.

Una GPU puede aplicar instrucciones en una misma operación sobre un cojunto de datos, en este caso, y debido a la naturaleza de las aplicaciones gráficas, la mayor parte de ellas involucran operaciones sobre matrices y vectores. 

Una GPU está formada por varios núcleos pequeños y diseñados de forma eficiente para realizar estos cálculos en forma paralela. Las características propias de las GPUs debido a su estructura de programación altamente en paralelo, y la incorporación de múltiples núcleos y memoria RAM dedicada, se han utilizado, y utilizarán para algo más que desarrollo de procesamiento gráfico.

GPGU. General Purpose GPU

Este acrónimo hace referencia al uso de una GPU como una alternativa a un microprocesador "común", usando su poder de procesamiento en paralelo, para su aplicación en otras tareas de propósito general, y que no involucren precisamente trabajos sobre gráficos.

El cálculo de operaciones en la GPU ofrece un rendimiento a nivel de teraflops, al descargar en las GPU las partes de la aplicación que requieren gran capacidad computacional, mientras que el resto del código sigue ejecutándose en la CPU.

Esto consituye una combinación de CPU  + GPU, usando la CPU formada por varios núcleos para el procesamiento en serie, mientras que la GPU consta también de varios núcleos optimizados para procesamiento en paralelo. 

Las partes en serie del código se ejecutan en la CPU, mientras que las paralelas en la GPU

NVIDIA CUDA. Compute Unified Device Architecture

Los primeros usos de las GPU como GPGPU, en los años 1999 y 2000, precisaban el uso de APIs de programació de gráficos como OpenGL y Cg para programar las GPU. Debido a esa limitación, NVIDIA ofreció un ambiente de programación transparente para desarolladores con lenguajes familiares como C, C++ y Fortran.

El primer SDK se publicó en el año 2007, y actualmente se puede instalar en plataformas Linux, Mac o Windows. La versión más nueva del SDK es Cuda 5, publicada en octubre del 2012.

Links interesantes

• Nvidia Cuda Zone