Evaluating Application Characteristics for GPU Portability Layer Selection

Este artículo presenta un estudio que analiza aplicaciones heterogéneas representativas de los principales experimentos de Física de Altas Energías para identificar las características clave que influyen en el rendimiento a través de diversas capas de portabilidad de GPU, guiando así a los desarrolladores en la selección de la tecnología más adecuada para sus necesidades específicas.

Lo esencial

Imagina que eres un chef intentando cocinar un banquete masivo. Tienes tres tipos diferentes de hornos de alta potencia en tu cocina: uno fabricado por NVIDIA, uno por AMD y uno por Intel. Cada horno cocina la comida de manera diferente, utiliza diferentes perillas y requiere diferentes recetas para funcionar de la mejor manera.

Si escribes tu receta específicamente para el horno NVIDIA (usando un lenguaje llamado CUDA), no puedes simplemente meter esa misma receta en los hornos de AMD o Intel. Tendrías que reescribir todo el proceso. Esto es un problema porque no siempre sabes qué horno tendrás en tu cocina mañana.

Para resolver esto, el artículo trata sobre las "capas de portabilidad". Piensa en ellas como traductores universales o adaptadores inteligentes. Permiten escribir una receta maestra que el traductor convierte al lenguaje específico que cada horno entiende. El artículo analiza varios de estos traductres (como Kokpos, SYCL, OpenMP y Alpaka) para ver cuál es el que mejor se adapta a diferentes tipos de tareas de cocina.

Aquí está lo que los autores descubrieron cuando probaron estos traductores con "recetas" reales de experimentos de física de altas energías (como los que se usan para estudiar partículas subatómicas):

1. El problema del "Tiempo de Arranque"

Encender una GPU (un horno) no es instantáneo. Toma unos milisegundos despertarse y estar listo.

• El Problema: Algunos traductores son lentos para iniciar el proceso de cocción. Por ejemplo, Kokkos puede añadir un retraso significativo al usar hornos AMD. Si tu tarea de cocina es muy corta (como hervir un huevo durante 10 segundos), y el traductor tarda 5 segundos solo en encender la estufa, habrás desperdiciado la mitad de tu tiempo.
• La Lección: Si tus tareas son diminutas y rápidas, evita los traductores que hagan que el arranque sea lento.

2. El problema de la "Cocina Atestada"

En un laboratorio de física real, la GPU no trabaja sola. Es parte de un sistema más grande donde muchas personas (hilos/threads) intentan usar el horno al mismo tiempo.

• El Problema: Algunos traductores son malos manejando multitudes. Kokkos, por ejemplo, tiene una regla que dice: "Solo una persona puede hablar con el horno a la vez", lo que causa un embotellamiento si múltiples chefs intentan lanzar tareas simultáneamente. SYCL es un poco inconsistente; a veces deja que todos cocinen a la vez y otras veces los obliga a esperar en fila, dependiendo de qué versión del traductor estés usando.
• La Lección: Si tu aplicación necesita que muchas personas trabajen a la vez, necesitas un traductor que sepa gestionar una cocina concurrida sin bloquear las puertas.

3. El problema de la "Compatibilidad de Herramientas"

Las recetas de física suelen utilizar herramientas especiales (librerías como ROOT o Eigen) que ayudan con las matemáticas y los datos.

• El Problema: Algunas de estas herramientas no se llevan bien con los traductores. Por ejemplo, una herramienta matemática muy popular llamada Eigen a menudo falla cuando se usa con el compilador de NVIDIA, el cual muchos traductores utilizan. Además, intentar usar dos compiladores diferentes (uno para la CPU y otro para la GPU) en el mismo proyecto es como intentar construir una casa con dos juegos de planos diferentes que no coinciden: hace que la construcción (la creación del software) sea una pesadilla.
• La Lección: Antes de elegir un traductor, verifica si tus herramientas favoritas encajarán dentro de él.

4. El problema del "Acomodo de los Muebles"

A las GPUs les encantan los diseños simples y planos. Prefieren que los datos estén dispuestos como una fila ordenada de cajas. Sin embargo, los datos de física suelen venir en formas complejas y desordenadas (como una pila de maletas de diferentes tamaños).

