Evaluating Application Characteristics for GPU Portability Layer Selection

Diese Arbeit präsentiert eine Studie zur Analyse repräsentativer heterogener Anwendungen aus großen Hochenergiephysik-Experimenten, um die wesentlichen Merkmale zu identifizieren, die die Leistung über verschiedene GPU-Portabilitätslayer hinweg beeinflussen, und damit Entwickler bei der Auswahl der am besten geeigneten Technologie für ihre spezifischen Anforderungen zu leiten.

Das Wesentliche

Stellen Sie sich vor, Sie sind ein Chefkoch, der versucht, ein riesiges Bankett zu kochen. Sie haben drei verschiedene Arten von Hochleistungsöfen in Ihrer Küche: einen von NVIDIA, einen von AMD und einen von Intel. Jeder Ofen kocht das Essen anders, verwendet andere Knöpfe und erfordert unterschiedliche Rezepte, um seine beste Leistung zu erbringen.

Wenn Sie ein Rezept speziell für den NVIDIA-Ofen schreiben (unter Verwendung einer Sprache namens CUDA), können Sie dasselbe Rezept nicht einfach in die AMD- oder Intel-Öfen geben. Sie müssten das ganze Rezept umschreiben. Das ist ein Problem, denn man weiß ja nicht immer, welchen Ofen man morgen in seiner Küche stehen hat.

Um dieses Problem zu lösen, diskutiert das Paper über „Portabilitäts-Schichten“ (Portability Layers). Denken Sie an diese als universelle Übersetzer oder intelligente Adapter. Sie lassen Sie ein einziges Master-Rezept schreiben, das der Übersetzer dann in die spezifische Sprache umwandelt, die jeder Ofen versteht. Das Paper untersucht mehrere dieser Übersetzer (wie Kokkos, SYCL, OpenMP und Alpaka), um zu sehen, welcher für welche Art von Kochaufgaben am besten geeignet ist.

Hier ist das, was die Autoren herausgefunden haben, als sie diese Übersetzer mit echten „Rezepten“ aus Experimenten der Hochenergiephysik (wie sie zur Untersuchung von subatomaren Teilchen verwendet werden) getestet haben:

1. Das „Startzeit“-Problem

Das Einschalten eines GPU (eines Ofens) geschieht nicht instantan. Es dauert einige Millisekunden, bis er aufgewacht ist und bereit ist.

• Das Problem: Einige Übersetsetzer sind langsam beim Starten des Kochvorgangs. Kokkos kann beispielsweise eine erhebliche Verzögerung verursachen, wenn AMD-Öfen verwendet werden. Wenn Ihre Kochaufgabe sehr kurz ist (wie das Kochen eines Eies für 10 Sekunden), und der Übersetzer allein 5 Sekunden braucht, um den Herd zu starten, haben Sie die Hälfte Ihrer Zeit verschwendet.
• Die Lektion: Wenn Ihre Aufgaben winzig und schnell sind, vermeiden Sie Übersetzer, die den Start langsam machen.

2. Das „Überfüllte Küche“-Problem

In einem echten Physiklabor arbeitet die GPU nicht allein. Sie ist Teil eines größeren Systems, in dem viele Menschen (Threads) gleichzeitig den Ofen benutzen wollen.

• Das Problem: Einige Übersetzer sind schlecht darin, mit Menschenmengen umzugehen. Kokkos hat zum Beispiel eine Regel, die besagt: „Nur eine Person darf gleichzeitig mit dem Ofen sprechen“, was zu einem Verkehrsstau führt, wenn mehrere Köche gleichzeitig Aufgaben starten wollen. SYCL ist etwas inkonsistent; manchmal lässt es alle gleichzeitig kochen, und manchmal zwingt es sie, in einer Schlange zu warten, je nachdem, welche Version des Übersetzers man verwendet.
• Die Lektion: Wenn Ihre Anwendung benötigt, dass viele Menschen gleichzeitig arbeiten, brauchen Sie einen Übersetzer, der weiß, wie man eine belebte Küche verwaltet, ohne die Türen zu verriegeln.

