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.
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.
No hay comentarios:
Publicar un comentario