No es “Speed Up” todo lo que reluce

Cuando realizamos optimizaciones en nuestro código, debemos ser conscientes de las limitaciones que impone la arquitectura, en muchos casos lo que puede parecer una mejora acaba siendo un cuello de botella.

 Aunque la técnica de “unroll” o desenrollado de bucles descrita en el post anterior parezca la “Panacea[1]”, hay que ser cautelosos, porque no siempre será así. Es necesario hacer un estudio detallado de cada caso para aplicar factores de desenrollado apropiados a cada uno de los algoritmos que se pretenden computar, ya que un desenrollado de bucles agresivo, puede aumentar demasiado las tareas de planificación (scheduling) y aumentar demasiado el uso de registro en el cuerpo del bucle, que en las CPU puede implicar una migración de los valores de los registros a memoria y esto puede desembocar en una degradación del rendimiento del programa. Por otro lado, si se utilizan factores de desenrollado muy altos, el tamaño del cuerpo del bucle resultante podría desbordar la caché de instrucciones, seguido de pérdidas de caché, y en consecuencia reduciéndose el rendimiento.

En programación GPGPU nos encontramos además con una serie de restricciones que limitan el rendimiento de los programas. Antes de proceder a la implementación deberemos estudiar el impacto que supone el desenrollado de bucles para los programas GPGPU en el contexto de las restricciones de esos recursos y hallar el factor de “unroll” óptimo para ese programa y dispositivo.

ILP y Factor de Ocupación (OF): como ya se ha visto, el nivel de paralelismo por instrucción juega un papel importante a la hora de aumentar el rendimiento, ya que el tamaño de la caché de instrucciones “I-cache” es limitado. Es importante no sobrecargar el cuerpo de la función con demasiadas instrucciones que provoquen un fallo de caché y la consiguiente penalización en el rendimiento.  Por otro lado en las GPU el factor de ocupación puede ser igual o más importante que el anterior a la hora de ocultar las latencias de acceso a memoria y de flujo de instrucciones. Debido a la interrelación que hay entre ILP y OF, se hace muy complicado estimar los valores óptimos. A continuación se muestra un caso concreto que sirve de ejemplo.

En el apartado B de esta sección se expuso brevemente el modelo de programación de CUDA compuesto por Mallas (Grids), Bloques (Blocks) e Hilos (Threads), ahora vamos a compararlo con la arquitectura de las GPUs para comprender mejor el problema del factor de ocupación. Una GPU está compuesta por x SM “stream multiprocessor”, cada uno dispone de c SP “stream processor”, en cada uno hay w Warps[2], que a su vez pueden ejecutar t Threads. La clave para comprender el problema del factor de ocupación reside en la diferencia real que existe entre la capacidad computacional de la arquitectura que utilicemos y la capacidad que solicitamos en la llamada al kernel. La Figura 1 muestra las capacidades para distintos dispositivos con arquitectura CUDA.

Device Name	GeForce GTX 480	Tesla C2050	GeForce GTX 295
Compute capability	2.0	2.0	1.3
Memory Informat.
Total global mem	1.536 MB	2.687 MB	896 MB
Total constant Mem	64 Kb	64 Kb	64 Kb
Max mem pitch	2048 MB	2048 MB	2048 MB
MP Information
Multiprocessor count	15	14	30
Shared mem per mp	48	48	16
Registers per mp	32768	32768	16384
Threads in warp	32	32	32
Warps/MP	48	48	32
Blocks/MP	8	8	8
Max threads/block	1024	1024	512
Max thread dim	(1024,1024, 64)	(1024,1024, 64)	(512, 512, 64)
Max grid dim	(65535, 65535, 65535)	(65535,  65535, 65535)	(65535, 65535, 1)
Warps/block in SP	6	6	4
threads in block	192	192	128
total pararell threads	2880	2688	3840

Figura 1. Capacidades de algunos dispositivos