3. Das „Werkzeugkasten-Kompatibilitätsproblem“

Physik-Rezepte verwenden oft spezielle Werkzeuge (Bibliotheken wie ROOT oder Eigen), die bei der Mathematik und den Daten helfen.

• Das Problem: Manche dieser Werkzeuge spielen nicht gut mit den Übersetzern zusammen. Zum Beispiel bricht ein beliebtes Mathematik-Tool namens Eigen oft ab, wenn es mit dem NVIDIA-Compiler verwendet wird, auf den viele Übersetzer angewiesen sind. Auch der Versuch, zwei verschiedene Compiler (einen für die CPU und einen für die GPU) im selben Projekt zu verwenden, ist wie der Versuch, ein Haus mit zwei verschiedenen Bausätzen zu bauen, die nicht zusammenpassen – das macht den Bau (das Erstellen der Software) zu einem Albtraum.
• Die Lektion: Bevor Sie einen Übersetzer auswählen, prüfen Sie, ob Ihre Lieblingswerkzeuge darin hineinpassen.

4. Das „Möbelanordnungsproblem“

GPUs lieben einfache, flache Layouts. Sie bevorzugen Daten, die wie eine ordentliche Reihe von Boxen angeordnet sind. Die Physik-Daten kommen jedoch oft in komplexen, unordentlichen Formen (wie einem Haufen verschieden großer Koffer).

• Das Problem: Übersetzer versuchen, dieses Chaos zu beheben, indem sie die Daten in spezielle Container einwickeln. Während dies die Daten portabel macht, fügt es „Overhead“ hinzu – so als würde man jeden einzelnen Gegenstand erst in einen Koffer packen, bevor man ihn bewegt, selbst wenn man nur eine einzige Socke bewegen muss. Das verlangsamt den Vorgang. Außerdem sind keine der Übersetzer besonders gut darin, mit „zackigen“ (jagged) Daten umzugehen (Reihen unterschiedlicher Länge), was in der Physik sehr verbreitet ist.
• Die Lektion: Wenn Ihre Daten komplex und unordentlich sind, könnte der Übersetzer Sie ausbremsen, während er versucht, Ordnung zu schaffen.

5. Das „Spezialisierte Werkzeuge“-Problem

Manchmal benötigen Sie ein spezielles Werkzeug, wie einen Zufallszahlengenerator (RNG) oder eine Schnelle Fourier-Transformation (FFT).

• Das Problem: Jeder Hardware-Hersteller hat seine eigenen, superschnellen, spezialisierten Versionen dieser Werkzeuge. Die universellen Übersetzer enthalten diese spezialisierten Versionen oft nicht oder nutzen ihre eigenen, langsameren Versionen. Man kann zwar den Übersetzer dazu zwingen, das native Werkzeug des Ofens zu verwenden, aber das bricht die „Portabilität“, da dieses Werkzeug nur auf diesem spezifischen Ofen funktioniert.
• Die Lektion: Wenn Sie stark auf diese spezifischen Werkzeuge angewiesen sind, müssen Sie sich vielleicht zwischen Geschwindigkeit (Verwendung des nativen Werkzeugs des Ofens) oder Portabilität (Verwendung des generischen Werkzeugs des Übersetzers) entscheiden.

6. Die „Bauzeit“- und „Umzugsprobleme“

• Das Rezept bauen: Einige Übersetzer machen die „Kochzeit“ (Kompilierzeit) viel länger. Bei riesigen Projekten kann die Nutzung bestimmter Übersetzer den Build-Prozess von Minuten auf Stunden verlängern.
• Die Küche umziehen: Wenn Sie Ihre Software für einen spezifischen Ofen bauen (z. B. einen NVIDIA V100), funktioniert sie vielleicht nicht auf einem neueren (z. B. einem NVIDIA A100). Einige Übersetzer erfordern, dass Sie für jeden einzelnen Typ von Ofen, den Sie eventuell antreffen, eine eigene Version bauen. Dies erzeugt einen massiven logistischen Aufwand für die Verteilung der Software an verschiedene Labore.