• El Problema: Los traductores intentan arreglar este desorden envolviendo los datos en contenedores especiales. Aunque esto hace que el código sea portátil, añade "sobrecarga" (overhead)—como poner cada objeto en una maleta antes de moverlo, incluso si solo necesitas mover un calcetín. Esto ralentiza las cosas. Además, ninguno de los traductores es muy bueno manejando datos "dentados" (filas de diferentes longitudes), lo cual es muy común en la física.
• La Lección: Si tus datos son complejos y desordenados, el traductor podría ralentizarte intentando poner orden.

5. El problema de las "Herramientas Especializadas"

A veces necesitas una herramienta específica, como un Generador de Números Aleatorios (RNG) o una Transformada Rápida de Fourier (FFT).

• El Proble El: Cada fabricante de hornos tiene sus propias versiones especializadas y súper rápidas de estas herramientas. Los traductores universales a menudo no incluyen estas versiones especializadas, o utilizan sus propias versiones más lentas. Aunque puedes forzar al traductor a usar la herramienta nativa del horno, esto rompe la "portabilidad" porque esa herramienta solo funciona en ese horno específico.
• La Lección: Si dependes fuertemente de estas herramientas específicas, es posible que tengas que elegir entre velocidad (usar la herramienta nativa del horno) o portabilidad (usar la herramienta genérica del traductor).

6. Los problemas de "Tiempo de Construcción" y "Día de Mudanza"

• Construir la Receta: Algunos traductores hacen que el "tiempo de cocción" (tiempo de compilación) sea mucho más largo. Para proyectos enormes, usar ciertos traductores puede hacer que el proceso de construcción tarde horas en lugar de minutos.
• Mover la Cocina: Si construyes tu software para un horno específico (por ejemplo, un NVIDIA V100), podría no funcionar en uno más nuevo (un NVIDIA A100). Algunos traductores requieren que construyas una versión separada para cada tipo de horno que puedas encontrar. Esto crea un dolor de cabeza logístico masivo para distribuir el software a diferentes laboratorios.

El Veredicto Final

El artículo concluye que no existe un traductor "perfecto".

• Kokkos es excelente para algunas cosas pero tiene dificultades con la concurrencia y los tiempos de arranque en cierto hardware.
• SYCL es potente pero puede ser inconsistente dependiendo de la versión del compilador.
• OpenMP y otros tienen sus propias fortalezas y debilidades con respecto a cómo manejan la memoria y el hardware diferente.

La Conclusión: No puedes elegir un traductor solo porque es popular. Tienes que mirar tu "receta" específica (tu aplicación). Si tu código es corto y rápido, elige un traductor con un bajo tiempo de arranque. Si tu código es complejo y utiliza muchas herramientas, elige uno que juegue bien con esas herramientas.

Los autores también señalan que estas tecnologías evolucionan rápidamente, como si salieran nuevos modelos de hornos cada año. Lo que funciona mejor hoy podría cambiar mañana, por lo que los desarrolladores deben seguir vigilando el panorama. En el futuro, nuevos estándares podrían facilitar estas elecciones, pero por ahora, realizar pruebas cuidadosas es la única forma de encontrar el ajuste adecuado.

Resumen técnico

Resumen Técnico: Evaluación de las Características de las Aplicaciones para la Selección de Capas de Portabilidad de GPU

Declaración del Problema

La computación de la Física de Altas Energías (HEP) depende cada vez más de las GPU para obtener rendimiento, pero el panorama del hardware se está diversificando más allá de NVIDIA para incluir también a AMD e Intel. Si bien los lenguajes específicos de cada proveedor, como CUDA, ofrecen un alto rendimiento, restringen el despliegue a hardware específico, lo que reduce la portabilidad. Por el contrario, las capas de portabilidad como Kokkos, SYCL, HIP, OpenMP, Alpaka y std::par prometen la ejecución en diversas arquitecturas de CPU y GPU. Sin embargo, estas capas no son funcionalmente equivalentes; exhiben características distintas en cuanto a rendimiento, soporte de concurrencia y compatibilidad con bases de código existentes. El problema central abordado es la falta de un marco claro para que los desarrolladores seleccionen la capa de portabilidad más apropiada basándose en las características específicas de la aplicación, la estructura de los kernels y los entornos de ejecución.

