Evaluating Application Characteristics for GPU Portability Layer Selection

Questo articolo presenta uno studio che analizza applicazioni eterogenee rappresentative di importanti esperimenti di Alta Energia per identificare le caratteristiche chiave che influenzano le prestazioni attraverso vari livelli di portabilità GPU, guidando così gli sviluppatori nella selezione della tecnologia più adatta alle loro specifiche esigenze.

L'essenziale

Immagina di essere uno chef che cerca di cucinare un banchetto enorme. Hai tre diversi tipi di forni ad alta potenza nella tua cucina: uno prodotto da NVIDIA, uno da AMD e uno da Intel. Ognuno di questi forni cucina il cibo in modo diverso, usa manopole diverse e richiede ricette differenti per funzionare al meglio.

Se scrivi la tua ricetta specificamente per il forno NVIDIA (usando un linguaggio chiamato CUDA), non puoi semplicemente mettere quella stessa ricetta nei forni AMD o Intel. Dovresti riscrivere tutto il pezzo. Questo è un problema perché non sempre sai quale forno avrai nella tua cucina domani.

Per risolvere questo problema, il documento parla di "livelli di portabilità". Immaginali come traduttori universali o adattatori intelligenti. Ti permettono di scrivere una ricetta maestra che il traduttore converte nel linguaggio specifico che ogni forno comprende. Il documento esamina diversi di questi traduttori (come Kokkos, SYCL, OpenMP e Alpaka) per vedere quale sia il più adatto per diversi tipi di compiti culinari.

Ecco cosa hanno scoperto gli autori quando hanno testato questi traduttori con vere "ricette" provenienti da esperimenti di fisica delle alte energie (come quelli usati per studiare le particelle subatomiche):

1. Il problema del "Tempo di Avvio"

Accendere una GPU (il forno) non è istantaneo. Richiede alcuni millisecondi per risvegliarsi e prepararsi.

• Il Problelo: Alcuni traduttori sono lenti a iniziare il processo di cottura. Ad esempio, Kokkos può aggiungere un ritardo significativo quando si utilizzano forni AMD. Se il tuo compito culinario è molto breve (come bollire un uovo per 10 secondi), e il traduttore impiega 5 secondi solo per accendere il fornello, hai sprecato metà del tuo tempo.
• La Lezione: Se i tuoi compiti sono minuscoli e rapidi, evita i traduttori che rendono lento l'avvio.

2. Il problema della "Cucina Affollata"

In un vero laboratorio di fisica, la GPU non lavora da sola. Fa parte di un sistema più grande dove molte persone (thread) cercano di usare il forno contemporaneamente.

• Il Problema: Alcuni traduttori sono scarsi nel gestire la folla. Kokkos, ad esempio, ha una regola che dice: "Solo una persona può parlare con il forno alla volta", il che causa un ingorgo se più chef provano a lanciare compiti simultaneamente. SYCL è un po' incoerente; a volte permette a tutti di cucinare insieme, e a volte li costringe ad aspettare in fila, a seconda di quale versione del traduttore stai usando.
• La Lezione: Se la tua applicazione ha bisogno di molte persone che lavorano contemporaneamente, hai bisogno di un traduttore che sappia gestire una cucina affollata senza bloccare le porte.

3. Il problema della "Compatibilità degli Strumenti"

Le ricette di fisica spesso usano strumenti speciali (librerie come ROOT o Eigen) che aiutano con la matematica e i dati.

• Il Problema: Alcuni di questi strumenti non vanno d'accordo con i traduttori. Ad esempio, uno strumento matematico molto popolare chiamato Eigen spesso si rompe quando viene usato con il compilatore NVIDIA, su cui molti traduttori fanno affidamento. Inoltre, cercare di usare due compilatori diversi (uno per la CPU e uno per la GPU) nello stesso progetto è come cercare di costruire una casa con due diversi set di progetti che non corrispondono: rende la costruzione (la creazione del software) un incubo.
• La Lezione: Prima di scegliere un traduttore, controlla se i tuoi strumenti preferiti ci staranno dentro.

4. Il problema della "Disposizione dei Mobili"

Le GPU amano layout semplici e piatti. Preferiscono che i dati siano disposti come una fila ordinata di scatole. Tuttavia, i dati della fisica spesso arrivano in forme complesse e disordinate (come un mucchio di valigie di diverse dimensioni).

