Evaluating Application Characteristics for GPU Portability Layer Selection

Dit artikel presenteert een studie die representatieve heterogene applicaties van grote High Energy Physics-experimenten analyseert om de belangrijkste kenmerken te identificeren die de prestaties over verschillende GPU-portabiliteitslagen beïnvloeden, om zo ontwikkelaars te begeleiden bij het selecteren van de meest geschikte technologie voor hun specifieke behoeften.

De kern

Stel je voor dat je een chef-kok bent die probeert een enorm banket te bereiden. Je hebt drie verschillende soorten krachtige ovens in je keuken: één gemaakt door NVIDIA, één door AMD en één door Intel. Elke oven kookt voedsel op een andere manier, gebruikt andere knoppen en vereist verschillende recepten om optimaal te functionen.

Als je je recept specifiek schrijft voor de NVIDIA-oven (met een taal genaamd CUDA), kun je datzelfde recept niet zomaar in de AMD- of Intel-ovens stoppen. Je zou het hele ding opnieuw moeten schrijven. Dit is een probleem, want je weet niet altijd welke oven er morgen in je keuken staat.

Om dit op te lossen, bespreekt het artikel "portability layers" (draagbaarheidslagen). Beschouw deze als universele vertalers of slimme adapters. Ze laten je één meesterrecept schrijven dat de vertaler vervolgens omzet naar de specifieke taal die elke oven begrijpt. Het artikel onderzoekt verschillende van deze vertalers (zoals Kokkos, SYCL, OpenMP en Alpaka) om te zien welke het beste past bij verschillende soorten kooktaken.

Dit is wat de auteurs vonden toen ze deze vertalers testten met echte "recepten" uit experimenten in de hoge-energiefysica (zoals die gebruikt worden om subatomaire deeltjes te bestudelen):

1. Het "Opstarttijd"-probleem

Het aanzetten van een GPU (de oven) gaat niet direct. Het duurt een paar milliseconden voordat hij wakker is en klaar is om te beginnen.

• Het probleem: Sommige vertalers zijn traag bij het starten van het kookproces. Kokkos kan bijvoorbeeld een aanzienlijke vertraging veroorzaken bij het gebruik van AMD-ovens. Als je kooktaak erg kort is (zoals een ei koken voor 10 seconden), en de vertaler heeft 5 seconden nodig om alleen al het fornuis aan de praat te krijgen, heb je de helft van je tijd verspild.
• De les: Als je taken klein en snel zijn, vermijd dan vertalers die de opstarttijd traag maken.

2. Het "Drukke Keuken"-probleem

In een echte natuurkundelab werkt de GPU niet alleen. Het is onderdeel van een groter systeem waar veel mensen (threads) tegelijkertijd de oven willen gebruiken.

• Het probleem: Sommige vertalers zijn slecht in het omgaan met drukte. Kokkos heeft bijvoorbeeld een regel die zegt: "Slechts één persoon kan tegelijk met de oven praten," wat voor een verkeersopstopping zorgt als meerdere chefs gelijktijdig taken willen starten. SYCL is een beetje inconsistent; soms laat het iedereen tegelijk koken, en soms dwingt het hen om in de rij te staan, afhankelijk van welke versie van de vertaler je gebruikt.
• De les: Als je applicatie nodig heeft dat veel mensen tegelijk werken, heb je een vertaler nodig die weet hoe hij een drukke keuken moet beheren zonder de deuren op slot te doen.

3. Het "Gereedschapskist Compatibiliteit"-probleem

Fysica-recepten gebruiken vaak speciale tools (bibliotheken zoals ROOT of Eigen) die helpen met wiskunde en data.

• Het probleem: Sommige van deze tools werken niet goed samen met de vertalers. Zo gaat een populaire wiskundige tool genaamd Eigen vaak kapot wanneer deze wordt gebruikt met de NVIDIA-compiler, waar veel vertalers op vertrouwen. Ook is het proberen te gebruiken van twee verschillende compilers (één voor de CPU en één voor de GPU) in hetzelfde project als het proberen te bouwen van een huis met twee verschillende bouwtekeningen die niet overeenkomen — het maakt de constructie (het bouwen van de software) een nachtmerrie.
• De les: Controleer voordat je een vertaler kiest of je favoriete tools erin passen.

4. Het "Meubelopstelling"-probleem

