CUDA

En 2003, un equipo de investigadores dirigido por Ian Buck presentó Brook, el primer modelo de programación ampliamente adoptado para extender C con construcciones paralelas de datos. Ian Buck luego se unió a NVIDIA y lideró el lanzamiento de CUDA en 2006, la primera solución del mundo para computación general en GPU. [6]

Ejemplo de flujo de procesamiento CUDA
1. Se copian los datos de la memoria principal a la memoria de la GPU
2. La CPU encarga el proceso a la GPU
3. La GPU lo ejecuta en paralelo en cada núcleo
4. Se copia el resultado de la memoria de la GPU a la memoria principal

CUDA presenta ciertas ventajas sobre otros tipos de computación sobre GPU utilizando APIs gráficas.

• Lecturas dispersas: se puede consultar cualquier posición de memoria.
• Memoria compartida: CUDA pone a disposición del programador un área de memoria de 16KB (o 48KB en la serie Fermi) que se compartirá entre hilos del mismo bloque. Dado su tamaño y rapidez puede ser utilizada como caché.
• Lecturas más rápidas de y hacia la GPU.
• Soporte para enteros y operadores a nivel de bit.

• No se pueden utilizar variables estáticas dentro de funciones o funciones con número de parámetros variable.
• No está soportado el renderizado de texturas.
• En precisión simple no soporta números desnormalizados o NaNs.
• Puede existir un cuello de botella entre la CPU y la GPU por los anchos de banda de los buses y sus latencias.
• Por razones de eficiencia, los threads o hilos de ejecución deben lanzarse en grupos de al menos 32, con miles de hilos en total.

CUDA intenta aprovechar el gran paralelismo, y el alto ancho de banda de la memoria en las GPU en aplicaciones con un gran coste aritmético frente a realizar numerosos accesos a memoria principal, lo que podría actuar de cuello de botella.

El modelo de programación de CUDA está diseñado para que se creen aplicaciones que de forma transparente escalen su paralelismo para poder incrementar el número de núcleos computacionales. Este diseño contiene tres puntos claves, que son la jerarquía de grupos de hilos, las memorias compartidas y las barreras de sincronización.

La estructura que se utiliza en este modelo está definido por un grid, dentro del cual hay bloques de hilos que están formados por como máximo 1024 hilos distintos.

Cada hilo en un bloque está identificado con un identificador único, que se accede con la variable threadIdx. Esta variable es muy útil para repartir el trabajo entre distintos hilos. threadIdx tiene 3 componentes (x, y, z), coincidiendo con las dimensiones de bloques de hilos. Así, cada elemento de una matriz, por ejemplo, lo podría tratar su homólogo en un bloque de hilos de dos dimensiones.

Al igual que los hilos, los bloques se identifican mediante blockIdx (en sus componentes x, y, z). Algunas otras funciones útiles son blockDim, para acceder al tamaño de bloque y gridDim, para acceder a la forma de la grid.

Un kernel es el código que se ejecuta en el dispositivo, la función que ejecutan los diferentes flujos durante la fase paralela. En CUDA un kernel se ejecuta mediante un conjunto de flujos, es decir, es una función la cual al ejecutarse lo hará en N distintos hilos en lugar de en secuencial. Se define incluyendo __global__ en la declaración. Por ejemplo:

//Definición del kernel
__global__ void f(int a, int b, int c)
{
}

Si nuestra función f queremos que calcule la diferencia entre dos vectores A y B y lo almacene en un tercero C:

__global__ void f(int* A, int* B, int* C)
{
	int i = threadIdx.x;
	C[i] = A[i] - B[i];
}

Esta función se ejecutaría una vez en cada hilo, reduciendo el tiempo total de ejecución en gran medida, y dividiendo su complejidad, O(n), por una constante directamente relacionada con el número de procesadores disponibles.

El mismo ejemplo con matrices sería:

__global__ void f(int** A, int** B, int** C)
{
	int i = threadIdx.x; //Columna del bloque que ocupa este determinado hilo
	int j= threadIdx.y; //Fila
	C[i][j] = A[i][j] - B[i][j];
}

Cuando se lanza un kernel se crea una malla de hilos donde todos ellos ejecutan el mismo programa. El kernel especifica las instrucciones que van a ser ejecutadas por cada hilo individual y se pueden lanzar todos los hilos en un único bloque, lanzar varios bloques con un solo hilo cada uno, o lanzar varios bloques con varios hilos en cada bloque. De esta forma, el GPU se encarga de ejecutar simultáneamente tantas copias del kernel como hilos se tenga en ejecución. A cada hilo se le asignan sus propios datos de forma que cada copia realice la misma operación pero con datos diferentes, aprovechando así el paralelismo que ofrecen las GPUs.

La memoria del procesador principal y del dispositivo son espacios de memoria completamente separados, lo que permite la computación simultánea tanto en la CPU como en la GPU sin competir por los recursos de memoria. Para ejecutar un kernel en el dispositivo GPU se siguen los siguientes pasos:

1. Reservar memoria en el dispositivo
2. Transferir los datos necesarios del procesador principal al espacio de memoria asignado al dispositivo.
3. Invocar la ejecución del kernel en cuestión.
4. Transferir los resultados del dispositivo al procesador principal y liberar la memoria del dispositivo una vez finalizada la ejecución del kernel.

En una llamada a un kernel, se le ha de pasar el tamaño de grid y de bloque, por ejemplo, en el main del ejemplo anterior podríamos añadir:

dim3 bloque(N,N); //Definimos un bloque de hilos de N*N
dim3 grid(M,M) //Grid de tamaño M*M

