Accelerating the Particle-In-Cell code ECsim with OpenACC

Dit artikel presenteert de optimalisatie van de semi-impliciete PIC-code ECsim voor exascale-architecturen met OpenACC, wat resulteert in een snelheidswinst van 5x en een energiebesparing van 3x op het Leonardo-systeem, met een sterke schaalbaarheid tot 1024 GPU's.

De kern

Stel je voor dat je een enorme, chaotische danszaal hebt vol met miljarden dansers. Elke danser vertegenwoordigt een deeltje in een plasma (zoals in een ster of een kernfusiereactor). Je wilt precies weten hoe ze bewegen, hoe ze elkaar duwen en trekken, en hoe ze samen een complex patroon vormen.

Dit is wat wetenschappers doen met een computerprogramma genaamd ECsim. Het is een "danssimulator" voor plasma. Maar hier is het probleem: om dit nauwkeurig te doen, moet de computer elke danser individueel volgen. Dat is zo veel rekenwerk dat het op een gewone computer (zoals een supercomputer met alleen maar CPU's) eeuwen kan duren.

In dit paper vertellen de auteurs hoe ze deze simulator hebben versneld door een nieuwe "dansvloer" te gebruiken: GPU's (de krachtige grafische kaarten die je ook in gaming-computers hebt).

Hier is de uitleg in simpele taal, met een paar creatieve vergelijkingen:

1. Het Probleem: De trage chef-kok

Stel je voor dat de CPU (de normale processor) een enkele, zeer slimme chef-kok is die probeert een gigantisch diner voor 100.000 gasten te bereiden. Hij moet elke pan één voor één controleren, elke groente één voor één snijden. Het resultaat is goed, maar het duurt heel lang.

De wetenschappers wilden dit versnellen. Ze hadden een nieuwe keuken nodig met duizenden kleine, snelle koks die allemaal tegelijk kunnen werken. Dat zijn de GPU's.

2. De Oplossing: OpenACC als de "Magische Chef"

Om de GPU's te gebruiken, moet je de code van het programma vaak helemaal herschrijven. Dat is als het bouwen van een compleet nieuw restaurantgebouw. Dat kost veel tijd en geld.

In plaats daarvan gebruikten deze onderzoekers een trucje genaamd OpenACC.

• De Analogie: Stel je voor dat je de oude chef-kok (de CPU-code) een magische hoed opzet. In plaats van het hele restaurant af te breken, zegt de chef tegen zijn team: "Jullie gaan nu in de nieuwe keuken werken!"
• Wat het doet: OpenACC is die magische hoed. Het is een set van simpele instructies (zoals post-itjes op de code) die de computer vertellen: "Hey, deze specifieke taak doe je op de snelle GPU-keuken."
• Het resultaat: Ze hoefden de code nauwelijks aan te passen, maar het programma werd ineens 5 keer sneller.

3. Wat gebeurde er precies? (De drie taken)

Het programma heeft drie hoofdtaken, en ze hebben ze allemaal geoptimaliseerd:

• De Deeltjesbeweging (De dansers): De computer moet elke danser een stap laten zetten. Dit deden ze op de GPU.
  • Resultaat: De dansers bewegen nu razendsnel.
• De Krachten meten (De zwaartekracht): De computer moet berekenen hoe de dansers elkaar beïnvloeden. Dit is het zwaarste werk (76% van de tijd!).
  • Resultaat: Dit werk werd 15 keer sneller op de GPU.
• De Energiebesparing: Omdat het programma nu zo snel klaar is, stopt de computer veel eerder.
  • De Analogie: Het is alsof je een auto hebt die 5 keer sneller rijdt, maar ook 3 keer minder benzine verbruikt per reis. De onderzoekers bespaarden dus enorm veel energie.

4. De Nieuwe Super-keuken: GH200 en Unify Memory

De onderzoekers testten hun programma op verschillende generaties van GPU's. Ze ontdekten iets interessants met de nieuwste, duurste GPU (de GH200).

• De Vergelijking:
  • Oude GPU's: De chef (CPU) en de koks (GPU) hadden elk hun eigen voorraadkast. Om te koken, moesten ze eten heen en weer slepen over een smalle brug (de PCIe-verbinding). Dat kostte tijd.
  • Nieuwe GH200: De chef en de koks delen nu één enorme, gezamenlijke voorraadkast. Ze hoeven niets meer te slepen; ze pakken gewoon wat ze nodig hebben direct uit dezelfde kast.
  • Het effect: Voor bepaalde taken was de nieuwe GPU 12 keer sneller dan de oude, puur omdat ze niet meer hoefden te wachten op het transport van data.

5. Schalen: Van één zaal naar een stadion

Tot slot keken ze of het systeem ook werkt als je het enorm groot maakt.

• Sterk schalen (Meer koks, zelfde taak): Als je 1024 GPU's gebruikt, werkt het nog steeds heel goed. Het is alsof je een stadion vult met duizenden koks die allemaal tegelijk een taak doen zonder dat ze in de weg lopen. De efficiëntie blijft boven de 70%.
• Zwak schalen (Grotere taak, meer koks): Als je een nog groter diner wilt bereiden, kun je gewoon meer koks toevoegen en blijft de tijd ongeveer hetzelfde.

Conclusie

Kortom: De onderzoekers hebben een oude, trage plasma-simulator (ECsim) opgefrist met een paar simpele instructies (OpenACC). Hierdoor draait het nu op de krachtigste supercomputers van Europa (Leonardo) met 5 keer de snelheid en 3 keer minder energie.

Ze hebben bewezen dat je niet altijd een compleet nieuw programma hoeft te bouwen om gebruik te maken van de nieuwste technologie; soms helpt het al om de oude code een beetje "slimmer" te maken met de juiste aanwijzingen. Dit opent de deur voor veel snellere en goedkopere simulaties van sterren, kernfusie en andere complexe natuurverschijnselen.

Technische samenvatting

Titel: Versnelling van de Particle-In-Cell-code ECsim met OpenACC

1. Het Probleem

De Particle-In-Cell (PIC) methode is een fundamentele computertechniek in de plasmafysica om kinetische effecten en complexe plasma-interacties te modelleren. Traditionele PIC-algoritmes gebruiken vaak expliciete tijdsdiscretisatie, wat strenge stabiliteitsvoorwaarden oplegt (zoals de Courant-Friedrichs-Lewy-voorwaarde). Dit vereist dat de kleinste schaal in het probleem wordt opgelost, wat leidt tot extreem hoge rekenkosten bij het modelleren van multischaal plasma-scenario's.

Semi-impliciete methoden bieden een compromis: ze behouden de voordelen van impliciete discretisatie (relaxatie van stabiliteitsvoorwaarden) zonder de hoge complexiteit en convergentieproblemen van volledig impliciete methoden. De code ECsim implementeert zo'n semi-impliciete, energiebehoudende methode. Hoewel ECsim al parallel is op CPU's met MPI, is het voorbereiden voor exascale-architecturen (supercomputers met GPU-versnellers) een uitdaging. Het herschrijven van de code naar een native GPU-taal (zoals CUDA) zou ingrijpende wijzigingen vereisen en de leesbaarheid en onderhoudbaarheid van de code in gevaar brengen.

2. Methodologie

De auteurs hebben gekozen voor een pragma-gebaseerde versnellingsstrategie met OpenACC. Deze aanpak maakt het mogelijk om bestaande C/C++-code te versnellen met minimale wijzigingen in de broncode, terwijl de oorspronkelijke structuur behouden blijft.

• Doelcode: ECsim, een MPI-gebaseerde code die momentopname (moment gathering), veldoplossing (field solver) en deeltjesbeweging (particle mover) uitvoert.
• Hardware: Tests zijn uitgevoerd op het Leonardo-supercomputer (CINECA, Italië), specifiek op de "Booster"-partitie met Intel Xeon Platinum CPU's en NVIDIA A100 GPU's. Vergelijkingen zijn gemaakt met andere GPU-generaties (V100, H100, GH200).
• Implementatie-Details:
  • Versnelling: De zwaarst belastende kernels (deeltjesbewerking en momentopname) zijn versneld met #pragma acc parallel loop.
  • Geheugenbeheer: Om data-overdrachten tussen host en device te optimaliseren, is gebruik gemaakt van managed memory in combinatie met cudaMemPrefetchAsync. Dit vermindert de overhead van expliciete data-migratie.
  • Optimalisaties:
    • 3D-pointers zijn omgezet naar 1D-arrays voor betere geheugenlokaliteit.
    • Geneste mini-loops voor interpolatie en rotatiematrices zijn handmatig "unrolled" om parallelisme te maximaliseren en loop-overhead te verkleinen.
    • Voor de berekening van de massamatrix is de toegangsvolgorde aangepast om cache-efficiëntie te verbeteren.
    • #pragma acc atomic update is gebruikt om race conditions te voorkomen bij het accumuleren van bijdragen van deeltjes naar het rooster.
  • Validatie: De resultaten van de GPU-versie zijn strikt vergeleken met de CPU-referentieversie op basis van energiebehoud en numerieke consistentie bij bekende instabiliteiten (tweestromingsinstabiliteit en stroom-filamentatie-instabiliteit).

3. Belangrijkste Bijdragen

1. OpenACC-Implementatie: De eerste succesvolle migratie van de semi-impliciete PIC-code ECsim naar GPU-versnelling met OpenACC, waarbij slechts minimale code-herstructurering nodig was.
2. Validatie en Benchmarking: Uitgebreide validatie van de numerieke nauwkeurigheid en prestatiebenchmarks in termen van "tijd tot oplossing" (time-to-solution) en "energie tot oplossing" (energy-to-solution).
3. Architectuurvergelijking: Een diepgaande analyse van de prestaties op verschillende NVIDIA GPU-generaties, met een specifieke focus op de voordelen van de GH200 Unified Memory-architectuur.
4. Schalingsanalyse: Uitgebreide sterke en zwakke schalingstests op het Leonardo-systeem tot 1024 GPU's.

4. Resultaten

De prestaties zijn gemeten op het Leonardo-systeem (1 node met 4x NVIDIA A100 GPU's vs. 1 node met 32 CPU-kernen):

• Snelheidswinst: De versnelde code bereikte een 5x snelheidswinst ten opzichte van de CPU-only referentie.
  • De "moment gathering" (de meest tijdrovende stap, oorspronkelijk ~76% van de tijd) werd 15x sneller.
  • De "particle mover" en I/O-routines werden respectievelijk ongeveer 2x en 2x sneller.
• Energie-efficiëntie: Er werd een 3x reductie in energieverbruik bereikt (van ~1295 kJ naar ~415 kJ per simulatie). Dit komt door de kortere rekentijd en het feit dat GPU's een hogere performance-per-watt bieden voor deze specifieke kernels.
• GPU-Generatie Vergelijking:
  • De prestaties verbeterden aanzienlijk van V100 naar A100 en H100, voornamelijk door hogere geheugenbandbreedte en efficiëntere atomische instructies.
  • De GH200 (met gedeeld CPU-GPU geheugen via NVLink-C2C) bood de grootste winst, vooral voor de "particle mover" (2x sneller dan V100). Hierdoor worden data-overdrachtkosten via PCIe geëlimineerd en worden prefetch-bewerkingen efficiënter.
• Schalbaarheid:
  • Sterke schaling: Bij 64 GPU's werd een parallelle efficiëntie van 70% bereikt. Bij grotere probleemgroottes (6720x6720 cellen) bleef de efficiëntie op 83% bij 1024 GPU's.
  • Zwakke schaling: De code behield een parallelle efficiëntie van 78% tot 1024 GPU's, wat aantoont dat communicatie- en geheugenoverhead beperkt blijven bij het vergroten van het probleem.

5. Betekenis en Conclusie

Dit werk demonstreert dat het gebruik van directive-based programming (OpenACC) een zeer effectieve strategie is om bestaande, complexe wetenschappelijke codes zoals ECsim klaar te maken voor exascale-computing zonder de code volledig te herschrijven.

• Efficiëntie: De combinatie van hoge snelheidswinst en aanzienlijke energiebesparing maakt ECsim zeer geschikt voor toekomstige grote schaal plasma-simulaties, zoals die nodig zijn voor fusie-energieonderzoek.
• Toekomstperspectief: Hoewel native CUDA-implementaties potentieel nog hogere prestaties kunnen bieden, biedt de OpenACC-aanpak nu al een robuuste en onderhoudbare oplossing. De auteurs werken momenteel aan een native CUDA-versie en onderzoeken het gebruik van cellulaire automata voor verdere optimalisatie.
• Hardware Impact: De resultaten benadrukken het belang van nieuwe hardware-architecturen zoals de GH200 Unified Memory voor data-intensieve wetenschappelijke toepassingen.

Kortom, de paper bewijst dat ECsim nu efficiënt kan draaien op moderne GPU-gebaseerde supercomputers, wat de weg vrijmaakt voor realistische, multischaal plasma-simulaties op exascale-niveau.