CUDA
# CUDA

CUDA son las siglas de Compute Unified Device Architecture (Arquitectura Unificada de Dispositivos de Cómputo) que hace referencia a una plataforma de computación en paralelo que incluye un compilador y un conjunto de herramientas de desarrollo creadas por Nvidia que permiten a los programadores usar una variación del lenguaje de programación C (CUDA C) para codificar algoritmos en GPU de Nvidia.

Por medio de wrappers se puede usar Python, Fortran, Julia y Java en vez de C/C++.

Funciona en todas las GPU Nvidia de la serie G8X en adelante, incluyendo GeForce, Quadro, ION y la línea Tesla. Nvidia afirma que los programas desarrollados para la serie GeForce 8 también funcionarán sin modificaciones en todas las futuras tarjetas Nvidia, gracias a la compatibilidad binaria provista por el conjunto de instrucciones PTX (Parallel Thread Execution).

CUDA intenta explotar las ventajas de las GPU frente a las CPU de propósito general utilizando el paralelismo que ofrecen sus múltiples núcleos, que permiten el lanzamiento de un altísimo número de hilos simultáneos. Por ello, si una aplicación está diseñada utilizando numerosos hilos que realizan tareas independientes (que es lo que hacen las GPU al procesar gráficos, su tarea natural), una GPU podrá ofrecer un gran rendimiento en campos que podrían ir desde la biología computacional a la criptografía, por ejemplo.

El primer SDK se publicó en febrero de 2007 en un principio para Windows, Linux, y más adelante en su versión 2.0 para macOS. Actualmente se ofrece para Windows XP/Vista/7/8/10, para Linux 32/64 bits y para macOS.

## Historia

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.

En 2006 Ian Buck se une a NVIDIA y lidera el lanzamiento de CUDA, la primera solución para el cómputo general en GPU.

En 2007 se hace público el SDK 1.0 de CUDA para Microsoft Windows y Linux. Soporta la micro-arquitectura Tesla.

En 2009 se hace público el SDK 1.0 de CUDA para Microsoft Windows y Linux. Soporta la micro-arquitectura Tesla. Desde entonces, el SDK se ha actualizado para soportar las nuevas arquitecturas que se desarrollan para las GPU.

Un aspecto importante en la historia de CUDA fue la relación que mantuvo con Apple, la cual estuvo marcada por varios conflictos y cambios en la dirección estratégica. Inicialmente, CUDA recibió soporte en macOS con su versión 2.0, la cual fue lanzada en el año 2007, pero este soporte terminó con la versión 10.12. La relación entre NVIDIA y Apple se vio afectada por múltiples desafíos, incluyendo problemas técnicos con chipsets gráficos y diferencias en la gestión de la memoria y las comunicaciones a través de la GPU. Estos problemas culminaron en la decisión de Apple de retirar el soporte de los controladores de GPU de NVIDIA en macOS Mojave y su posterior transición hacia su propia tecnología de gráficos, llamada “Metal”. Este cambio no solo mostró las tensiones entre ambas compañías, sino que también tuvo un impacto significativo en la comunidad de desarrolladores, especialmente aquellos que dependían de CUDA para aplicaciones de cálculo computacional intensivo. La falta de soporte de CUDA en los sistemas operativos de Apple marcó un cambio significativo en la industria, evidenciando las diferencias entre las estrategias de hardware y software de NVIDIA y Apple.

## Ventajas

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.

## Desventaja

* 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.

## El modelo CUDA

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.

## Kernel

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:

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

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:

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.

## Ejecución de un kernel

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:

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

## Invocaciones a un 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:

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.

## Sincronización

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.

## Compilación

Las funciones kernel se separan del código principal y son compiladas por los ensambladores y compiladores patentados por NVIDIA, mientras que el resto del código se compila con un compilador de C++ compatible con el Compilador Nvidia CUDA (NVCC), por ejemplo gcc, g++ o cl.exe. Al final, se incorporan las funciones de GPU compiladas al archivo principal del código.

## Jerarquía de memoria

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.

## Arquitectura CUDA

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).

## Diferencia entre una tarjeta gráfica y CUDA