De todos los valores de la tabla ahora mismo nos interesan: “Multiprocessor count”, que es el número de multiprocesadores que tiene el dispositivo, “Blocks/MP” bloques por multiprocesador, “”Warps/MP” que son agrupaciones de hilos dentro de cada bloque y “Threads in warp” que son el número de hilos que ejecutan en cada warp. Sin tener nada más en cuenta, podemos calcular el número total de hilos que cada dispositivo es capaz de ejecutar de forma concurrente: por ejemplo para el dispositivo GeForce GTX 480 tenemos 15 procesadores x 8 bloques x 6 warps x 48 hilos, en total 34.560 hilos concurrentes ejecutando el mismo kernel. Evidentemente la arquitectura nos está imponiendo además ciertas restricciones: “Max threads per block” está indicando un número (1024) mayor al teórico. Si en un bloque hay 6 warps y 32 hilos (192 hilos en total) ¿cómo es posible alcanzar esa cifra? CUDA puede utilizar hasta 8 bloques por procesador, es decir, que si utilizamos 192 hilos por bloque, dispondremos de 8 bloques por procesador, pero si necesitamos 256 hilos, dispondremos de 6 bloques, para 512 hilos solo 3 bloques, para 1024 solo un bloque por procesador. Cuando se solicitan los recursos en las llamadas a los kernels, CUDA realiza una segmentación y planificación de ejecución de todos los bloques solicitados, ahí es donde se pueden producir holguras si no hemos diseñado bien nuestra llamada, reduciéndose el factor de ocupación.

Otra cuestión a considerar, también muy importante es el uso de los registros que necesita nuestro Kernel, a mayor uso de registros, menor disponibilidad de bloques por procesador. Estas dos variables son las que van a modificar el factor de ocupación que, como hemos podido inferir, es el aprovechamiento físico de la arquitectura para cada tanda de instrucciones. El primero se maneja afinando los parámetros de llamada al núcleo y el segundo, mediante la creación de kernels lo más sencillos que se pueda (sin perder de vista el tiempo de latencia). Podemos saber el uso de registros que hace nuestro kernel añadiendo el parámetro (31) – Xptxas –v en la línea de compilación. 

nvcc -Xptxas –v ejemplo.cu ptxas info   : Compiling entry function 'acos_main' ptxas info   : Used 4 registers, 60+56 bytes lmem, 44+40 bytes smem,  20 bytes cmem[1],  12 bytes cmem[14]

Figura 2. Obteniendo información de nuestro kernel

---

[1] De panacea universal: Remedio que buscaban los antiguos alquimistas para curar todas las enfermedades

[2] El término “Warp” puede asimilarse al de “tanda” ya que hace referencia al conjunto de hilos de un procesador que se ejecutan en el mismo instante por pertenecer a la misma programación del scheduler. 

Unroll y Vencerás

Este, es el primero de una serie de 5 post en los cuales quiero presentar las claves básicas para desarrollar kernels CUDA de calidad.

Desde hace décadas, uno de los caballos de batalla en el mundo de los compiladores para CPU es el análisis de bucles y la reducción de iteraciones con el objetivo de optimizar el uso de recursos y mejorar los tiempos de ejecución. El resultado suele ser un bucle de menor número de repeticiones con instrucciones paralelas que son independientes unas de otras, en muchos casos esta simple acción puede reducir drásticamente el tiempo de procesamiento. El número de instrucciones serializadas, dentro del nuevo bucle coincide con el factor de desenrollado “unroll factor”. Este valor es clave a la hora de afinar y sintonizar los mejores valores para la paralelización como veremos más adelante. La Figura 1 muestra un ejemplo sencillo de esta técnica para UF = 4.

/* Antes del desenrollado */ for (i = 0; i < N; ++i) { c[i] = a[i] + b[i]; }

/* Despues del desenrollado */ for (i = 0; i < N - (4 - 1); i += 4) { c[i] = a[i] + b[i]; c[i+1] = a[i+1] + b[i+1]; c[i+2] = a[i+2] + b[i+2]; c[i+3] = a[i+3] + b[i+3]; } /* Bucle para el resto */ for (; i < N; ++i) { c[i] = a[i] + b[i]; }

                                                    Figura 1. Ejemplo de "unroll"

Los beneficios inmediatos de este proceso son: 

1. Se reduce el contador de instrucciones dinámicas, se reducen el número de comparaciones y operaciones ramificadas, para la misma carga de trabajo. 