• Il Problema: I traduttori cercano di correggere questo disordine avvolgendo i dati in contenitori speciali. Sebbene ciò renda il codice portabile, aggiunge "overhead" — come mettere ogni singolo oggetto in una valigia prima di spostarlo, anche se devi spostare solo un calzino. Questo rallenta le cose. Inoltre, nessuno dei traduttori è bravo a gestire i dati "frastagliati" (righe di lunghezze diverse), che sono molto comuni nella fisica.
• ** La Lezione:** Se i tuoi dati sono complessi e disordinati, il traduttore potrebbe rallentarti cercando di sistemarli.

5. Il problema degli "Strumenti Specializzati"

A volte serve uno strumento specifico, come un Generatore di Numeri Casuali (RNG) o una Trasformata di Fourier Veloce (FFT).

• Il Problema: Ogni produttore di forni ha le proprie versioni specializzate e super veloci di questi strumenti. I traduttori universali spesso non includono queste versioni specializzate, o usano le proprie versioni più lente. Sebbene sia possibile forzare il traduttore a usare lo strumento nativo del forno, questo rompe la "portabilità" perché tale strumento funziona solo su quel forno specifico.
• La Lezione: Se fai affidamento pesantemente su questi strumenti specifici, potresti dover scegliere tra velocità (usando lo strumento nativo del forno) o portabilità (usando lo strumento generico del traduttore).

6. I problemi del "Tempo di Costruzione" e del "Giorno del Trasloco"

• Costruire la Ricetta: Alcuni traduttori rendono molto più lungo il "tempo di cottura" (tempo di compilazione). Per progetti enormi, l'uso di certi traduttori può far sì che il processo di build duri ore invece di minuti.
• Traslocare la Cucina: Se costruisci il tuo software per un forno specifico (ad esempio, un NVIDIA V100), potrebbe non funzionare su uno più recente (un NVIDIA A100). Alcuni traduttori richiedono di costruire una versione separata per ogni singolo tipo di forno che potresti incontrare. Questo crea un enorme incubo logistico per la distribuzione del software ai diversi laboratori.

Il Verdetto Finale

Il documento conclude che non esiste un traduttore "perfetto".

• Kokkos è ottimo per alcune cose ma fatica con la concorrenza e i tempi di avvio su determinati hardware.
• SYCL è potente ma può essere incoerente a seconda della versione del compilatore.
• OpenMP e altri hanno i propri punti di forza e di debolezza riguardo alla gestione della memoria e di hardware diversi.

Il Punto Chiave: Non puoi scegliere un traduttore solo perché è popolare. Devi guardare la tua specifica "ricetta" (la tua applicazione). Se il tuo codice è breve e veloce, scegli un traduttore con un basso tempo di avvio. Se il tuo codice è complesso e usa molti strumenti, scegline uno che funzioni bene con quegli strumenti.

Gli autori notano anche che queste tecnologie si stanno evolvendo rapidamente, come se uscissero nuovi modelli di forni ogni anno. Ciò che funziona meglio oggi potrebbe cambiare domani, quindi gli sviluppatori devono continuare a monitorare il panorama. In futuro, nuovi standard potrebbero rendere queste scelte più facili, ma per ora, un test accurato è l'unico modo per trovare la combinazione giusta.

Sommario tecnico

Riepilogo Tecnico: Valutazione delle Caratteristiche delle Applicazioni per la Selezione dello Strato di Portabilità GPU

Definizione del Problema

L'informatica per l'Alta Energia (HEP) si sta affidando sempre più alle GPU per le prestazioni, ma il panorama dell'hardware si sta diversificando oltre NVIDIA per includere anche AMD e Intel. Sebbene i linguaggi specifici dei vendor come CUDA offrano alte prestazioni, essi limitano l'implementazione a hardware specifici, riducendo la portabilità. Al contrario, gli strati di portabilità come Kokkos, SYCL, HIP, OpenMP, Alpaka e std::par promettono l'esecuzione su diverse architetture CPU e GPU. Tuttavia, questi strati non sono funzionalmente equivalenti; essi presentano caratteristiche distinte per quanto riguarda prestazioni, supporto alla concorrenza e compatibilità con le basi di codice esistenti. Il problema centrale affrontato è la mancanza di un quadro chiaro per permettere agli sviluppatori di selezionare lo strato di portabilità più appropriato in base alle specifiche caratteristiche dell'applicazione, alla struttura dei kernel e agli ambienti di runtime.