GPU's houden van eenvoudige, platte lay-outs. Ze geven de voorkeur aan data die gerangschikt is als een nette rij dozen. Echter, fysica-data komt vaak in complexe, rommelige vormen (zoals een stapel verschillende maten koffers).

• Het probleem: Vertalers proberen dit probleem op te lossen door de data in speciale containers te verpakken. Hoewel dit de code draagbaar maakt, voegt het "overhead" toe — zoals elk enkel item in een koffer stoppen voordat je het verplaatst, zelfs als je alleen maar één sok nodig hebt om te verplaatsen. Dit vertraagt de boel. Ook zijn geen van de vertalers erg goed in het afhandelen van "gekartelde" data (rijen van verschillende lengtes), wat zeer gebruikelijk is in de natuurkunde.
• De les: Als je data complex en rommelig is, kan de vertaler je vertragen door te proberen alles netjes op te ruimen.

5. Het "Gespecialiseerde Tools"-probleem

Soms heb je een specifieke tool nodig, zoals een willekeurige getallengenerator (RNG) of een snelle Fourier-transformatie (FFT).

• Het probleem: Elke ovenfabrikant heeft zijn eigen super-snelle, gespecialiseerde versies van deze tools. De universele vertalers bevatten deze gespecialiseerde versies vaak niet, of ze gebruiken hun eigen tragere versies. Hoewel je de vertaler kunt dwingen om de native tool van de oven te gebruiken, verbreekt dit de "portabiliteit" omdat die tool alleen werkt op die specifieke oven.
• De les: Als je zwaar leunt op deze specifieke tools, moet je misschien kiezen tussen snelheid (gebruik van de native tool van de oven) of portabiliteit (gebruik van de generieke tool van de vertaler).

6. De "Constructietijd" en "Verhuisdag"-problemen

• Het recept bouwen: Sommige vertalers maken de "kooktijd" (compilatieduur) veel langer. Voor enorme projecten kan het gebruik van bepaalde vertalers het bouwproces uren laten duren in plaats van minuten.
• De keuken verhuizen: Als je je software bouwt voor een specifieke oven (bijv. een NVIDIA V100), werkt het misschien niet op een nieuwere (een NVIDIA A100). Sommige vertalers vereisen dat je een aparte versie bouwt voor elk type oven dat je mogelijk tegenkomt. Dit creëert een enorme logistieke hoofdpijn voor het distribueren van de software naar verschillende laboratoria.

Het Eindvonnis

Het artikel concludeert dat er geen "perfecte" vertaler bestaat.

• Kokkos is goed voor sommige zaken, maar worstelt met gelijktijdigheid en opstarttijden op bepaalde hardware.
• SYCL is krachtig, maar kan inconsistent zijn afhankelijk van de compilerversie.
• OpenMP en anderen hebben hun eigen sterktes en zwaktes wat betreft het beheer van geheugen en verschillende hardware.

De kernboodschap: Je kunt niet zomaar een vertaler kiezen omdat hij populair is. Je moet kijken naar je specifieke "recept" (je applicatie). Als je code kort en snel is, kies dan een vertaler met een lage opstarttijd. Als je code complex is en veel tools gebruikt, kies dan een vertaler die goed samenwerkt met die tools.

De auteurs merken ook op dat deze technologieën snel evolueren, net zoals er elk jaar nieuwe modellen ovens uitkomen. Wat vandaag het beste werkt, kan morgen veranderd zijn, dus ontwikkelaars moeten de ontwikkelingen nauwlettend blijven volgen. In de toekomst kunnen nieuwe standaarden deze keuzes makkelijker maken, maar voor nu is zorgvuldige testen de enige manier om de juiste match te vinden.

Technische samenvatting

Technische Samenvatting: Evaluatie van Applicatiekenmerken voor de Selectie van GPU-portabiliteitslagen

Probleemstelling

High Energy Physics (HEP) computing is steeds meer afhankelijk van GPU's voor prestaties, maar het hardwarelandschap diversifieert buiten NVIDIA om naar inclusief AMD en Intel. Hoewel vendor-specifieke talen zoals CUDA hoge prestaties bieden, beperken ze de inzetbaarheid tot specifieke hardware, wat de portabiliteit vermindert. Daartegenover staan portabiliteitslagen zoals Kokkos, SYCL, HIP, OpenMP, Alpaka en std::par, die uitvoering over diverse CPU- en GPU-architecturen beloven. Deze lagen zijn echter niet functioneel equivalent; ze vertonen onderscheidende kenmerken met betrekking tot prestaties, concurrency-ondersteuning en compatibiliteit met bestaande codebases. Het kernprobleem dat wordt geadresseerd, is het gebrek aan een duidelijk kader voor ontwikkelaars om de meest geschikte portabiliteitslaag te selecteren op basis van specifieke applicatiekenmerken, kernelstructuren en runtime-omgevingen.