2. Aumentan las posibles combinaciones para la planificación, debido a que surgen más instrucciones independientes, aumentando así el nivel de paralelización por instrucción (ILP). 

3. Aumenta la oportunidad para explotar la herencia de los registros y memoria local cuando los bucles externos se desenrollan y los internos se fusionan en el mismo juego de instrucciones iguales para todos.

En las arquitecturas GPGPU esta técnica también es habitualmente usada, aunque con otro enfoque, la diferencia de arquitecturas y modelos de programación necesita de una nueva forma de contemplar el impacto del desenrollado de bucles en el rendimiento en los programas para GPGPU.

A diferencia de los compiladores para CPU, no es tan fácil poder realizar optimizaciones automáticas de código usando esta técnica, por lo que el programador es el que deberá diseñar un buen código paralelizado en GPU. El modelo de programación CUDA y las restricciones de los dispositivos nos dan la guía para realizar la paralelización de bucles, como no es el objetivo de este post exponer los pormenores de la arquitectura y su forma de uso, sólo haremos una breve reseña con objeto de aportar claridad a esta exposición, para ampliar conceptos puede referirse el lector a las guías que acompañan el toolkit de CUDA.

 Básicamente podemos decir que un Kernel[1] equivale a una iteración del cuerpo de un bucle (en programación secuencial). La GPU tiene la habilidad de ejecutar de forma paralela el mismo kernel tantas veces como se defina en su llamada. Cada kernel corre en un hilo, y existen unas variables locales a cada uno de los kernels, pudiendo así cada kernel disponer de la información del índice o contador de iteración equivalente en caso de los bucles secuenciales. Los hilos (Thread), están organizados en bloques (Block), y los bloques en mallas (Grid). Los hilos pueden llegar a tener hasta un máximo de 3 dimensiones, y los bloques y las mallas hasta 2 o 3 dependiendo de la arquitectura. La siguiente línea de código muestra como un kernel puede obtener su identificador único (utilizando una dimensión en cada nivel):

 tid = threadIdx.x + blockIdx.x * blockDim.x;

Cómo pueden compartir y acceder a la memoria cada uno de los hilos y su ámbito se describe en el post "Los diversos usos de la memoria".


 El tamaño de la malla a utilizar y del bloque de hilos, es una elección que se deja al programador en función del problema a resolver. Esta decisión tendrá impacto en el rendimiento final del programa, que como veremos a continuación, está ligada al problema y sujeta a las restricciones que impone la arquitectura GPU y los recursos propios del dispositivo que se utilice.

---

[1] El Kernel es una porción de código que es invocada desde un proceso en CPU (host) y se ejecuta en un dispositivo GPU.

Estado del Arte III (CUDA)

CUDA[1] es una arquitectura de cálculo paralelo de nVidia que aprovecha la gran potencia de la GPU (unidad de procesamiento gráfico) para proporcionar un incremento extraordinario del rendimiento del sistema. CUDA ha sido recibida con entusiasmo por la comunidad científica. Por ejemplo, se está utilizando para acelerar AMBER, un simulador de dinámica molecular empleado por más de 60.000 investigadores del ámbito académico y farmacéutico de todo el mundo para acelerar el descubrimiento de nuevos medicamentos. En el mercado financiero, Numerix y CompatibL introdujeron soporte de CUDA para una nueva aplicación de cálculo de riesgo de contraparte y, como resultado, se ha multiplicado por 18 la velocidad de la aplicación.

CUDA Toolkit es el conjunto de herramientas y librerías que comprenden el entorno software de la arquitectura. Los componentes básicos son: El compilador C/C++, herramientas para depurar, librerías matemáticas optimizadas para GPU, Primitivas de rendimiento optimizado para GPU, todo ello acompañado de una extensa documentación y abundantes ejemplos. El paquete está disponible para los sistemas operativos más habituales: Windows, Linux y MacOS. Además el desarrollador cuenta con un enorme repertorio de servicios en línea, Webminars, whitepapers y numerosas publicaciones electrónicas y como no, los foros de la comunidad de desarrolladores.