Metodologia

L'High Energy Physics Center for Computational Excellence (HEP-CCE) ha condotto uno studio comparativo utilizzando applicazioni eterogenee rappresentative dei principali esperimenti HEP:

• ATLAS: FastCaloSim (simulazione del calorimetro parametrizzato).
• CMS: Patatrack (ricostruzione dei pixel eterogenea con oltre 40 kernel, multithreading e pool di memoria) e P2R (ricostruzione della traccia a singolo kernel).
• DUNE: Wire-Cell Toolkit (rasterizzazione, scatter-add e kernel di convoluzione FFT).
• Generatori di Eventi: Porting GPU di Sherpa (Pepper) e Madgraph (utilizzando codice auto-generato).

Questi testbed sono stati portati da C++ seriale a varie tecnologie di portabilità per identificare le caratteristiche del codice che influenzano significativamente le prestazioni e la compatibilità. Lo studio ha valutato queste applicazioni attraverso molteplici dimensioni, tra cui la latenza di lancio del kernel, la gestione della concorrenza, la compatibilità con le librerie, la gestione della memoria e l'overhead di compilazione.

Risultati Chiave e Conclusioni

1. Latenza di Lancio del Kernel

Il lancio di un kernel non è istantaneo e può variare da microsecondi a decine di microsecondi.

• Kokkos: È stato osservato che aggiunge un overhead significativo (decine di microsecondi), in particolare sulle GPU AMD, che possiedono già latenze di lancio native più elevate rispetto a NVIDIA o Intel.
• SYCL, HIP, OpenMP, Alpaka: Non hanno mostrato penalità aggiuntive rilevanti rispetto alle implementazioni native.
• Specifiche AMD: Chiamare rocRAND su GPU AMD ha causato latenze estreme alla prima chiamata (superiori a 1 secondo), suggerendo che i kernel a breve esecuzione possono soffrire di un grave degrado delle prestazioni se utilizzano strati di portabilità che esacerbano l'overhead di lancio.

2. Concorrenza e Thread Pool

Le applicazioni HEP spesso girano all'interno di framework più ampi che utilizzano il multithreading (ad esempio, Intel TBB).

• Kokkos: Ha dimostrato incompatibilità significative con i thread pool esterni. Il suo backend seriale utilizza lock che serializzano le chiamate, e il suo backend per i thread proibisce le chiamate da thread esterni. Il lancio concorrente di kernel è attualmente praticabile solo con i backend CUDA/HIP utilizzando API specifiche per l'architettura, limitando la portabilità.
• SYCL: Le implementazioni sono state incoerenti. Alcune hanno serializzato le chiamate concorrenti da thread diversi a seconda del produttore del compilatore e della versione. Le architetture AMD sembravano mancare di una soluzione SYCL concorrente al momento dello studio.

3. Compatibilità con Librerie Esterne e Compilatori