Methodologie

Het High Energy Physics Center for Computational Excellence (HEP-CCE) heeft een vergelijkende studie uitgevoerd met behulp van representatieve heterogene applicaties uit grote HEP-experimenten:

• ATLAS: FastCaloSim (geparameteriseerde calorimeter simulatie).
• CMS: Patatrack (heterogene pixel reconstructie met 40+ kernels, multithreading en geheugenpools) en P2R (single-kernel track reconstructie).
• DUNE: Wire-Cell Toolkit (rasterisatie, scatter-add en FFT convolutie kernels).
• Event Generators: GPU-ports van Sherpa (Pepper) en Madgraph (gebruikmakend van auto-gegenereerde code).

Deze testbedden zijn vanuit seriële C++ geport naar diverse portabiliteitstechnologieën om codekenmerken te identificeren die de prestaties en compatibiliteit significant beïnvloeden. De studie evalueerde deze applicaties over meerdere dimensies, waaronder kernel launch-latentie, concurrency-afhandeling, bibliotheekcompatibiliteit, geheugenbeheer en compilatie-overhead.

Belangrijkste Bevindingen en Resultaten

1. Kernel Launch Latency

Het lanceren van een kernel is niet instant en kan variëren van microseconden tot tientallen microseconden.

• Kokkos: Werd waargenomen als een bron van significante overhead (tientallen microseconden), met name op AMD GPU's, die al hogere native launch-latenties hebben dan NVIDIA of Intel.
• SYCL, HIP, OpenMP, Alpaka: Vertoonden geen grote extra boetes vergeleken met native implementaties.
• AMD Specifieke zaken: Het aanroepen van rocRAND op AMD GPU's resulteerde in extreme eerste-oproep-latenties (meer dan 1 seconde), wat suggereert dat kortlopende kernels ernstige prestatievermindering kunnen ondervinden als ze gebruikmaken van portabiliteitslagen die de launch-overhead verergeren.

2. Concurrency en Thread Pools

HEP-applicaties draaien vaak binnen grotere frameworks die gebruikmaken van multithreading (bijv. Intel TBB).

• Kokkos: Toonde significante incompatibiliteiten met externe thread pools. De seriële backend gebruikt locks die aanroepen serialiseren, en de threads-backend verbiedt externe thread-aanroepen. Gelijktijdige kernel-launches zijn momenteel alleen haalbaar met CUDA/HIP-backends via architectuur-specifieke API's, wat de portabiliteit beperkt.
• SYCL: Implementaties waren inconsistent. Sommige implementaties serialiseerden gelijktijdige aanroepen vanuit verschillende threads, afhankelijk van de fabrikant van de compiler en de versie. AMD-architecturen leken op het moment van de studie geen concurrente SYCL-oplossing te missen.

3. Externe Bibliotheek en Compiler Compatibiliteit