Das Endurteil

Das Paper kommt zu dem Schluss, dass es keinen „perfekten“ Übersetzer gibt.

• Kokkos ist gut für manche Dinge, kämpft aber mit Nebenläufigkeit (Concurrency) und Startzeiten auf bestimmter Hardware.
• SYCL ist leistungsstark, kann aber je nach Compiler-Version inkonsistent sein.
• OpenMP und andere haben ihre eigenen Stärken und Schwächen in Bezug auf die Handhabung von Speicher und unterschiedlicher Hardware.

Das Fazit: Man kann einen Übersetzer nicht einfach deshalb wählen, weil er populär ist. Man muss sich sein spezifisches „Rezept“ (seine Anwendung) ansehen. Wenn Ihr Code kurz und schnell ist, wählen Sie einen Übersetzer mit geringer Startzeit. Wenn Ihr Code komplex ist und viele Werkzeuge nutzt, wählen Sie einen, der gut mit diesen Werkzeugen harmoniert.

Die Autoren merken zudem an, dass sich diese Technologien rasant entwickeln, so als kämen jedes Jahr neue Modelle von Öfen auf den Markt. Was heute am besten funktioniert, kann sich morgen schon ändern, daher müssen Entwickler die Landschaft weiterhin genau beobachten. In Zukunft könnten neue Standards diese Entscheidungen erleichtern, aber für den Moment ist sorgfältiges Testen der einzige Weg, um die richtige Passform zu finden.

Technische Zusammenfassung

Technischer Bericht: Evaluierung von Anwendungscharakteristika zur Auswahl einer GPU-Portabilitätschicht

Problemstellung

Das High Energy Physics (HEP) Computing stützt sich zunehmend auf GPUs zur Leistungssteigerung, doch die Hardwarelandschaft diversifiziert sich über NVIDIA hinaus auch in Richtung AMD und Intel. Während herstellerspezifische Sprachen wie CUDA eine hohe Leistung bieten, schränken sie den Einsatz auf spezifische Hardware ein, was die Portabilität reduziert. Im Gegensatz dazu versprechen Portabilitätschichten wie Kokkos, SYCL, HIP, OpenMP, Alpaka und std::par eine Ausführung über diverse CPU- und GPU-Architekturen hinweg. Diese Schichten sind jedoch nicht funktional äquivalent; sie weisen unterschiedliche Merkmale hinsichtlich Leistung, Nebenläufigkeitsunterstützung und Kompatibilität mit bestehenden Codebasen auf. Das Kernproblem besteht im Mangel eines klaren Rahmens für Entwickler, um die am besten geeignete Portabilitäts-Schicht basierend auf spezifischen Anwendungscharakteristika, Kernel-Strukturen und Laufzeitumgebungen auszuwählen.

Methodik

Das High Energy Physics Center for Computational Excellence (HEP-CCE) führte eine Vergleichsstudie unter Verwendung repräsentativer heterogener Anwendungen aus großen HEP-Experimenten durch:

• ATLAS: FastCaloSim (parametrisierte Kalorimetersimulation).
• CMS: Patatrack (heterogene Pixel-Rekonstruktion mit über 40 Kernels, Multithreading und Speicherpools) und P2R (Single-Kernel-Track-Rekonstruktion).
• DUNE: Wire-Cell Toolkit (Rasterisierung, Scatter-Add und FFT-Faltungskernel).
• Ereignisgeneratoren: GPU-Ports von Sherpa (Pepper) und Madgraph (unter Verwendung von automatisch generiertem Code).

Diese Testbeds wurden von serieller C++ zu verschiedenen Portabilitätstechnologien portiert, um Codecharakteristika zu identifizieren, die die Leistung und Kompatibilität signifikant beeinflussen. Die Studie evaluierte diese Anwendungen über mehrere Dimensionen hinweg, einschließlich Kernel-Launch-Latenz, Nebenläufigkeitshandhabung, Bibliothekskompatibilität, Speicherverwaltung und Kompilierungsaufwand.

Zentrale Erkenntnisse und Ergebnisse

1. Kernel-Launch-Latenz