Versiones del toolkit: desde 2006 hasta la fecha han aparecido varias versiones de CUDA que han supuesto un avance en cada una de ellas, la Figura siguiente muestra dicha evolución a lo largo del tiempo.

Evolución de CUDA

En mayo de 2011 se ha publicado la versión 4.0 del toolkit, el cual supone grandes avances en la facilidad de uso y programación, entre otras destacan las siguientes novedades: Se simplifica la portabilidad de aplicaciones: se pueden compartir múltiples GPUs a través de múltiples hilos, se pueden utilizar todas las GPUs de forma concurrente desde el mismo hilo del host. Reserva de memoria de host más rápida, implementación de las funciones C++ new/delete y soporte para funciones virtuales, soporte para código ensamblador inline PTX, librerías para la gestión de vectores y matrices optimizadas, y más.


El siguiente post, tratará de argumentar la elección de CUDA como plataforma de desarrollo frente a otras alternativas existentes.

---

[1] son las siglas de Compute Unified Device Architecture

Bienvenidos a desarrollo CUDA

Bienvenidos!

Con este blog quiero compartir mi experiencia paralelizando algoritmos mediante el uso de GPU's. Trataré de presentar aportes cortos que permitan obtener soluciones concretas.

 La computación sobre Unidades de Procesamiento Gráfico (GPU), es sin lugar a dudas, uno de los avances tecnológicos más importantes de los últimos años, que aporta un nuevo paradigma en el campo de la computación en paralelo y la computación de altas prestaciones. Este tipo de hardware (GPU), inicialmente diseñado para representar imágenes en el ordenador, tiene unas características peculiares que ofrecen una capacidad de cómputo inusitada. A diferencia de los actuales procesadores centrales (CPU) en los cuales podemos encontrar configuraciones de unos cuantos procesadores con 2, 4 u 8 núcleos de proceso, las GPU ofrecen configuraciones de  hasta 30 procesadores con hasta 8 núcleos de procesamientos cada uno, permitiendo ejecutar hasta 1536 hilos cada procesador.

Este potencial inicialmente era explotado por librerías gráficas de renderizado y trazadores de rayos para obtener imágenes realistas aplicando geometría computacional. El objetivo fundamental era obtener ciertos los valores para cada punto (en el plano o en el espacio para sistemas tridimensionales) como son: visible o no visible, color y profundidad de color. Así, con objeto de producir imágenes en movimiento se diseñaron sistemas que permitiesen realizar el cálculo de esos valores para cada punto de forma paralela y obtener el resultado de la imagen en tiempos tales que permitan la producción de una ilusión óptica de movimiento al presentar imágenes con una frecuencia mínima de 24 imágenes por segundo. Se trataba de poder calcular todos los valores de la matriz de puntos que forman la imagen en paralelo y de forma independiente del procesador central el cual estaría realizando otro tipo de tareas. Este conjunto de técnicas fueron denominados como Aceleración Gráfica por Hardware. Para poder explotar estas capacidades del nuevo hardware, deberá haber asociado a éste, un controlador que permita un repertorio específico de instrucciones, que serán llamadas utilizando un determinado API[1]. Términos tales como OpenGL® o DitectX® empezaron a sonar como librerías gráficas que eran capaces de explotar las nuevas capacidades de las GPUs.

Con el paso del tiempo cada vez fueron mayores las prestaciones de estas GPUs, aumentando tanto la capacidad de memoria como de los procesadores. Esto permitió aumentar el tipo de operaciones a realizar en dichos dispositivos de modo que además de poder realizar operaciones de dibujo, también se empezaron a realizar operaciones de codificación y decodificación, lo que permitió liberar a los procesadores centrales de tales tareas.

El empleo de procesadores gráficos para la resolución de problemas, comienza en el mundo académico en un principio dando forma gráfica al problema a solucionar, ya que las primitivas y repertorio de instrucciones con el que se contaba eran muy básicos y demasiado orientados al mundo de los gráficos. Actualmente se ha conseguido evolucionar estas librerías para proporcionar soporte a sistemas de desarrollo de propósito general que permiten aplicar con gran éxito la computación sobre GPU a cualquier tipo de problema sea o no de tipo gráfico. 

---

[1] Application Programming Interface