Capítulo 2: Modelos de Programación en GPU
Las Unidades de Procesamiento Gráfico (GPUs) han evolucionado de aceleradores gráficos de funcionalidad fija a motores de cálculo altamente paralelos y programables, capaces de acelerar una amplia gama de aplicaciones. Para permitir que los programadores aprovechen efectivamente el paralelismo masivo de las GPUs, se han desarrollado varios modelos y APIs de programación paralela, como NVIDIA CUDA, OpenCL y DirectCompute. Estos modelos de programación proporcionan abstracciones que permiten a los programadores expresar el paralelismo en sus aplicaciones, al mismo tiempo que ocultan los detalles de bajo nivel del hardware de la GPU.
En este capítulo, exploraremos los conceptos y principios clave detrás de los modelos de programación paralela para GPUs, centrándonos en el modelo de ejecución, las arquitecturas de conjunto de instrucciones (ISA) de GPU, las ISA de GPU de NVIDIA y la ISA de Graphics Core Next (GCN) de AMD. También proporcionaremos ejemplos para ilustrar cómo se aplican estos conceptos en la práctica.
Modelo de Ejecución
El modelo de ejecución de los modernos modelos de programación de GPU se basa en el concepto de kernels, que son funciones que se ejecutan en paralelo por un gran número de hilos en la GPU. Al lanzar un kernel, el programador especifica el número de hilos a crear y cómo se organizan en una jerarquía de grillas, bloques (o matrices de hilos cooperativos - CTAs) e hilos individuales.
- Una grilla representa todo el espacio del problema y consta de uno o más bloques.
- Un bloque es un grupo de hilos que pueden cooperar y sincronizarse entre sí a través de la memoria compartida y las barreras. Los hilos dentro de un bloque se ejecutan en el mismo núcleo de GPU (llamado multiprocesador de flujo o unidad de cómputo).
- Cada hilo tiene un ID único dentro de su bloque y grilla, que se puede usar para calcular direcciones de memoria y tomar decisiones de flujo de control.
Esta organización jerárquica permite a los programadores expresar tanto el paralelismo de datos (donde la misma operación se aplica a múltiples elementos de datos) como el paralelismo de tareas (donde se ejecutan diferentes tareas en paralelo).
FiguraA continuación se presenta la traducción al español del archivo Markdown proporcionado, manteniendo el código sin traducir y solo traduciendo los comentarios:
e 2.1 ilustra la jerarquía de hilos en el modelo de ejecución de la GPU.
Grid
________________
/ / / / /
/ / / / /
/ / / / /
/ / / / /
/__/__/__/__/__/
| | | |
| | Bloque |
| | | |
Hilo Hilo ...Figura 2.1: Jerarquía de hilos en el modelo de ejecución de la GPU.
Ejecución SIMT
Los modelos de programación de GPU como CUDA y OpenCL siguen un modelo de ejecución de Una Instrucción, Múltiples Hilos (SIMT). En el modelo SIMT, los hilos se ejecutan en grupos llamados warps (terminología de NVIDIA) o wavefronts (terminología de AMD). Todos los hilos dentro de un warp ejecutan la misma instrucción al mismo tiempo, pero cada hilo opera con diferentes datos.
Sin embargo, a diferencia del modelo tradicional de Una Instrucción, Múltiples Datos (SIMD), donde todos los elementos de procesamiento se ejecutan en sincronía, SIMT permite que los hilos tengan rutas de ejecución independientes y divergir en las instrucciones de ramificación. Cuando un warp encuentra una instrucción de ramificación, el hardware de la GPU evalúa la condición de ramificación para cada hilo en el warp. Si todos los hilos siguen la misma ruta (convergente), el warp continúa la ejecución normalmente. Si algunos hilos siguen diferentes rutas (divergente), el warp se divide en dos o más subwarps, cada uno siguiendo una ruta diferente. El hardware de la GPU serializa la ejecución de las rutas divergentes, enmascarando los hilos inactivos en cada subwarp. Cuando todas las rutas se completan, los subwarps se reconvergen y continúan la ejecución en sincronía.
La Figura 2.2 ilustra la ejecución SIMT con un flujo de control divergente.
Warp
________________
/ / / / /
/ / / / /
/ / / / /
| | |
| Ramificación |
| | |
/ \ / \ / \
/ X \ \
/ / \ \ \
/ \ \
/ \ \
/ \ \
/ \ \
\
\
\
ReconvergenciaFigura 2.2: Ejecución SIMT con flujo de control divergente.
Este mecanismo de manejo de divergencia permite que SIMT admita un flujo de control más flexibleAquí está la traducción al español:
Los programadores deben esforzarse por minimizar la divergencia dentro de un warp para lograr un rendimiento óptimo.
Jerarquía de memoria
Las GPU tienen una jerarquía de memoria compleja para admitir los requisitos de alto ancho de banda y baja latencia de las cargas de trabajo paralelas. La jerarquía de memoria suele constar de:
- Memoria global: el espacio de memoria más grande pero más lento, al que pueden acceder todos los hilos de un kernel. La memoria global suele implementarse mediante memoria GDDR o HBM de alto ancho de banda.
- Memoria compartida: un espacio de memoria rápido y en chip compartido por todos los hilos de un bloque. La memoria compartida se utiliza para la comunicación entre hilos y el intercambio de datos dentro de un bloque.
- Memoria constante: un espacio de memoria de solo lectura utilizado para difundir datos de solo lectura a todos los hilos.
- Memoria de textura: un espacio de memoria de solo lectura optimizado para la localidad espacial y al que se accede a través de las memorias caché de textura. La memoria de textura se utiliza más comúnmente en cargas de trabajo gráficas.
- Memoria local: un espacio de memoria privado para cada hilo, utilizado para el desbordamiento de registros y estructuras de datos grandes. La memoria local suele asignarse a la memoria global.
La utilización eficaz de la jerarquía de memoria es crucial para lograr un alto rendimiento en las GPU. Los programadores deben tratar de maximizar el uso de la memoria compartida y minimizar los accesos a la memoria global para reducir la latencia de la memoria y los cuellos de botella del ancho de banda.
La Figura 2.3 ilustra la jerarquía de memoria de la GPU.
____________
| |
| Global |
| Memory |
____________
|
____________
| |
| Constant |
| Memory |
____________
|
____________
| |
| Texture |
| Memory |
____________
|
|
____________
| |
| Shared |
| Memory |
____________
|
____________
| |
| Local |
| Memory |
____________FigAquí está la traducción al español del archivo markdown, con los comentarios de código traducidos:
Arquitecturas de Conjunto de Instrucciones de GPU
Las arquitecturas de conjunto de instrucciones de GPU (ISA, por sus siglas en inglés) definen la interfaz de bajo nivel entre el software y el hardware. Ellas especifican las instrucciones, los registros y los modos de direccionamiento de memoria soportados por la GPU. Comprender las ISA de GPU es esencial para desarrollar código de GPU eficiente y optimizar el rendimiento.
En esta sección, exploraremos las ISA de dos principales proveedores de GPU: NVIDIA y AMD. Nos enfocaremos en las ISA Parallel Thread Execution (PTX) y SASS de NVIDIA, y la ISA Graphics Core Next (GCN) de AMD.
ISA de GPU de NVIDIA
Las GPU de NVIDIA soportan dos niveles de ISA: PTX (Parallel Thread Execution) y SASS (Streaming ASSembler). PTX es una ISA virtual que proporciona un objetivo estable para los compiladores de CUDA, mientras que SASS es la ISA nativa de las GPU de NVIDIA.
PTX (Parallel Thread Execution)
PTX es una ISA virtual de bajo nivel diseñada para las GPU de NVIDIA. Es similar a LLVM IR o Java bytecode en el sentido de que proporciona un objetivo estable e independiente de la arquitectura para los compiladores. Los programas de CUDA suelen compilarse a código PTX, que luego es traducido a las instrucciones SASS nativas por el controlador de la GPU de NVIDIA.
PTX soporta una amplia gama de instrucciones aritméticas, de memoria y de control de flujo. Tiene un número ilimitado de registros virtuales y soporta la predicción, lo que permite una implementación eficiente del control de flujo. PTX también proporciona instrucciones especiales para la sincronización de hilos, las operaciones atómicas y el muestreo de texturas.
Aquí hay un ejemplo de código PTX para un kernel de adición de vectores simple:
.version 7.0
.target sm_70
.address_size 64
.visible .entry vecAdd(
.param .u64 vecAdd_param_0,
.param .u64 vecAdd_param_1,
.param .u64 vecAdd_param_2,
.param .u32 vecAdd_param_3
)
{
.reg .b32 %r<4>;
.reg .b64 %rd<6>;
// Cargar los parámetros de entrada
ld.param.u64 %rd1, [vecAdd_param_0];
ld.param.u64 %rd2, [vecAdd_param_1];
ld.param.u64 %rd3, [vecAdd_param_2];
ld.param.u32 %r1, [vecAdd_param_3];
// Convertir los punteros a direcciones globales
cvta.to.global.u64 %rd4, %rd1;
cvta
```Aquí está la traducción al español del archivo Markdown, con los comentarios traducidos y el código sin traducir:
.to.global.u64 %rd5, %rd2; mov.u32 %r2, %tid.x; mul.wide.u32 %rd6, %r2, 4; add.s64 %rd7, %rd4, %rd6; add.s64 %rd8, %rd5, %rd6;
ld.global.f32 %f1, [%rd7]; ld.global.f32 %f2, [%rd8]; add.f32 %f3, %f1, %f2;
cvta.to.global.u64 %rd9, %rd3; add.s64 %rd10, %rd9, %rd6; st.global.f32 [%rd10], %f3;
ret; }
Este código PTX define una función de kernel `vecAdd` que toma cuatro parámetros: punteros a los vectores de entrada y salida, y el tamaño de los vectores. El kernel calcula el ID de hilo global, carga los elementos correspondientes de los vectores de entrada, realiza la adición y almacena el resultado en el vector de salida.
#### SASS (Streaming ASSembler)
SASS es el ISA nativo de las GPU de NVIDIA. Es un ISA de bajo nivel y específico de la máquina que se asigna directamente al hardware de la GPU. Las instrucciones SASS se generan mediante el controlador de la GPU de NVIDIA a partir del código PTX y no suelen ser visibles para los programadores.
Las instrucciones SASS se codifican en un formato compacto para reducir el ancho de banda de memoria y la huella de la caché de instrucciones. Admiten una amplia gama de tipos de operandos, incluidos registros, valores inmediatos y varios modos de direccionamiento para el acceso a la memoria.
Aquí hay un ejemplo de código SASS para el kernel de adición de vectores:
```sass
code_version_number 90
@P0 LDG.E R2, [R8];
@P0 LDG.E R0, [R10];
@P0 FADD R0, R2, R0;
@P0 STG.E [R12], R0;
EXIT;Este código SASS se corresponde con el código PTX mostrado anteriormente. Carga los elementos del vector de entrada desde la memoria global (LDG.E), realiza la adición (FADD), almacena el resultado de vuelta en la memoria global (STG.E) y sale del kernel.
AMD Graphics Core Next ISA
Las GPU de AMD utilizan la arquitectura y el ISA Graphics Core Next (GCN). GCN es un ISA basado en RISC que admite cargas de trabajo tanto gráficas como de cálculo. Está diseñado para un alto rendimiento, escalabilidad y eficiencia energética.
GCN introduce varias características clave, como:
- Un Aquí está la traducción al español de este archivo Markdown, con los comentarios traducidos, pero sin traducir el código:
alar ALU para una ejecución eficiente de operaciones escalares y control de flujo.
- Una ALU vectorial para la ejecución en paralelo de operaciones de paralelismo de datos.
- Un sistema de memoria de gran ancho de banda con soporte para operaciones atómicas y acceso de baja latencia a la memoria compartida.
- Un modo de direccionamiento flexible para operaciones de memoria, que admite direccionamiento base+desplazamiento y direccionamiento escalar+vectorial.
Aquí hay un ejemplo de código de la ISA de GCN para un kernel de adición de vectores:
.text
.globl vecAdd
.p2align 2
.type vecAdd,@function
vecAdd:
.set DPTR, 0
# Cargar los argumentos del kernel desde la memoria
s_load_dwordx4 s[0:3], s[4:5], 0x0
s_load_dword s4, s[4:5], 0x10
s_waitcnt lgkmcnt(0)
# Preparar los operandos de la adición
v_lshlrev_b32 v0, 2, v0
v_add_u32 v1, vcc, s1, v0
v_mov_b32 v3, s3
v_addc_u32 v2, vcc, s2, v3, vcc
flat_load_dword v1, v[1:2]
v_add_u32 v3, vcc, s0, v0
v_mov_b32 v5, s3
v_addc_u32 v4, vcc, s2, v5, vcc
flat_load_dword v0, v[3:4]
# Realizar la adición de los vectores
v_add_f32 v0, v0, v1
flat_store_dword v[3:4], v0
s_endpgmEste código de GCN carga los elementos de los vectores de entrada usando flat_load_dword, realiza la adición usando v_add_f32 y almacena el resultado de vuelta en la memoria usando flat_store_dword. Las instrucciones s_load_dwordx4 y s_load_dword se utilizan para cargar los argumentos del kernel desde la memoria.
Asignación de Algoritmos a Arquitecturas GPU
Asignar eficientemente los algoritmos a la arquitectura GPU es crucial para lograr un alto rendimiento. Las consideraciones clave incluyen:
Exponer Suficiente Paralelismo
El algoritmo debe descomponerse en muchos hilos de grano fino que puedan ejecutarse simultáneamente para aprovechar al máximo las capacidades de procesamiento paralelo de la GPU. A menudo, esto implica identificar las porciones de paralelismo de datos del algoritmo que se pueden ejecutar de forma independiente en diferentes elementos de datos.
Minimizar la Divergencia de Ramificación
El flujo de control divergente dentro de una warp/wavefront puede conducir a la serialización y a una eficiencia SIMD reducida. Los algoritmos deben estructurarse para minimizar la divergencia de ramificación cuando sea posible. Esto se puede lograr reduciendo el uso de flujos de control dependientes de los datos.Aquí está la traducción al español del archivo Markdown, con los comentarios del código traducidos:
Explotando la jerarquía de memoria
Acceder a la memoria global es costoso. Los algoritmos deben maximizar el uso de la memoria compartida y los registros para reducir los accesos a la memoria global. Los datos también deben organizarse en la memoria para permitir accesos de memoria coalescentes, donde los hilos en una warp accedan a ubicaciones de memoria contiguas. El uso eficaz de la jerarquía de memoria puede reducir significativamente la latencia de la memoria y los cuellos de botella de ancho de banda.
Equilibrando cálculo y accesos a la memoria
Los algoritmos deben tener una alta relación de operaciones aritméticas a operaciones de memoria para ocultar eficazmente la latencia de la memoria y lograr un alto rendimiento computacional. Esto se puede lograr maximizando la reutilización de datos, la búsqueda anticipada de datos y superponiendo el cálculo con los accesos a la memoria.
Minimizando las transferencias de datos entre host y dispositivo
Transferir datos entre la memoria del host (CPU) y del dispositivo (GPU) es lento. Los algoritmos deben minimizar dichas transferencias realizando la mayor cantidad de cálculos posible en la GPU. Los datos deben transferirse a la GPU en lotes grandes y mantenerse en el dispositivo durante el tiempo necesario para amortizar el costo de la transferencia.
Se utilizan varios patrones de diseño de algoritmos paralelos comúnmente cuando se desarrollan kernels de GPU:
- Mapeo: Cada hilo realiza la misma operación en un elemento de datos diferente, lo que permite un procesamiento paralelo sencillo de grandes conjuntos de datos.
- Reducción: La reducción paralela se usa para calcular de manera eficiente un solo valor (por ejemplo, suma, máximo) a partir de un gran conjunto de datos de entrada. Los hilos realizan reducciones locales, que luego se combinan para producir el resultado final.
- Exploración: También conocida como suma de prefijos, la exploración se usa para calcular la suma acumulada de los elementos de una matriz. Los algoritmos eficientes de exploración paralela son componentes clave para muchas aplicaciones aceleradas por GPU.
- Plantilla: Cada hilo calcula un valor en función de los elementos de datos vecinos. Los cálculos de plantilla son comunes en simulaciones científicas y aplicaciones de procesamiento de imágenes.Here is the Spanish translation of the provided markdown file, with the code comments translated but the code itself left unchanged:
Aplicaciones
- Gather/Scatter: Los hilos leen de (gather) o escriben en (scatter) ubicaciones arbitrarias en la memoria global. Se requiere un diseño de datos y patrones de acceso cuidadosos para una eficiencia.
La Figura 3.20 ilustra un ejemplo del patrón de asignación, donde cada hilo aplica una función (por ejemplo, raíz cuadrada) a un elemento diferente del array de entrada.
Array de entrada:
| | | | | | | |
v v v v v v v v
______________________________
Hilos: | | | | | | | |
|____|____|____|____|____|____|____|
| | | | | | |
v v v v v v v
Array de salida: Figura 3.20: Ejemplo del patrón de asignación en la programación de GPU.
Conclusión
Los modelos de programación de GPU como CUDA y OpenCL exponen las capacidades de procesamiento paralelo de las GPU modernas a los desarrolladores, permitiéndoles acelerar una amplia gama de aplicaciones. Estos modelos de programación proporcionan abstracciones que permiten que las cargas de trabajo paralelas de grano fino se asignen de manera eficiente al hardware de la GPU.
Comprender el modelo de ejecución, la jerarquía de memoria y los primitivos de sincronización proporcionados por estos modelos de programación es esencial para escribir código de GPU de alto rendimiento. Los desarrolladores deben considerar cuidadosamente factores como la organización de hilos, la divergencia de ramas, los patrones de acceso a la memoria y el diseño de algoritmos para aprovechar al máximo el poder computacional de las GPU.
A medida que las arquitecturas de GPU continúan evolucionando, los modelos y las herramientas de programación también deben avanzar para permitir que los desarrolladores utilicen de manera efectiva las nuevas características y capacidades del hardware. La investigación continua en áreas como el diseño de lenguajes de programación, la optimización del compilador y el autoajuste será crucial para mejorar la productividad de los programadores y la portabilidad del rendimiento en la era de la computación heterogénea.