• Eigen: Molte versioni sono incompatibili con nvcc di NVIDIA, influenzando gli strati di portabilità che dipendono da esso (Kokkos, Alpaka). È inoltre attualmente incompatibile con l'offload OpenMP.
• ROOT: Non può essere compilato interamente con nvc++ di NVIDIA. La compilazione ibrida (usando g++ e nvc++) è possibile ma complica le regole di build e richiede una gestione attenta dell'allocazione della memoria (i dati GPU devono essere allocati da nvc++).
• Complessità di Build: Supportare più strati di portabilità in un unico progetto richiede regole di build complesse (uso di #ifdef, separazione dei file) per evitare conflitti di header tra strati come Kokkos e CUDA.

4. Strutture Dati e Trasferimenti di Memoria

• Layout Orientati agli Oggetti vs Flat: Le strutture dati complesse e orientate agli oggetti, comuni in HEP, si mappano male sulle GPU, che preferiscono layout piatti o Structure of Arrays (SoA).
• Overhead di Portabilità: Strati come Kokkos e SYCL forniscono astrazioni (ad esempio, Kokkos::Views, sycl::buffers) che aumentano l'overhead di allocazione e trasferimento, specialmente per piccoli oggetti.
• Inizializzazione: L'inizializzazione predefinita della memoria in Kokkos::Views può essere un collo di bottiglia per le prestazioni; questo può essere mitigato usando Kokkos::ViewAllocateWithoutInitializing.
• Unified Shared Memory (USM): I trasferimenti automatici USM (usati da std::par) sono stati osservati essere più lenti dei trasferimenti espliciti e potrebbero non attivarsi al momento ottimale.
• Array Jagged (Dentellati): Nessuna delle API supporta nativamente vettori multidimensionali di dimensioni variabili (array dentellati), richiedendo la creazione manuale, la riduzione delle dimensioni o il padding.

5. RNG, FFT e Atomi

• Librerie Native: Le librerie specifiche dei vendor (cuRAND, rocRAND, oneMKL) offrono prestazioni ottimizzate ma mancano di portabilità. Sebbene Kokkos fornisca implementazioni cross-architettura di RNG/FFT, queste potrebbero non eguagliare le prestazioni delle versioni native.
• Contributi HEP-CCE: Il gruppo ha sviluppato una libreria header-only che avvolge gli RNG nativi del dispositivo con un'API uniforme e backend hardware-native per oneMKL, per consentire a SYCL di chiamare le implementazioni native tra diversi vendor.
• Atomi: Le operazioni atomiche si mappano male sull'hardware GPU. Le prestazioni variano enormemente in base al compilatore e all'hardware (ad esempio, atomic<int> su NVIDIA con OpenMP/clang13 era 35 volte più lento di CUDA). Gli atomi interi generalmente performano peggio rispetto agli atomi float/double.

6. Tempo di Compilazione

• Overhead del Compilatore: nvc++ è generalmente da 2 a 3 volte più lento di g++. Kokkos aggiunge un overhead del 10%-20%.
• Codice Pesante di Template: Madgraph, che utilizza codice auto-generato pesantemente basato su template, ha sperimentato tempi di compilazione con il backend icpx di SYCL superiori di 1000 volte rispetto ad altri backend per kernel di grandi dimensioni.

7. Provisioning del Runtime

• Kokkos: Richiede la build con identificativi hardware specifici. Un binario costruito per una GPU (es. V100) potrebbe non girare su un'altra (es. A100). Supportare librerie condivise e molteplici architetture crea una matrice complessa di configurazioni di build, sfidando la distribuzione dei binari.
• Alpaka: Richiede la compilazione separata per ogni backend del compilatore.
• SYCL/OneAPI: Binari multi-piattaforma stanno emergendo ma sono spesso limitati a specifiche architetture o richiedono l'identificazione al tempo della compilazione.
• OpenMP: Alcune implementazioni (LLVM Clang) permettono di elencare più architetture al tempo della compilazione e selezionare dinamicamente il backend al runtime.

Significato e Conclusione

Il documento conclude che nessun singolo modulo di portabilità GPU è ottimale per tutti gli scenari. La scelta di uno strato di portabilità deve essere guidata da un'analisi attenta delle caratteristiche specifiche dell'applicazione, inclusi il runtime del kernel, i requisiti di concorrenza, la complessità delle strutture dati e le dipendenze dalle librerie esterne.

Gli autori sottolineano che l'ecosistema è in rapida evoluzione. Sia gli strati di portabilità (Kokkos, SYCL, ecc.) che i compilatori nativi (clang, nvcc, hipcc) stanno migliorando, con tendenze verso un maggiore supporto delle funzionalità e prestazioni ottimizzate. Il documento raccomanda di monitorare frequentemente queste tecnologie per garantire il miglior abbinamento tra codebase e strati di portabilità. Inoltre, viene notato che gli emergenti standard C++ (come P2300 per l'esecuzione asincrona) potrebbero semplificare queste scelte in futuro, sebbene il supporto hardware diffuso rimanga a diversi anni di distanza.

In definitiva, selezionare uno strato di portabilità senza comprendere appieno la natura del codice, le dipendenze esterne e l'ambiente di runtime rischia di causare problemi significativi sia nelle prestazioni che nella flessibilità di distribuzione.