En la arquitectura clásica de una tarjeta gráfica podemos encontrar la presencia de dos tipos de procesadores, los procesadores de vértices y los procesadores de fragmentos, dedicados a tareas distintas e independientes dentro del entorno gráfico, y con repertorios de instrucciones diferentes.

Esto presenta dos problemas importantes, por un lado el desequilibrio de carga que aparece entre ambos procesadores y por otro la diferencia entre sus respectivos repertorios de instrucciones. De este modo, la evolución natural en la arquitectura de una GPU ha sido la búsqueda de una arquitectura unificada donde no se distingue entre ambos tipos de procesadores. Así se llegó a la arquitectura CUDA, donde todos los núcleos de ejecución necesitan el mismo repertorio de instrucciones.

En la arquitectura CUDA cuenta con unas unidades de ejecución denominadas Streaming Multiprocessors (SM), 8 unidades en el ejemplo de la figura, que están interconectadas entre sí por una zona de memoria común. Cada SM está compuesto a su vez por unos núcleos de cómputo llamados “núcleos CUDA” que son los encargados de ejecutar las instrucciones y que en nuestro ejemplo vemos que hay 32 núcleos por cada SM, lo que hace un total de 256 núcleos de procesamiento. Este diseño hardware permite la programación sencilla de los núcleos de la GPU utilizando un lenguaje de alto nivel como puede ser el lenguaje C para CUDA. De este modo, el programador simplemente escribe un programa secuencial dentro del cual se llama a lo que se conoce como kernel, que puede ser una simple función o un programa completo. Este kernel se ejecuta de forma paralela dentro de la GPU como un conjunto de hilos y que el programador organiza dentro de una jerarquía en la que pueden agruparse en bloques, y que a su vez se pueden distribuir formando una malla.

Así pues, un bloque de hilos es un conjunto de hilos concurrentes que pueden cooperar entre ellos a través de mecanismos de sincronización y compartir accesos a un espacio de memoria exclusivo de cada bloque. Y una malla es un conjunto de bloques que pueden ser ejecutados independientemente y que por lo tanto pueden ser lanzados en paralelo en los SM.

Cuando se invoca un kernel, el programador especifica el número de hilos por bloque y el número de bloques que conforman la malla. Una vez en la GPU, a cada hilo se le asigna un ´único número de identificación dentro de su bloque, y cada bloque recibe un identificador dentro de la malla.

Esto permite que cada hilo decida sobre qué datos tiene que trabajar, lo que simplifica enormemente el direccionamiento de memoria cuando se trabaja con datos multidimensionales, como es el caso del procesado de imágenes o la resolución de ecuaciones diferenciales en dos y tres dimensiones. Otro aspecto a destacar en la arquitectura CUDA es la presencia de una unidad de distribución de trabajo que se encarga de distribuir los bloques entre los SM disponibles.

Los hilos dentro de cada bloque se ejecutan concurrentemente y cuando un bloque termina, la unidad de distribución lanza nuevos bloques sobre los SM libres. Los SM mapean cada hilo sobre un núcleo SP, y cada hilo se ejecuta de manera independiente con su propio contador de programa y registros de estado. Dado que cada hilo tiene asignados sus propios registros, no existe penalización por los cambios de contexto, pero en cambio sí existe un límite en el número máximo de hilos activos debido a que cada SM tiene un número determinado de registros.

Una característica particular de la arquitectura CUDA es la agrupación de los hilos en grupos de 32. Un grupo de 32 hilos recibe el nombre de warp, y se puede considerar como la unidad de ejecución en paralelo, ya que todos los hilos de un mismo warp se ejecutan físicamente en paralelo y por lo tanto comienzan en la misma instrucción. Así, cuando se selecciona un bloque para su ejecución dentro de un SM, el bloque se divide en warps, se selecciona uno que esté listo para ejecutarse y se emite la siguiente instrucción a todos los hilos que forman el warp. Dado que todos ellos ejecutan la misma instrucción al unísono, la máxima eficiencia se consigue cuando todos los hilos coinciden en su ruta de ejecución.

