Kapitel 2: GPU-Programmiermodelle
Grafikprozessoren (GPUs) haben sich von festen Grafikbeschleunigern zu hochgradig parallelen, programmierbaren Rechenmaschinen entwickelt, die in der Lage sind, eine breite Palette von Anwendungen zu beschleunigen. Um Programmierer in die Lage zu versetzen, die massive Parallelität in GPUs effektiv zu nutzen, wurden mehrere parallele Programmiermodelle und APIs entwickelt, wie NVIDIA CUDA, OpenCL und DirectCompute. Diese Programmiermodelle bieten Abstraktionen, die es Programmierern ermöglichen, Parallelität in ihren Anwendungen auszudrücken, während die Low-Level-Details der GPU-Hardware verborgen bleiben.
In diesem Kapitel werden wir die Schlüsselkonzepte und -prinzipien hinter parallelen Programmiermodellen für GPUs untersuchen, wobei wir uns auf das Ausführungsmodell, die Befehlssatzarchitekturen (ISAs) von GPU-Chips, die NVIDIA GPU-ISAs und die Graphics Core Next (GCN)-ISA von AMD konzentrieren. Wir werden auch Beispiele geben, um zu veranschaulichen, wie diese Konzepte in der Praxis angewendet werden.
Ausführungsmodell
Das Ausführungsmodell moderner GPU-Programmiermodelle basiert auf dem Konzept der Kernel, die Funktionen sind, die von einer großen Anzahl von Threads parallel auf der GPU ausgeführt werden. Beim Starten eines Kernels gibt der Programmierer die Anzahl der zu erstellenden Threads und deren Organisation in einer Hierarchie von Gittern, Blöcken (oder kooperativen Thread-Arrays - CTAs) und einzelnen Threads an.
- Ein Gitter repräsentiert den gesamten Problemraum und besteht aus einem oder mehreren Blöcken.
- Ein Block ist eine Gruppe von Threads, die miteinander kooperieren und über gemeinsamen Speicher und Barrieren synchronisieren können. Threads innerhalb eines Blocks werden auf demselben GPU-Kern (genannt Streaming-Multiprozessor oder Recheneinheit) ausgeführt.
- Jeder Thread hat innerhalb seines Blocks und Gitters eine eindeutige ID, die zum Berechnen von Speicheradressen und zum Treffen von Kontrollflussentscheidungen verwendet werden kann.
Diese hierarchische Organisation ermöglicht es Programmierern, sowohl Datenparallelität (bei der dieselbe Operation auf mehrere Datenelemente angewendet wird) als auch Aufgabenparallelität (bei der verschiedene Aufgaben parallel ausgeführt werden) auszudrücken.
FigurHere is the German translation for the provided markdown file, with the code comments translated but the code itself left untranslated:
e 2.1 veranschaulicht die Thread-Hierarchie im GPU-Ausführungsmodell.
Grid
________________
/ / / / /
/ / / / /
/ / / / /
/ / / / /
/__/__/__/__/__/
| | | |
| | Block |
| | | |
Thread Thread ...Abbildung 2.1: Thread-Hierarchie im GPU-Ausführungsmodell.
SIMT-Ausführung
GPU-Programmiermodelle wie CUDA und OpenCL folgen einem Single-Instruction, Multiple-Thread (SIMT)-Ausführungsmodell. Im SIMT-Modell werden Threads in Gruppen ausgeführt, die als Warps (NVIDIA-Terminologie) oder Wavefronts (AMD-Terminologie) bezeichnet werden. Alle Threads innerhalb eines Warps führen die gleiche Anweisung gleichzeitig aus, aber jeder Thread arbeitet mit unterschiedlichen Daten.
Im Gegensatz zum herkömmlichen Single-Instruction, Multiple-Data (SIMD)-Modell, bei dem alle Verarbeitungselemente im Gleichschritt ausgeführt werden, erlaubt SIMT Threads, unabhängige Ausführungspfade zu haben und bei Verzweigungsanweisungen zu divergieren. Wenn ein Warp auf eine Verzweigungsanweisung trifft, wertet die GPU-Hardware die Verzweigungsbedingung für jeden Thread im Warp aus. Wenn alle Threads den gleichen Pfad einschlagen (konvergiert), setzt der Warp die Ausführung normal fort. Wenn einige Threads unterschiedliche Pfade einschlagen (divergiert), wird der Warp in zwei oder mehr Subwarps aufgeteilt, die jeweils einem anderen Pfad folgen. Die GPU-Hardware serialisiert die Ausführung der divergenten Pfade und blendet die inaktiven Threads in jedem Subwarp aus. Wenn alle Pfade abgeschlossen sind, rekombinieren die Subwarps und setzen die Ausführung im Gleichschritt fort.
Abbildung 2.2 veranschaulicht die SIMT-Ausführung mit divergenter Kontrollfluss.
Warp
________________
/ / / / /
/ / / / /
/ / / / /
| | |
| Branch |
| | |
/ \ / \ / \
/ X \ \
/ / \ \ \
/ \ \
/ \ \
/ \ \
/ \ \
\
\
\
ReconvergenceAbbildung 2.2: SIMT-Ausführung mit divergentem Kontrollfluss.
Dieser Divergenzsteuerungsmechanismus ermöglicht es SIMT, flexiblere Kontrollflüsse zu unterstützen.Hier ist die deutsche Übersetzung der Markdown-Datei. Für den Codeabschnitt wurden nur die Kommentare übersetzt, der Code selbst blieb unverändert.
Han SIMD, aber es kommt auf Kosten einer reduzierten SIMD-Effizienz bei Divergenz. Programmierer sollten versuchen, die Divergenz innerhalb eines Warps zu minimieren, um eine optimale Leistung zu erzielen.
Speicherhierarchie
GPUs haben eine komplexe Speicherhierarchie, um die hohe Bandbreite und geringe Latenz paralleler Workloads zu unterstützen. Die Speicherhierarchie besteht in der Regel aus:
- Globaler Speicher: Der größte, aber langsamste Speicherbereich, der von allen Threads in einem Kernel zugänglich ist. Der globale Speicher wird in der Regel mit hochbandbreitigen GDDR- oder HBM-Speichern implementiert.
- Gemeinsamer Speicher: Ein schneller, on-chip-Speicherbereich, der von allen Threads in einem Block gemeinsam genutzt wird. Der gemeinsame Speicher wird für die Kommunikation zwischen Threads und den Datenaustausch innerhalb eines Blocks verwendet.
- Konstanter Speicher: Ein schreibgeschützter Speicherbereich, der zum Senden von schreibgeschützten Daten an alle Threads verwendet wird.
- Texturspeicher: Ein schreibgeschützter Speicherbereich, der für räumliche Lokalität optimiert und über Texturspeichercaches zugegriffen wird. Der Texturspeicher wird häufiger in Grafikanwendungen verwendet.
- Lokaler Speicher: Ein privater Speicherbereich für jeden Thread, der für das Auslagern von Registern und große Datenstrukturen verwendet wird. Der lokale Speicher ist in der Regel auf den globalen Speicher abgebildet.
Eine effektive Nutzung der Speicherhierarchie ist entscheidend für eine hohe Leistung auf GPUs. Programmierer sollten darauf achten, die Verwendung des gemeinsamen Speichers zu maximieren und den Zugriff auf den globalen Speicher zu minimieren, um Speicherlatenz und Engpässe bei der Bandbreite zu reduzieren.
Abbildung 2.3 veranschaulicht die GPU-Speicherhierarchie.
____________
| |
| Global |
| Memory |
____________
|
____________
| |
| Constant |
| Memory |
____________
|
____________
| |
| Texture |
| Memory |
____________
|
|
____________
| |
| Shared |
| Memory |
____________
|
____________
| |
| Local |
| Memory |
____________Abb.Here is the German translation of the markdown file, with the code comments translated but the code itself left untranslated:
GPU-Befehlssatzarchitekturen
GPU-Befehlssatzarchitekturen (Instruction Set Architectures, ISAs) definieren die Niedrigstebene-Schnittstelle zwischen Software und Hardware. Sie legen die unterstützten Befehle, Register und Speicheradressierungsmodi fest. Das Verständnis von GPU-ISAs ist entscheidend für die Entwicklung effizienter GPU-Code und die Optimierung der Leistung.
In diesem Abschnitt werden wir die ISAs zweier großer GPU-Hersteller erkunden: NVIDIA und AMD. Wir werden uns auf NVIDIAs Parallel Thread Execution (PTX) und SASS-ISAs sowie auf AMDs Graphics Core Next (GCN)-ISA konzentrieren.
NVIDIA GPU-ISAs
NVIDIA-GPUs unterstützen zwei Ebenen von ISAs: PTX (Parallel Thread Execution) und SASS (Streaming ASSembler). PTX ist eine virtuelle ISA, die einen stabilen Zielort für CUDA-Compiler bietet, während SASS die native ISA von NVIDIA-GPUs ist.
PTX (Parallel Thread Execution)
PTX ist eine Niedrigstebene-virtuelle ISA, die für NVIDIA-GPUs entwickelt wurde. Sie ähnelt LLVM IR oder Java-Bytecode, da sie ein stabiles, architekturunabhängiges Ziel für Compiler bietet. CUDA-Programme werden in der Regel in PTX-Code kompiliert, der dann vom NVIDIA-GPU-Treiber in die nativen SASS-Anweisungen übersetzt wird.
PTX unterstützt eine breite Palette an arithmetischen, Speicher- und Kontrollflussanweisungen. Es hat eine unbegrenzte Anzahl von virtuellen Registern und unterstützt Prädikation, was eine effiziente Implementierung von Kontrollfluss ermöglicht. PTX bietet auch spezielle Anweisungen für Thread-Synchronisation, atomare Operationen und Textursampling.
Hier ist ein Beispiel für PTX-Code für einen einfachen Vektor-Additions-Kernel:
.version 7.0
.target sm_70
.address_size 64
.visible .entry vecAdd(
.param .u64 vecAdd_param_0,
.param .u64 vecAdd_param_1,
.param .u64 vecAdd_param_2,
.param .u32 vecAdd_param_3
)
{
.reg .b32 %r<4>;
.reg .b64 %rd<6>;
ld.param.u64 %rd1, [vecAdd_param_0];
ld.param.u64 %rd2, [vecAdd_param_1];
ld.param.u64 %rd3, [vecAdd_param_2];
ld.param.u32 %r1, [vecAdd_param_3];
cvta.to.global.u64 %rd4, %rd1;
cvta
}
```Hier ist die deutsche Übersetzung der Markdown-Datei, wobei nur die Kommentare übersetzt wurden und der Code selbst unverändert bleibt:
.to.global.u64 %rd5, %rd2; mov.u32 %r2, %tid.x; mul.wide.u32 %rd6, %r2, 4; add.s64 %rd7, %rd4, %rd6; add.s64 %rd8, %rd5, %rd6;
ld.global.f32 %f1, [%rd7]; ld.global.f32 %f2, [%rd8]; add.f32 %f3, %f1, %f2;
cvta.to.global.u64 %rd9, %rd3; add.s64 %rd10, %rd9, %rd6; st.global.f32 [%rd10], %f3;
ret; }
Diese PTX-Code definiert eine Kernel-Funktion `vecAdd`, die vier Parameter übernimmt: Zeiger auf die Eingabe- und Ausgabevektoren sowie die Größe der Vektoren. Der Kernel berechnet die globale Thread-ID, lädt die entsprechenden Elemente aus den Eingabevektoren, führt die Addition durch und speichert das Ergebnis im Ausgabevektor.
#### SASS (Streaming ASSembler)
SASS ist die native ISA von NVIDIA-GPUs. Es ist eine low-level, maschinespezifische ISA, die direkt auf die GPU-Hardware abgebildet ist. SASS-Anweisungen werden vom NVIDIA-GPU-Treiber aus PTX-Code generiert und sind normalerweise für Programmierer nicht sichtbar.
SASS-Anweisungen sind in einem kompakten Format codiert, um die Speicherbandbreite und den Instruktions-Cache-Verbrauch zu reduzieren. Sie unterstützen eine Vielzahl von Operanden-Typen, einschließlich Register, Immediate-Werte und verschiedene Adressierungsmodi für den Speicherzugriff.
Hier ein Beispiel für SASS-Code für den Vektor-Additions-Kernel:
```sass
code_version_number 90
@P0 LDG.E R2, [R8];
@P0 LDG.E R0, [R10];
@P0 FADD R0, R2, R0;
@P0 STG.E [R12], R0;
EXIT;Dieser SASS-Code entspricht dem zuvor gezeigten PTX-Code. Er lädt die Eingabevektorelemente aus dem globalen Speicher (LDG.E), führt die Addition (FADD) durch, speichert das Ergebnis zurück in den globalen Speicher (STG.E) und beendet den Kernel.
AMD Graphics Core Next ISA
AMD-GPUs verwenden die Graphics Core Next (GCN)-Architektur und -ISA. GCN ist eine RISC-basierte ISA, die sowohl Grafik- als auch Rechenaufgaben unterstützt. Sie ist für hohe Leistung, Skalierbarkeit und Energieeffizienz ausgelegt.
GCN führt mehrere Schlüsselmerkmale ein, wie:
- Eine skalaHere is the German translation of the provided Markdown file, with comments translated but code left untranslated:
alar ALU für die effiziente Ausführung von Skalaroperationen und Flusskontrolle.
- Eine Vektor-ALU für die parallele Ausführung von datenparallelen Operationen.
- Ein Hochleistungs-Speichersystem mit Unterstützung für Atomoperationen und Zugriff mit niedriger Latenz auf gemeinsamen Speicher.
- Ein flexibler Adressierungsmodus für Speicheroperationen, der Base+Offset- und Skalar+Vektor-Adressierung unterstützt.
Hier ist ein Beispiel für GCN ISA-Code für einen Vektor-Additions-Kernel:
.text
.globl vecAdd
.p2align 2
.type vecAdd,@function
vecAdd:
.set DPTR, 0
# Lädt die Kernel-Argumente aus dem Speicher
s_load_dwordx4 s[0:3], s[4:5], 0x0
s_load_dword s4, s[4:5], 0x10
s_waitcnt lgkmcnt(0)
# Berechnung der Adresse für den Vektorzugriff
v_lshlrev_b32 v0, 2, v0
v_add_u32 v1, vcc, s1, v0
v_mov_b32 v3, s3
v_addc_u32 v2, vcc, s2, v3, vcc
flat_load_dword v1, v[1:2]
v_add_u32 v3, vcc, s0, v0
v_mov_b32 v5, s3
v_addc_u32 v4, vcc, s2, v5, vcc
flat_load_dword v0, v[3:4]
# Durchführung der Vektoraddition
v_add_f32 v0, v0, v1
flat_store_dword v[3:4], v0
s_endpgmAbbildung von Algorithmen auf GPU-Architekturen
Das effiziente Abbilden von Algorithmen auf die GPU-Architektur ist entscheidend für eine hohe Leistung. Wichtige Überlegungen sind:
Genügend Parallelismus offenlegen
Der Algorithmus sollte in viele feinkörnige Threads zerlegt werden, die parallel ausgeführt werden können, um die parallelen Verarbeitungsfähigkeiten der GPU vollständig zu nutzen. Dies beinhaltet oft das Identifizieren von datenparallelen Teilen des Algorithmus, die unabhängig voneinander auf verschiedenen Datenelementen ausgeführt werden können.
Minimierung der Zweigdivergenz
Divergenter Kontrollfluss innerhalb eines Warps/einer Wavefront kann zu Serialisierung und reduzierter SIMD-Effizienz führen. Algorithmen sollten so strukturiert werden, dass die Zweigdivergenz möglichst gering ist. Dies kann durch Reduzierung der Verwendung von datenabhängigen### Ausnutzen der Speicherhierarchie
Der Zugriff auf den globalen Speicher ist teuer. Algorithmen sollten die Verwendung von gemeinsamem Speicher und Registern maximieren, um Zugriffe auf den globalen Speicher zu reduzieren. Die Daten sollten auch so im Speicher angeordnet sein, dass die Threads einer Warp aufeinanderfolgende Speicherorte zugreifen können. Eine effektive Nutzung der Speicherhierarchie kann die Speicherlatenzen und Engpässe bei der Bandbreite erheblich reduzieren.
Ausgewogene Berechnung und Speicherzugriffe
Algorithmen sollten ein hohes Verhältnis von arithmetischen Operationen zu Speicheroperationen aufweisen, um Speicherlatenzen zu verstecken und einen hohen Durchsatz an Berechnungen zu erreichen. Dies kann durch Maximierung der Datenwiederverwertung, Vorabrufen von Daten und Überlagerung von Berechnungen mit Speicherzugriffen erreicht werden.
Minimierung der Host-Gerät-Datenübertragungen
Der Datentransfer zwischen Host-Speicher (CPU) und Geräte-Speicher (GPU) ist langsam. Algorithmen sollten solche Übertragungen minimieren, indem möglichst viele Berechnungen auf der GPU durchgeführt werden. Daten sollten in großen Batches auf die GPU übertragen und so lange wie nötig auf dem Gerät gehalten werden, um die Übertragungskosten zu amortisieren.
Bei der Entwicklung von GPU-Kernels werden häufig folgende parallele Algorithmus-Designmuster verwendet:
- Map: Jeder Thread führt die gleiche Operation auf einem anderen Datenelement aus, was die einfache parallele Verarbeitung großer Datensätze ermöglicht.
- Reduce: Parallele Reduktion wird verwendet, um effizient einen einzelnen Wert (z.B. Summe, Maximum) aus einem großen Eingabedatensatz zu berechnen. Threads führen lokale Reduktionen durch, die dann kombiniert werden, um das Endergebnis zu erhalten.
- Scan: Auch als Präfixsumme bekannt, wird Scan verwendet, um die laufende Summe der Elemente in einem Array zu berechnen. Effiziente parallele Scan-Algorithmen sind Schlüsselbausteine für viele GPU-beschleunigte Anwendungen.
- Stencil: Jeder Thread berechnet einen Wert basierend auf benachbarten Datenelementen. Stencil-Berechnungen sind in wissenschaftlichen Simulationen und Bildverarbeitungsanwendungen weit verbreitet.Here is the German translation of the provided Markdown file, with the code comments translated and the code left unchanged:
Anwendungen.
- Gather/Scatter: Threads lesen aus (Gather) oder schreiben auf (Scatter) beliebige Speicherlokationen im globalen Speicher. Sorgfältige Datenorganisation und Zugriffsmuster sind für die Effizienz erforderlich.
Abbildung 3.20 veranschaulicht ein Beispiel für das Map-Muster, bei dem jeder Thread eine Funktion (z.B. Quadratwurzel) auf ein anderes Element des Eingabearrays anwendet.
Eingabearray:
| | | | | | | |
v v v v v v v v
______________________________
Threads: | | | | | | | |
|____|____|____|____|____|____|____|
| | | | | | |
v v v v v v v
Ausgabearray: Abbildung 3.20: Beispiel für das Map-Muster in der GPU-Programmierung.
Schlussfolgerung
GPU-Programmiermodelle wie CUDA und OpenCL machen die parallele Verarbeitungsfähigkeit moderner GPUs für Entwickler zugänglich und ermöglichen es ihnen, eine Vielzahl von Anwendungen zu beschleunigen. Diese Programmiermodelle bieten Abstraktionen, die es ermöglichen, feingranulare parallele Arbeitslasten effizient auf die GPU-Hardware abzubilden.
Das Verständnis des Ausführungsmodells, der Speicherhierarchie und der Synchronisationsprimitiven, die von diesen Programmiermodellen bereitgestellt werden, ist für das Schreiben von GPU-Code mit hoher Leistung unerlässlich. Entwickler müssen sorgfältig Faktoren wie Thread-Organisation, Zweigsteuerungsdivergenzen, Speicherzugriffsmustern und Algorithmendesign berücksichtigen, um die Rechenleistung von GPUs vollständig auszuschöpfen.
Da sich GPU-Architekturen weiterentwickeln, müssen sich auch Programmiermodelle und Tools weiterentwickeln, um Entwickler in die Lage zu versetzen, neue Hardware-Funktionen und -Fähigkeiten effektiv zu nutzen. Laufende Forschung in Bereichen wie Programmiersprachen-Design, Compiler-Optimierung und Autotuning werden entscheidend sein, um die Produktivität von Programmierern und die Leistungsportabilität im Zeitalter des heterogenen Rechnens zu verbessern.