Metodología

El Centro para la Excelencia en Computación de la Física de Altas Energías (HEP-CCE) llevó a cabo un estudio comparativo utilizando aplicaciones heterogéneas representativas de los principales experimentos de HEP:

• ATLAS: FastCaloSim (simulación de calorímetro parametrizado).
• CMS: Patatrack (reconstrucción de píxeles heterogénea con más de 40 kernels, multihilo y pools de memoria) y P2R (reconstrucción de trazas de un solo kernel).
• DUNE: Wire-Cell Toolkit (rasterización, scatter-add y kernels de convolución FFT).
• Generadores de Eventos: Puertos a GPU de Sherpa (Pepper) y Madgraph (utilizando código autogenerado).

Estos bancos de pruebas fueron portados de C++ serial a diversas tecnologías de portabilidad para identificar las características del código que impactan significamente en el rendimiento y la compatibilidad. El estudio evaluó estas aplicaciones a través de múltiples dimensiones, incluyendo la latencia de lanzamiento de kernels, el manejo de la concurrencia, la compatibilidad de librerías, la gestión de memoria y la sobrecarga de compilación.

Hallazgos Clave y Resultados

1. Latencia de Lanzamiento de Kernels

El lanzamiento de un kernel no es instantáneo y puede oscilar entre microsegundos y decenas de microsegundos.

• Kokkos: Se observó que añade una sobrecarga significativa (decenas de microseg segundos), particularmente en GPUs AMD, que ya poseen latencias de lanzamiento nativas más altas que NVIDIA o Intel.
• SYCL, HIP, OpenMP, Alpaka: No exhibieron penalizaciones adicionales grandes en comparación con las implementaciones nativas.
• Especificidades de AMD: Llamar a rocRAND en GPUs AMD resultó en latencias extremas en la primera llamada (superando 1 segundo), lo que sugiere que los kernels de ejecución corta pueden sufrir una degradación severa de rendimiento si utilizan capas de portabilidad que exacerban la sobrecarga de lanzamiento.

2. Concurrencia y Pools de Hilos

Las aplicaciones de HEP suelen ejecutarse dentro de marcos de trabajo (frameworks) más grandes que utilizan multihilo (por ejemplo, Intel TBB).

• Kokkos: Demostró incompatibilidades significativas con pools de hilos externos. Su backend serial utiliza bloqueos que serializan las llamadas, y su backend de hilos prohíbe las llamadas de hilos externos. Los lanzamientos de kernels concurrentes actualmente solo son viables con backends CUDA/HIP utilizando APIs específicas de la arquitectura, lo que limita la portabilidad.
• SYCL: Las implementaciones fueron inconsistentes. Algunas serializaron llamadas concurrentes desde diferentes hilos dependiendo del fabricante del compilador y la versión. Las arquitecturas AMD parecían carecer de una solución SYCL concurrente al momento del estudio.

3. Compatibilidad de Librerías Externas y Compiladores