Aunque el programador puede ignorar este comportamiento, conviene tenerlo en cuenta si se pretende optimizar alguna aplicación. En cuanto a la memoria, durante su ejecución los hilos pueden acceder a los datos desde diferentes espacios dentro de una jerarquía de memoria. Así, cada hilo tiene una zona privada de memoria local y cada bloque tiene una zona de memoria compartida visible por todos los hilos del mismo bloque.

## Tarjetas Soportadas

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

| Nvidia GeForce GeForce RTX 3090 GeForce RTX 3080 GeForce RTX 3070 GeForce RTX TITAN GeForce RTX 2080 SUPER GeForce RTX 2080TI GeForce RTX 2080 GeForce RTX 2070 SUPER GeForce RTX 2070 GeForce RTX 2060 SUPER GeForce RTX 2060 GeForce GTX 1660 SUPER GeForce GTX 1660TI GeForce GTX 1660 GeForce GTX 1650 SUPER GeForce GTX 1650TI GeForce GTX 1650 GeForce GTX TITAN V GeForce GTX TITAN XP GeForce GTX 1080Ti GeForce GTX 1080 GeForce GTX 1070TI GeForce GTX 1070 GeForce GTX 1060 GeForce GTX 1050Ti GeForce GTX 1050 GeForce GTX TITAN X GeForce GTX 980Ti GeForce GTX 980 GeForce GTX 970 GeForce GTX 960 GeForce GTX 950 GeForce GTX TITAN Z GeForce GTX TITAN GeForce GTX 780Ti GeForce GTX 780 GeForce GTX 770 GeForce GTX 760 GeForce GTX 750Ti GeForce GTX 750 GeForce GTX 690 GeForce GTX 680 GeForce GTX 670 GeForce GTX 650 GeForce GT 640 GeForce GTX 590 GeForce GTX 580 GeForce GTX 570 GeForce GTX 560 GeForce GTX 550 GeForce GTX 480 GeForce GTX 470 GeForce GTX 465 GeForce GTX 460 GeForce GTS 450 GeForce GT 440 GeForce GT 430 GeForce GTX 295 GeForce GTX 285 GeForce GTX 280 GeForce GTX 275 GeForce GTX 260 GeForce GTS 250 GeForce GTS 240 GeForce GT 240 GeForce GT 220 GeForce 210/G210 GeForce 9800 GX2 GeForce 9800 GTX+ GeForce 9800 GTX GeForce 9800 GT GeForce 9600 GSO GeForce 9600 GT GeForce 9500 GT GeForce 9400 GT GeForce 9400 m GeForce 9300 m GeForce 9100 m GeForce 8800 Ultra GeForce 8800 GTX GeForce 8800 GTS GeForce 8800 GT GeForce 8800 GS GeForce 8600 GTS GeForce 8600 GT GeForce 8600 m GeForce 8500 GT GeForce 8400 GS GeForce 8300 m GeForce 8200 m GeForce 8100 m | Nvidia GeForce Mobile GeForce RTX 2080 MaxQ GeForce RTX 2070 MaxQ GeForce RTX 2060 MaxQ GeForce GTX 1660 GeForce GTX 1650 GeForce GTX 1080 MaxQ GeForce GTX 1070 MaxQ GeForce GTX 1060 MaxQ GeForce GTX 1050 MaxQ GeForce GTX 880M GeForce GTX 780M GeForce GTX 770M GeForce GTX 765M GeForce GTX 760M GeForce GT 750M GeForce GT 745M GeForce GT 740M GeForce GT 735M GeForce GT 730M GeForce GTX 680MX GeForce GTX 680M GeForce GTX 675MX GeForce GTX 675M GeForce GTX 670MX GeForce GTX 670M GeForce GTX 660M GeForce GT 650M GeForce GT 645M GeForce GT 640M GeForce GTX 580M GeForce GTX 570M GeForce GTX 560M GeForce GT 555M GeForce GT 550M GeForce GT 540M GeForce GT 525M GeForce GT 520M GeForce GTX 480M GeForce GTX 285M GeForce GTX 280M GeForce GTX 260M GeForce GTS 360M GeForce GTS 350M GeForce GTS 260M GeForce GTS 250M GeForce GT 620M GeForce GT 335M GeForce GT 330M GeForce GT 325M GeForce GT 320M GeForce 310M GeForce GT 240M GeForce GT 230M GeForce GT 220M GeForce G210M GeForce GTS 160M GeForce GTS 150M GeForce GT 130M GeForce GT 120M GeForce G110M GeForce G105M GeForce G103M GeForce G102M GeForce G100 GeForce 9800M GTX GeForce 9800M GTS GeForce 9800M GT GeForce 9800M GS GeForce 9700M GTS GeForce 9700M GT GeForce 9650M GT GeForce 9650M GS GeForce 9600M GT GeForce 9600M GS GeForce 9500M GS GeForce 9500M G GeForce 9400M G GeForce 9300M GS GeForce 9300M G GeForce 9200M GS GeForce 9100M G GeForce 8800M GTX GeForce 8800M GTS GeForce 8700M GT GeForce 8600M GT GeForce 8600M GS GeForce 8400M GT GeForce 8400M GS GeForce 8400M G GeForce 8200M G | Nvidia Quadro Quadro RTX 8000 Quadro RTX 6000 Quadro RTX 5000 Quadro RTX 4000 Quadro GP100 Quadro GV100 Quadro P6000 Quadro P5000 Quadro P4000 Quadro P2200 Quadro P2000 Quadro P1000 Quadro P620 Quadro P600 Quadro 6000 Quadro 5000 Quadro 4000 Quadro FX 5800 Quadro FX 5600 Quadro FX 4800 Quadro FX 4700 X2 Quadro FX 4600 Quadro FX 3800 Quadro FX 3700 Quadro FX 1800 Quadro FX 1700 Quadro FX 580 Quadro FX 570 Quadro FX 380 Quadro FX 370 Quadro NVS 450 Quadro NVS 420 Quadro NVS 295 Quadro NVS 290 Quadro Plex 1000 Model IV Quadro Plex 1000 Model S4 Nvidia Quadro Mobile Quadro FX 3800M Quadro FX 3700M Quadro FX 3600M Quadro FX 2800M Quadro FX 2700M Quadro FX 1800M Quadro FX 1700M Quadro FX 1600M Quadro FX 880M Quadro FX 770M Quadro FX 570M Quadro FX 380M Quadro FX 370M Quadro FX 360M Quadro NVS 320M Quadro NVS 160M Quadro NVS 150M Quadro NVS 140M Quadro NVS 135M Quadro NVS 130M Nvidia Tesla Tesla C2050 Tesla S1070 Tesla M1060 Tesla C1060 Tesla C870 Tesla D870 Tesla S870 |
| - | - | - |