Ein Kernel-Launch erfolgt nicht instantan und kann im Bereich von Mikrosekunden bis hin zu Zehn-Mikrosekunden liegen.

• Kokkos: Es wurde ein signifikanter Overhead (zehn Mikrosekunden) beobachtet, insbesondere auf AMD-GPUs, die bereits höhere native Launch-Lätzenzen als NVIDIA oder Intel aufweisen.
• SYCL, HIP, OpenMP, Alpaka: Diese zeigten keine großen zusätzlichen Strafen im Vergleich zu nativen Implementierungen.
• AMD-Spezifika: Der Aufruf von rocRAND auf AMD-GPUs führte zu extremen Latenzen beim ersten Aufruf (über 1 Sekunde), was darauf hindeutet, dass kurz laufende Kernel unter erheblicher Leistungsdegradation leiden können, wenn sie Portabilitäts-Schichten verwenden, die den Launch-Overhead verschärfen.

2. Nebenläufigkeit und Thread-Pools

HEP-Anwendungen laufen oft innerhalb größerer Frameworks, die Multithreading nutzen (z. B. Intel TBB).

• Kokkos: Zeigte signifikante Inkompatibilitäten mit externen Thread-Pools. Sein serieller Backend verwendet Locks, die Aufrufe serialisieren, und sein Threads-Backend verbietet externe Thread-Aufrufe. Gleichzeitige Kernel-Launches sind derzeit nur mit CUDA/HIP-Backends unter Verwendung architekturspezifischer APIs realisierbar, was die Portabilität einschränkt.
• SYCL: Die Implementierungen waren inkonsistent. Einige serialisierten gleichzeitige Aufrufe von verschiedenen Threads, abhängig vom Hersteller des Compilers und der Version. AMD-Architekturen schienen zum Zeitpunkt der Studie keine gleichzeitige SYCL-Lösung zu besitzen.

3. Externe Bibliotheks- und Compiler-Kompatibilität