• Eigen: Veel versies zijn incompatibel met NVIDIA's nvcc, wat portabiliteitslagen die op nvcc vertrouwen (Kokkos, Alpaka) beïnvloedt. Het is ook momenteel incompatibel met OpenMP offload.
• ROOT: Kan niet volledig worden gecompileerd met NVIDIA's nvc++. Hybride compilatie (met gebruik van g++ en nvc++) is mogelijk, maar compliceert de build-regels en vereist zorgvuldig geheugenallocatiemanagement (GPU-data moet door nc++ worden gealloceerd).
• Build Complexiteit: Het ondersteunen van meerdere portabiliteitslagen in één project vereist complexe build-regels (bijv. #ifdef gebruik, scheiding van bestanden) om header-conflicten tussen lagen zoals Kokkos en CUDA te voorkomen.

4. Datastructuren en Geheugenoverdrachten

• Object-Georiënteerd vs. Flat Layouts: Complexe, object-georiënteerde datastructuren die gebruikelijk zijn in HEP mappen slecht naar GPU's, die de voorkeur geven aan platte layouts of Structures of Arrays (SoA).
• Portabiliteits-overhead: Lagen zoals Kokkos en SYCL bieden abstracties (bijv. Kokkos::Views, sycl::buffers) die de allocatie- en transfer-overhead verhogen, vooral voor kleine objecten.
• Initialisatie: Standaard geheugeninitialisatie in Kokkos::Views kan een prestatiebottleneck zijn; dit kan worden verzacht met Kokkos::ViewAllocateWithoutInitializing.
• Unified Shared Memory (USM): Automatische USM-transfers (gebruikt door std::par) werden waargenomen als trager dan expliciete transfers en worden mogelijk niet op optimale momenten getriggerd.
• Jagged Arrays: Geen van de API's ondersteunt native multidimensionale vectoren van variërende grootte (jagged arrays), wat handmatige constructie, dimensiereductie of padding vereist.

5. RNGs, FFTs en Atomics

• Native Bibliotheken: Vendor-specifieke bibliotheken (cuRAND, rocRAND, oneMKL) bieden getunede prestaties maar missen portabiliteit. Hoewel Kokkos cross-architectuur RNG/FFT-implementaties biedt, kunnen deze de native prestaties mogelijk niet evenaren.
• HEP-CCE Bijdragen: De groep heeft een header-only bibliotheek ontwikkeld die device-native RNG's omhult met een uniforme API en hardware-native backends voor oneMKL om SYCL in staat te stellen native implementaties over verschillende vendors aan te roepen.
• Atomics: Atomaire operaties mappen slecht naar GPU-hardware. Prestaties variëren sterk per compiler en hardware (bijv. atomic<int> op NVIDIA met OpenMP/clang13 was 35x trager dan CUDA). Integer atomics presteren over het algemeen slechter dan float/double atomics.

6. Compilatietijd

• Compiler Overhead: nvc++ is over het algemeen 2x–3x trager dan g++. Kokkos voegt een overhead van 10%–20% toe.
• Template Heavy Code: Madgraph, dat gebruikmaakt van zwaar getemplateerde auto-gegenereerde code, ervoer compilatietijden met de SYCL icpx backend die meer dan 1000x trager waren dan andere backends voor grote kernels.

7. Runtime Provisioning

• Kokkos: Vereist builden met specifieke hardware-identificatoren. Een binary gebouwd voor één GPU (bijv. V100) kan mogelijk niet draaien op een andere (bijv. A100). Het ondersteunen van shared libraries en meerdere architecturen creëert een complexe matrix van build-configuraties, wat de distributie van binaries bemoeilijkt.
• Alpaka: Vereist aparte compilatie voor elke backend-compiler.
• SYCL/OneAPI: Multi-platform binaries komen opzetten, maar zijn vaak beperkt tot specifieke architecturen of vereisen identificatie tijdens compile-time.
• OpenMP: Sommige implementaties (LLVM Clang) staan toe om meerdere architecturen tijdens compile-time op te sommen en de backend tijdens runtime dynamisch te selecteren.

Betekenis en Conclusie

Het paper concludeert dat geen enkele GPU-portabiliteitsoplossing optimaal is voor alle scenario's. De keuze voor een portabiliteitslaag moet worden gedreven door een zorgvuldige analyse van de specifieke kenmerken van de applicatie, inclusioneel de kernel runtime, concurrency-vereisten, complexiteit van datastructuren en externe bibliotheekafhankelijkheden.

De auteurs benadrukken dat het ecosysteem snel evolueert. Zowel portabiliteitslagen (Kokkos, SYCL, etc.) als native compilers (clang, nvcc, hipcc) verbeteren, waarbij de trends bewegen naar verhoogde feature-ondersteuning en geoptimaliseerde prestaties. Het paper pleit voor frequente monitoring van deze technologieën om de beste match tussen codebases en portabiliteitslagen te waarborgen. Daarnaast merkt het paper op dat opkomende C++-standaarden (zoals P2300 voor asynchrone executie) deze keuzes in de toekomst mogelijk vereenvoudigen, hoewel brede hardware-ondersteuning nog jaren weg is.

Uiteindelijk brengt het selecteren van een portabiliteitslaag zonder de aard van de code, externe afhankelijkheden en de runtime-omgeving volledig te begrijpen, het risico met zich mee van aanzienlijke problemen in zowel prestaties als de flexibiliteit van de inzetbaarheid.