## IDEs

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

## Campos de utilización

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.

* Síntesis de voz:

Con CUDA podemos utilizar la librería TTS (Text to Speech) de COQUI que implementa herramientas para medir la calidad del conjunto de datos de modelos ya preentrenados y utilizarlos en la síntesis de voz. Actualmente se usa en más de 20 idiomas para productos y proyectos de investigación.

## Compatibilidad

El Driver de CUDA mantiene la compatibilidad con versiones anteriores para soportar aplicaciones creadas en Toolkits más antiguos. Con una versión de Driver secundaria compatible, las aplicaciones creadas en CUDA Toolkit 11 y posteriores son compatibles con cualquier controlador dentro de la versión principal correspondiente.

Con el paquete CUDA Forward Compatibility, los administradores del sistema pueden ejecutar aplicaciones creadas con un conjunto de herramientas más nuevo incluso cuando se instala en el sistema un Driver más antiguo que no cumple con la versión mínima requerida. Esta compatibilidad “hacia adelante” permite que las implementaciones de CUDA se beneficien de la cadencia de lanzamiento más rápida y de las funciones y el rendimiento más recientes del CUDA Toolkit.

## Bibliografía

* 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.
* Misael Angeles Arce, Georgina Flores Becerra, Antonio M. Vidal (2016). Implementación CPU-GPU y comparativa de las bibliotecas BLAS-CUBLAS, LAPACK-CULA. Recuperado el 15/11/2019, de Universidad Politécnica de Valencia.
* 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.