• Eigen: Viele Versionen sind inkompatibel mit NVIDIAs nvcc, was Portabilitäts-Schichten betrifft, die darauf angewiesen sind (Kokkos, Alpaka). Es ist zudem derzeit inkompatibel mit OpenMP Offload.
• ROOT: Kann nicht vollständig mit NVIDIAs nvc++ kompiliert werden. Eine hybride Kompilierung (unter Verwendung von g++ und nvc++) ist möglich, erschwert jedoch die Build-Regeln und erfordert ein sorgfältiges Management der Speicherallokation (GPU-Daten müssen durch nc++ alloziert werden).
• Build-Komplexität: Die Unterstützung mehrerer Portabilitäts-Schichten in einem Projekt erfordert komplexe Build-Regeln (z. B. #ifdef-Verwendung, Dateitrennung), um Header-Konflikte zwischen Schichten wie Kokkos und CUDA zu vermeiden.

4. Datenstrukturen und Speicherübertragungen

• Objektorientierte vs. flache Layouts: Komplexe, objektorientierte Datenstrukturen, die in der HEP üblich sind, bilden sich schlecht auf GPUs ab, welche flache Layouts oder Structures of Arrays (SoA) bevorzugen.
• Portabilitäts-Overhead: Schichten wie Kokkos und SYCL bieten Abstraktionen (z. B. Kokkos::Views, sycl::buffers), die den Allokations- und Transfer-Overhead erhöhen, insbesondere bei kleinen Objekten.
• Initialisierung: Die Standard-Speicherinitialisierung in Kokkos::Views kann ein Leistungsengpass sein; dies kann durch Kokkos::ViewAllocateWithoutInitializing gemildert werden.
• Unified Shared Memory (USM): Automatische USM-Übertragungen (verwendet von std::par) wurden als langsamer als explizite Übertragungen beobachtet und werden möglicherweise nicht zum optimalen Zeitpunkt ausgelöst.
• Jagged Arrays: Keines der APIs unterstützt nativ mehrdimensionale Vektoren variabler Größe (Jagged Arrays), was manuelles Erstellen, Dimensionsreduktion oder Padding erfordert.

5. RNGs, FFTs und Atomics

• Native Bibliotheken: Herstellerspezifische Bibliotheken (cuRAND, rocRAND, oneMKL) bieten optimierte Leistung, lassen aber die Portabilität vermissen. Während Kokkos plattformübergreifende RNG/FFT-Implementierungen bereitstellt, erreichen diese möglicherweise nicht die native Performance.
• HEP-CCE Beiträge: Die Gruppe entwickelte eine Header-only Bibliothek, die herstellerspezifische RNGs mit einer einheitlichen API und hardware-nativen Backends für oneMKL umschließt, um es SYCL zu ermöglichen, native Implementierungen über verschiedene Hersteller hinweg aufzurufen.
• Atomics: Atomare Operationen bilden sich schlecht auf die GPU-Hardware ab. Die Leistung variiert stark je nach Compiler und Hardware (z. B. war atomic<int> auf NVIDIA mit OpenMP/clang13 35-mal langsamer als CUDA). Integer-Atomics performen generell schlechter als Float/Double-Atomics.

6. Kompilierungszeit

• Compiler-Overhead: nvc++ ist im Allgemeinen 2- bis 3-mal langsamer als g++. Kokkos fügt einen Overhead von 10 % bis 20 % hinzu.
• Template-lastiger Code: Madgraph, das stark auf hochgradig templatierten, automatisch generierten Code setzt, erlebte mit dem icpx-Backend von SYCL Kompilierungszeiten, die bei großen Kernels über 1000-mal langsamer waren als andere Backends.

7. Runtime-Bereitstellung

• Kokkos: Erfordert das Bauen mit spezifischen Hardware-Identifikatoren. Ein für eine bestimmte GPU (z. B. V100) gebautes Binary kann möglicherweise nicht auf einer anderen (z. B. A100) laufen. Die Unterstützung von Shared Libraries und mehreren Architekturen erzeugt eine komplexe Matrix von Build-Konfigurationen, was die Verteilung von Binaries erschwert.
• Alpaka: Erfordert eine separate Kompilierung für jeden Backend-Compiler.
• SYCL/OneAPI: Multi-Plattform-Binaries entstehen, sind aber oft auf spezifische Architekturen beschränkt oder erfordern eine Identifikation zur Kompilierzeit.
• OpenMP: Einige Implementierungen (LLVM Clang) erlauben es, mehrere Architekturen zur Kompilierzeit aufzulisten und das Backend zur Laufzeit dynamisch auszuwählen.

Bedeutung und Fazit

Das Paper kommt zu dem Schluss, dass keine einzelne GPU-Portabilitätslösung für alle Szenarien optimal ist. Die Wahl der Portabilitäts-Schicht muss durch eine sorgfältige Analyse der spezifischen Charakteristika der Anwendung getrieben werden, einschließlich Kernel-Laufzeit, Nebenläufigkeitsanforderungen, Komplexität der Datenstrukturen und Abhängigkeiten zu externen Bibliotheken.

Die Autoren betonen, dass sich das Ökosystem schnell entwickelt. Sowohl Portabilitäts-Schichten (Kokkos, SYCL, etc.) als auch native Compiler (clang, nvcc, hipcc) verbessern sich stetig, wobei der Trend zu erhöhter Feature-Unterstützung und optimierter Leistung geht. Das Paper plädiert für eine häufige Überwachung dieser Technologien, um die beste Übereinstimmung zwischen Codebasen und Portabilitäts-Schichten sicherzustellen. Zusätzlich merkt das Paper an, dass aufkommende C++-Standards (wie P2300 für asynchrone Ausführung) diese Entscheidungen in Zukunft vereinfachen könnten, obwohl die breite Hardware-Unterstützung noch Jahre entfernt ist.

Letztlich birgt die Auswahl einer Portabilitäts-Schicht ohne vollständiges Verständnis der Art des Codes, seiner externen Abhängigkeiten und seiner Laufzeitumgebung das Risiko erheblicher Probleme hinsichtlich der Leistung und der Flexibilität bei der Bereitstellung.