f<<<grid, bloque>>>(A, B, C);

En el momento que se invoque esta función, los bloques de un grid se enumerarán y distribuirán por los distintos multiprocesadores libres.

Como los distintos hilos colaboran entre ellos y pueden compartir datos, se requieren unas directivas de sincronización. En un kernel, se puede explicitar una barrera incluyendo una llamada a __syncthreads(), en la que todos los hilos se esperarán a que los demás lleguen a ese mismo punto.

Los hilos en CUDA pueden acceder a distintas memorias, unas compartidas y otras no.

• En primer lugar, está la memoria privada de cada hilo, solamente accesible desde él mismo.
• Cada bloque de hilos posee también un espacio de memoria, compartida en este caso por los hilos del bloque y con un ámbito de vida igual que el del propio bloque.
• Todos los hilos pueden acceder a una memoria global.

Además, existen otros dos espacios de memoria más, que son de solo lectura y accesibles por todos los hilos. Son la memoria constante y la de texturas. Todas las memorias de acceso global persisten mientras esté el kernel en ejecución.

Un multiprocesador contiene ocho procesadores escalares, dos unidades especiales para funciones trascendentales, una unidad multihilo de instrucciones y una memoria compartida. El multiprocesador crea y maneja los hilos sin ningún tipo de overhead por la planificación, lo cual unido a una rápida sincronización por barreras y una creación de hilos muy ligera, consigue que se pueda utilizar CUDA en problemas de muy baja granularidad, incluso asignando un hilo a un elemento por ejemplo de una imagen (un píxel).

Lista completa, con la versión soportada de CUDA: https://developer.nvidia.com/cuda-gpus

Visual Studio: CUDA proporciona un pequeño conjunto de extensiones para lenguajes de programación estándar, como C, que permiten una implementación sencilla de algoritmos paralelos. Con CUDA C/C++, los programadores pueden concentrarse en la tarea de paralelizar los algoritmos en lugar de dedicar tiempo a su implementación.

Para usar CUDA en su sistema, necesitará lo siguiente instalado:

• Una GPU compatible con CUDA
• Una versión compatible de Microsoft Windows
• Una versión compatible de Microsoft Visual
• Studio (2015, 2017, 2019, 2022)
• El kit de herramientas NVIDIA CUDA[7]

Los campos de utilización de CUDA son muy variados y van en aumento con el paso del tiempo. Entre los más populares se encuentran:

• Animación y modelado:

El motor gráfico Unreal Engine de Epic Games utiliza una GPU única y un único nodo. Usa CUDA para hacer un renderizado acelerado por GPU en OpenGL, DirectX y Vulkan. También para implementar Phys-X.

• Astronomía y Astrofísica:

La GPU-AH de Universidade do Porto utiliza una GPU única y un nodo único. Usa CUDA para calcular la densidad y la velocidad promedio de la red.

• Big Data y minería de datos:

La ADVANCE.AI utiliza multi GPU y multi nodos. Usa CUDA para el reconocimiento de documentos de identidad, incorporación digital de clientes, eliminar el fraude/robo de identidad y la verificación de identidad sin contacto. También los servicios de pagos vinculados.

• Bioinformática:

El DEEPCHAIN de InstaDeep Ltd. usa multi GPU pero un nodo único. Implementa CUDA para calcular la probabilidad de transición de cualquier mutación utilizando modelos de lenguaje de última generación, experimentos de mutagénesis in-silico a partir de archivos PDB de proteínas, simulaciones de dinámica molecular a gran escala y visualizaciones 3D.

• Modelado climático, meteorológico y oceánico:

COSMO de COSMO consortium utiliza multi GPU y multi nodos. Usan CUDA para todas las funciones en la rama MCH utilizadas para el pronóstico del tiempo operativo.[6]

• Sanders, Jason. (2016). CUDA by Example: An Introduction to General-Purpose GPU Programming. 2a ed. Upper Saddle River, New Jersey : Addison-Wesley.
• Iván Rivera; M. Vargas-Lombardo. (Diciembre de 2012). Principios y Campos de Aplicación en CUDA Programación paralela y sus potencialidades. Nexo. Revista Científica, 25, 39-46.
• Francesc Guim, Ivan Rodero. (2015). Arquitecturas basadas en computación gráfica (GPU). Recuperado el 15/11/2019, de Universitat Oberta de Catalunya Sitio web: https://www.exabyteinformatica.com/uoc/Informatica/Arquitecturas-de-computadores-avanzadas/Arquitecturas-de-computadores-avanzadas-(Modulo-5).pdf
• Misael Angeles Arce, Georgina Flores Becerra, and Antonio M. Vidal. (2016). Implementación CPU-GPU y comparativa de las bibliotecasBLAS-CUBLAS, LAPACK-CULA.. Recuperado el 15/11/2019, de Universidad Politécnica de Valencia Sitio web: https://riunet.upv.es/bitstream/handle/10251/11735/INCO2-2011-01.pdf
• César Represa Pérez José María Cámara Nebreda Pedro Luis Sánchez Ortega . (2016). UNIVERSIDAD DE BURGOS Área de Tecnología Electrónica INTRODUCCIÓN A LA PROGRAMACIÓN EN CUDA. Recuperado el 15/11/2019, de UNIVERSIDAD DE BURGOS Sitio web: https://riubu.ubu.es/bitstream/handle/10259/3933/Programacion-en-CUDA.pdf;jsessionid=4F61BF7B394035CC9E8553901D643C16?sequence=1

• OpenCL
• GPGPU
• GPU

• Zona CUDA de Nvidia
• Guía de programación para CUDA C (en inglés)