House of ROC: AMDs Alternative zu CUDA

Laufzeitumgebung ROCr

Auf der Kernel-Infrastruktur baut die ROCr-Laufzeitumgebung auf. Sie implementiert nicht nur die Laufzeitspezifikation der HSA-Foundation (Heterogeneous System Architecture), sondern enthält diverse Zusätze für den Multi-GPU-Betrieb – gerade für Deep-Learning-Anwendungen und anspruchsvolle Simulationen ein wichtiges Feature. Des Weiteren bieten die Macher APIs zur "Device Discovery", Nebenläufigkeit von CPU- und GPU-Prozessen und GPU-Multitasking, atomare Speichertransaktionen und Signale sowie User-Space-Warteschlangen und flache Speicheradressierung an.

Entwickler interagieren mit der ROCm-Plattform hauptsächlich über den mitgelieferten Compiler hcc (Heterogenous Compute Compiler). Er basiert auf der LLVM/Clang-Compiler-Infrastruktur und bietet Frontends zur Verarbeitung von Quelltext in OpenCL, HIP und HC, auf die der Artikel später genauer eingehen wird, sowie künftig OpenMP4 zur Programmierung der GPU. Der OpenCL-Support ist erst kürzlich und auf Drängen der Community in ROCm-Version 1.4 hinzugekommen. Im Backend generiert der Compiler native GCN-ISA-Instruktionen, die auf der GPU ausgeführt werden. Die ROCm bietet auch hierfür Werkzeuge zur weiteren Erforschung des Codes (Assembler, Disassembler) und ein offenes Code-Objekt-Format ("hsaco"). Die Beschreibung der GCN-ISA-Spezifikation können Interessierte auf AMDs GPUOpen-Site() einsehen.

Für erfahrene GPU-Programmierer liefern OpenCL und OpenMP Standardansätze zur Beschreibung von Parallelität auf Multi- und Many-Core-Architekturen abseits von CUDA. Heterogeneous-compute Interface for Portability (HIP) und HC sind jedoch Eigenentwicklungen aus dem Haus AMD. HIP ist eine CUDA-nahe Sprache und soll Entwicklern – gegebenenfalls mithilfe dedizierter Werkzeuge – die Überführung von CUDA-Anwendungen und -Bibliotheken in eine auf AMD- und Nvidia-Hardware lauffähige Form erleichtern. HIP selbst bietet C++ in Kernel-Konstrukten an und kennt C++11, C++14 und teilweise C++17. Die Laufzeitprogrammierung auf der CPU geschieht wie bei CUDA über eine C-API, die ebenso in C++ nutzbar ist. HIP unterstützt die meist genutzten CUDA-Features wie Speichermanagement auf dem Device, Streams zur Verwaltung von CPU-GPU-Nebenläufigkeit inklusive Ereignissen zur Synchronisation sowie die Profiling-API. Die HIP-Kompilate erfreuen sich einer hundertprozentigen Tool-Unterstützung aller Profiler und Debugger aus dem CUDA- wie dem dem ROCm-Universum.

Benchmark zum Vergleich

Zur Illustration der weiteren Ausführungen soll der quelloffene BabelStream-Benchmark der Universität Bristol dienen. Er nutzt C++11 und baut auf dem klassischen HPC-Speicherbandbreiten-Benchmark STREAM von John McAlpin auf. BabelStream implementiert die vier Vektor-Operationen (copy, multiply, add und triad) sowie das Skalarprodukt (dot):

c[:]    = a[:]               /* copy  */
b[:] = scalar*b[:] /* mul */
c[:] = a[:] + b[:] /* add */
a[:] = b[:] + scalar*c[:] /* triad */
scalar = dot(a[:],b[:]) /* dot */

Der Benchmark führt diese Operationen wiederholt auf synthetisch gefüllten Feldern aus, deren Größe durch die Nutzer konfigurierbar ist. Die Messung der Laufzeit der Operationen zeigt nicht nur die Werte der Speicherbandbreite auf der GPU, sondern auch der Effizienz eines Programmierparadigmas und der benutzten Compiler-Infrastruktur.

Der add-Kernel in der CUDA-Implementierung
sieht folgendermaßen aus:

__global__ void add_kernel(const T * a,
const T * b,
T * c){
const int i = blockDim.x * blockIdx.x + threadIdx.x;
c[i] = a[i] + b[i];}
void CUDAStream<T>::add(){
add_kernel<<<array_size/TBSIZE, TBSIZE>>>(d_a, d_b, d_c);
check_error(); //..
}

Wie das Listing zeigt, bildet BabelStream den Gleitkomma-Datentyp des Benchmarks als Template ab, um die Genauigkeit ohne zusätzlichen Code wechseln zu können. Die achte Zeile (add_kernel)zeigt den Aufruf des Device-Kernel in vereinfachter Form. Die drei Zeiger d_a, d_b, d_c sind dabei allozierte Felder auf der GPU mit der Anzahl array_size an Elementen. CUDA verlangt für jeden Kernel-Aufruf eine doppelte, virtuelle Partitionierung des Indexraumes aller zu bearbeitenden Elemente. In diesem Fall ist der Indexraum eindimensional, der in array_size/TBSIZE Blöcke und TBSIZE Threads pro Block zerlegt wird -- beides sind ganzzahlige Skalare. Diese Notation umgeht das explizite Angeben einer for-Schleife über den Indexraum. Es erzwingt aber das manuelle Zusammensetzen des Index, wie in der vierten Zeile zu sehen ist. Dabei sind blockDim, blockIdx und threadIdx C-Strukturen, die die CUDA-Laufzeit-API zur Verfügung stellt.

HIP-pe Übersetzung

Die Übersetzung des obigen Listings mit dem Kommandozeilenwerkzeug hipify der ROCm-Plattform erzeugt folgenden HIP-Code:

__global__ void add_kernel(hipLaunchParm lp,
                           const T * a, const T * b,
T * c){
const int i = hipBlockDim_x * hipBlockIdx_x + hipThreadIdx_x;
c[i] = a[i] + b[i];
}
void HIPStream<T>::add(){
hipLaunchKernel(HIP_KERNEL_NAME(add_kernel),
dim3(array_size/TBSIZE), dim3(TBSIZE), 0, 0,
d_a, d_b, d_c); check_error(); //...
}

Das Listing ist strukturell identisch zu dem vorherigen. Das Tool hat die Laufzeit-Structs umbenannt (hipBlockDim_x, hipBlockIdx_x, hipThreadIdx_x). Den Aufruf und die Signatur des Kernel hat es durch HIP-spezifische Boilerplate-Argumente ergänzt, um eine vergleichbare Funktionsweise zu CUDA zu erreichen und dennoch Code zu generieren, den nvcc und hcc verarbeiten können.

HIP ist damit ein interessantes und ernst zu nehmendes Werkzeug zur Konvertierung von Legacy-CUDA-Anwendungen und sollte AMD ermöglichen, schnell auf AMD-GPUs lauffähige Projekte vorzeigen zu können. HIP-Code kann zudem mit einem Ökosystem spezialisierter Bibliotheken wie hipBlas, hipFFT und hipRNG interagieren.