• Eigen: Muchas versiones son incompatibles con nvcc de NVIDIA, lo que afecta a las capas de portabilidad que dependen de este. También es actualmente incompatible con OpenMP offload.
• ROOT: No puede compilarse completamente con nvc++ de NVIDIA. La compilación híbrida (usando g++ y nvc++) es posible pero complica las reglas de construcción y requiere una gestión cuidadosa de la asignación de memoria (los datos de la GPU deben ser asignados por nvc++).
• Complejidad de Construcción (Build): Soportar múltiples capas de portabilidad en un mismo proyecto requiere reglas de construcción complejas (uso de #ifdef, separación de archivos) para evitar conflictos de cabeceras entre capas como Kokkos y CUDA.

4. Estructuras de Datos y Transferencias de Memoria

• Estructuras Orientadas a Objetos vs. Planas: Las estructuras de datos complejas y orientadas a objetos, comunes en HEP, se mapean mal a las GPUs, que prefieren layouts planos o Estructuras de Arrays (SoA).
• Sobrecarga de Portabilidad: Capas como Kokkos y SYCL proporcionan abstracciones (por ejemplo, Kokkos::Views, sycl::buffers) que aumentan la sobrecarga de asignación y transferencia, especialmente para objetos pequeños.
• Inicialización: La inicialización de memoria por defecto en Kokkos::Views puede ser un cuello de botella de rendimiento; esto puede mitigarse usando Kokkos::ViewAllocateWithoutInitializing.
• Memoria Compartida Unificada (USM): Se observó que las transferencias automáticas de USM (usadas por std::par) son más lentas que las transferencias explícitas y pueden no activarse en el momento óptimo.
• Arrays Dentados (Jagged Arrays): Ninguna de las APIs soporta nativamente vectores multidimensionales de tamaños variables (jagged arrays), lo que requiere la creación manual, reducción de dimensiones o padding.

5. RNGs, FFTs y Atómicos

• Librerías Nativas: Las librerías específicas de cada proveedor (cuRAND, rocRAND, oneMKL) ofrecen un rendimiento optimizado pero carecen de portabilidad. Aunque Kokkos proporciona implementaciones de RNG/FFT multiplataforma, estas pueden no alcanzar el rendimiento de las nativas.
• Contribuciones del HEP-CCE: El grupo desarrolló una librería sin cabeceras (header-only) que envuelve los RNGs nativos del dispositivo con una API uniforme y backends nativos de hardware para oneMKL, permitiendo que SYCL llame a implementaciones nativas a través de distintos proveedores.
• Atómicos: Las operaciones atómicas se mapean mal al hardware de GPU. El rendimiento varía drásticamente según el compilador y el hardware (por ejemplo, atomic<int> en NVIDIA con OpenMP/clang13 fue 35 veces más lento que CUDA). Los atómicos de enteros generalmente rinden menos que los de punto flotante/doble.

6. Tiempo de Compilación

• Sobrecarga del Compilador: nvc++ es generalmente de 2 a 3 veces más lento que g++. Kokkos añade una sobrecarga del 10% al 20%.
• Código con Exceso de Plantillas (Templates): Madgraph, que utiliza código autogenerado fuertemente basado en plantillas, experimentó tiempos de compilación con el backend de SYCL icpx que fueron más de 1000 veces más lentos que otros backends para kernels grandes.

7. Aprovisionamiento en Tiempo de Ejecución (Runtime Provisioning)

• Kokkos: Requiere construirse con identificadores de hardware específicos. Un binario construido para una GPU (ej. V100) puede no ejecutarse en otra (ej. A100). Soportar librerías compartidas y múltiples arquitecturas crea una matriz compleja de configuraciones de construcción, lo que desafía la distribución de binarios.
• Alpaka: Requiere compilación separada para cada compilador de backend.
• SYCL/OneAPI: Los binarios multiplataforma están emergiendo, pero a menudo están limitados a arquitecturas específicas o requieren identificación en tiempo de compilación.
• OpenMP: Algunas implementaciones (LLVM Clang) permiten listar múltiples arquitecturas en tiempo de compilación y seleccionar dinámicamente el backend en tiempo de ejecución.

Significado y Conclusión

El artículo concluye que ninguna solución de portabilidad de GPU es óptima para todos los escenarios. La elección de una capa de portabilidad debe ser impulsada por un análisis cuidadoso de las características específicas de la aplicación, incluyendo el tiempo de ejecución de los kernels, los requisitos de concurrencia, la complejidad de las estructuras de datos y las dependencias de librerías externas.

Los autores enfatizan que el ecosistema está evolucionando rápidamente. Tanto las capas de portabilidad (Kokkos, SYCL, etc.) como los compiladores nativos (clang, nvcc, hipcc) están mejorando, con tendencias hacia un mayor soporte de características y rendimiento optimizado. El artículo aboga por el monitoreo frecuente de estas tecnologías para asegurar el mejor ajuste entre las bases de código y las capas de portabilidad. Además, el documento señala que los estándares emergentes de C++ (como P2300 para ejecución asíncrona) podrían simplificar estas elecciones en el futuro, aunque el soporte de hardware generalizado aún tardará años.

En última instancia, seleccionar una capa de portabilidad sin comprender plenamente la naturaleza del código, sus dependencias externas y su entorno de ejecución conlleva riesgos significativos tanto en rendimiento como en la flexibilidad de despliegue.