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.