Hoofdstuk 2: GPU Programmeermodellen
Grafische Verwerkingseenheden (GPU's) zijn geëvolueerd van vaste-functie grafische versnellers naar hoogst parallelle, programmeerbare rekenkundige motoren die in staat zijn om een breed scala aan toepassingen te versnellen. Om programmeurs in staat te stellen om de massieve parallelliteit in GPU's effectief te benutten, zijn er verschillende parallelle programmeermodellen en API's ontwikkeld, zoals NVIDIA CUDA, OpenCL en DirectCompute. Deze programmeermodellen bieden abstracties die programmeurs in staat stellen om parallellisme in hun toepassingen uit te drukken, terwijl de lage-niveau details van de GPU-hardware verborgen blijven.
In dit hoofdstuk zullen we de belangrijkste concepten en principes achter parallelle programmeringsmodellen voor GPU's verkennen, met de nadruk op het uitvoeringsmodel, GPU-instructieset-architecturen (ISA's), NVIDIA GPU-ISA's en AMD's Graphics Core Next (GCN) ISA. We zullen ook voorbeelden geven om te illustreren hoe deze concepten in de praktijk worden toegepast.
Uitvoeringsmodel
Het uitvoeringsmodel van moderne GPU-programmeermodellen is gebaseerd op het concept van kernels, die functies zijn die parallel worden uitgevoerd door een groot aantal threads op de GPU. Bij het opstarten van een kernel specificeert de programmeur het aantal threads dat moet worden gemaakt en hoe ze zijn georganiseerd in een hiërarchie van grids, blokken (of coöperatieve thread-arrays - CTAs) en individuele threads.
- Een grid vertegenwoordigt de hele probleemruimte en bestaat uit één of meer blokken.
- Een blok is een groep threads die kunnen samenwerken en synchroniseren met elkaar via gedeeld geheugen en barrières. Threads binnen een blok worden uitgevoerd op dezelfde GPU-kern (een streaming multiprocessor of rekenunit genoemd).
- Elke thread heeft een unieke ID binnen zijn blok en grid, die kan worden gebruikt om geheugen-adressen te berekenen en controle-flow beslissingen te nemen.
Deze hiërarchische organisatie stelt programmeurs in staat om zowel data-parallellisme (waar dezelfde bewerking wordt toegepast op meerdere data-elementen) als taak-parallellisme (waar verschillende taken parallel worden uitgevoerd) uit te drukken.
FiguurHier is de Nederlandse vertaling van het Markdown-bestand, waarbij alleen de opmerkingen in de code worden vertaald:
e 2.1 illustreert de thread-hiërarchie in het GPU-uitvoeringsmodel.
Grid
________________
/ / / / /
/ / / / /
/ / / / /
/ / / / /
/__/__/__/__/__/
| | | |
| | Block |
| | | |
Thread Thread ...Figuur 2.1: Thread-hiërarchie in het GPU-uitvoeringsmodel.
SIMT-uitvoering
GPU-programmeringsmodellen zoals CUDA en OpenCL volgen een Single-Instruction, Multiple-Thread (SIMT) uitvoeringsmodel. In het SIMT-model worden threads uitgevoerd in groepen die warps (NVIDIA-terminologie) of wavefronts (AMD-terminologie) worden genoemd. Alle threads binnen een warp voeren dezelfde instructie tegelijkertijd uit, maar elke thread werkt op verschillende gegevens.
In tegenstelling tot het traditionele Single-Instruction, Multiple-Data (SIMD)-model, waarbij alle verwerkingselementen synchroon worden uitgevoerd, staat SIMT threads toe om onafhankelijke uitvoerpaden te hebben en af te wijken bij takkinstructies. Wanneer een warp een takkinstructie tegenkomt, evalueert de GPU-hardware de takcondities voor elke thread in de warp. Als alle threads hetzelfde pad kiezen (geconvergeerd), gaat de warp normaal verder met de uitvoering. Als sommige threads verschillende paden kiezen (divergeren), wordt de warp opgesplitst in twee of meer subwarps, waarbij elk een ander pad volgt. De GPU-hardware serialiseert de uitvoering van de divergente paden en maskeert inactieve threads in elke subwarp. Wanneer alle paden zijn voltooid, reconvergeren de subwarps en gaan ze weer synchroon verder.
Figuur 2.2 illustreert SIMT-uitvoering met divergente controleflow.
Warp
________________
/ / / / /
/ / / / /
/ / / / /
| | |
| Tak |
| | |
/ \ / \ / \
/ X \ \
/ / \ \ \
/ \ \
/ \ \
/ \ \
/ \ \
\
\
\
ReconvergentieFiguur 2.2: SIMT-uitvoering met divergente controleflow.
Deze divergentie-afhandelingsmechanisme stelt SIMT in staat om meer flexibele controleflow tHier is de Nederlandse vertaling van het Markdown-bestand, waarbij alleen de opmerkingen in de code zijn vertaald:
Geheugenhi??rarchie
GPU's hebben een complexe geheugenhi??rarchie om te voldoen aan de hoge bandbreedte- en lage latentievereisten van parallelle werkbelastingen. De geheugenhi??rarchie bestaat doorgaans uit:
- Globaal geheugen: De grootste maar traagste geheugenruimte, toegankelijk voor alle threads in een kernel. Globaal geheugen wordt meestal ge??mplementeerd met behulp van high-bandwidth GDDR- of HBM-geheugen.
- Gedeeld geheugen: Een snel, on-chip geheugenruimte die door alle threads in een blok wordt gedeeld. Gedeeld geheugen wordt gebruikt voor communicatie tussen threads en gegevensuitwisseling binnen een blok.
- Constant geheugen: Een alleen-lezen geheugenruimte die wordt gebruikt voor het uitzenden van alleen-lezen gegevens naar alle threads.
- Tekstuurgeheugen: Een alleen-lezen geheugenruimte die is geoptimaliseerd voor ruimtelijke lokaliteit en wordt benaderd via tekstuurcaches. Tekstuurgeheugen wordt vaker gebruikt in grafische werkbelastingen.
- Lokaal geheugen: Een privégeheugenruimte voor elke thread, gebruikt voor registeruitwisseling en grote gegevensstructuren. Lokaal geheugen wordt meestal toegewezen aan het globale geheugen.
Effectief gebruik van de geheugenhi??rarchie is cruciaal voor het behalen van hoge prestaties op GPU's. Programmeurs moeten ernaar streven om het gebruik van gedeeld geheugen te maximaliseren en toegangen tot globaal geheugen te minimaliseren om geheugenlatentie en bandbreedte-knelpunten te verminderen.
Figuur 2.3 illustreert de GPU-geheugenhi??rarchie.
____________
| |
| Globaal |
| Geheugen |
____________
|
____________
| |
| Constant |
| Geheugen |
____________
|
____________
| |
| Tekstuur |
| Geheugen |
____________
|
|
____________
| |
| Gedeeld |
| Geheugen |
____________
|
____________
| |
| Lokaal |
| Geheugen |
____________FigVertaling naar het Nederlands:
GPU-instructieset-architecturen
GPU-instructieset-architecturen (ISA's) definiëren de laag-niveau-interface tussen software en hardware. Ze specificeren de instructies, registers en geheugenadressering die door de GPU worden ondersteund. Inzicht in GPU-ISA's is essentieel voor het ontwikkelen van efficiënte GPU-code en het optimaliseren van prestaties.
In deze sectie zullen we de ISA's van twee grote GPU-leveranciers verkennen: NVIDIA en AMD. We zullen ons richten op NVIDIA's Parallel Thread Execution (PTX) en SASS-ISA's, en AMD's Graphics Core Next (GCN) ISA.
NVIDIA GPU-ISA's
NVIDIA-GPU's ondersteunen twee niveaus van ISA's: PTX (Parallel Thread Execution) en SASS (Streaming ASSembler). PTX is een virtuele ISA die een stabiel doelwit biedt voor CUDA-compilers, terwijl SASS de native ISA van NVIDIA-GPU's is.
PTX (Parallel Thread Execution)
PTX is een low-level, virtuele ISA ontworpen voor NVIDIA-GPU's. Het lijkt op LLVM IR of Java-bytecode in die zin dat het een stabiel, architectuur-onafhankelijk doelwit biedt voor compilers. CUDA-programma's worden meestal gecompileerd naar PTX-code, die vervolgens wordt vertaald naar de native SASS-instructies door de NVIDIA GPU-driver.
PTX ondersteunt een breed scala aan rekenkundige, geheugen- en controlestroominstructies. Het heeft een onbeperkt aantal virtuele registers en ondersteunt predikatie, waardoor een efficiënte implementatie van controlestromen mogelijk is. PTX biedt ook speciale instructies voor draadsynchronisatie, atomaire bewerkingen en textuursampling.
Hier is een voorbeeld van PTX-code voor een eenvoudige vector-optellingskernel:
.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>;
// Laad parameters uit geheugen
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];
// Converteer adressen naar globale adressen
cvta.to.global.u64 %rd4, %rd1;
cvta
```Hier is de Nederlandse vertaling van het gegeven bestand, met de opmerkingen vertaald, maar zonder vertaling van de code-delen:
// Kopieer pointer naar globaal geheugen als 64-bit waarde naar %rd5, vanuit %rd2 .to.global.u64 %rd5, %rd2; // Laad thread-ID in X-dimensie in %r2 mov.u32 %r2, %tid.x; // Bereken offset in bytes op basis van thread-ID, vermenigvuldigd met 4 (grootte van float) mul.wide.u32 %rd6, %r2, 4; // Bereken pointer naar element in eerste invoervector, door basis-pointer %rd4 op te tellen met offset %rd6 add.s64 %rd7, %rd4, %rd6; // Bereken pointer naar element in tweede invoervector, door basis-pointer %rd5 op te tellen met offset %rd6 add.s64 %rd8, %rd5, %rd6;
// Laad element uit eerste invoervector in %f1 ld.global.f32 %f1, [%rd7]; // Laad element uit tweede invoervector in %f2 ld.global.f32 %f2, [%rd8]; // Bereken som van elementen in %f1 en %f2, opslaan in %f3 add.f32 %f3, %f1, %f2;
// Kopieer pointer naar globaal geheugen als 64-bit waarde naar %rd9, vanuit %rd3 cvta.to.global.u64 %rd9, %rd3; // Bereken pointer naar element in uitvoervector, door basis-pointer %rd9 op te tellen met offset %rd6 add.s64 %rd10, %rd9, %rd6; // Schrijf waarde in %f3 naar geheugen op pointer %rd10 st.global.f32 [%rd10], %f3;
// Functie afsluiten ret; }
Dit PTX-code definieert een kernelfunctie `vecAdd` die vier parameters accepteert: pointers naar de invoervectoren en de uitvoervector, en de grootte van de vectoren. De kernel berekent de globale thread-ID, laadt de bijbehorende elementen uit de invoervectoren, voert de optelling uit, en slaat het resultaat op in de uitvoervector.
#### SASS (Streaming ASSembler)
SASS is de native ISA van NVIDIA-GPU's. Het is een laag-niveau, machine-specifieke ISA die direct overeenkomt met de GPU-hardware. SASS-instructies worden gegenereerd door de NVIDIA GPU-driver vanuit PTX-code en zijn meestal niet zichtbaar voor programmeurs.
SASS-instructies zijn gecodeerd in een compact formaat om bandbreedte en instructie-cache verbruik te verminderen. Ze ondersteunen een breed scala aan operandtypes, waaronder registers, directe waarden en verschillende adressermodi voor geheugentoegang.
Hier is een voorbeeld van SASS-code voor de vector-optel-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;Deze SASS-code komt overeen met de eerdere PTX-code. Het laadt de invoervector-elementen uit het globale geheugen (LDG.E), voert de optelling uit (FADD), slaat het resultaat terug naar het globale geheugen op (STG.E) en sluit de kernel af.
AMD Graphics Core Next ISA
AMD-GPU's gebruiken de Graphics Core Next (GCN) architectuur en ISA. GCN is een RISC-gebaseerde ISA die zowel graphics- als rekenkundige werklasten ondersteunt. Het is ontworpen voor hoge prestaties, schaalbaarheid en energie-efficiëntie.
GCN introduceert verschillende belangrijke kenmerken, zoals:
- Een schaal...
alar ALU voor efficiënte uitvoering van scalaire bewerkingen en stroombesturing.
- Een vector ALU voor parallelle uitvoering van data-parallelle bewerkingen.
- Een geheugensubsysteem met hoge bandbreedte met ondersteuning voor atomaire bewerkingen en lage latentie-toegang tot gedeeld geheugen.
- Een flexibele adressingsmodus voor geheugenoperaties, met ondersteuning voor basis+offset- en scalar+vector-adressering.
Hier is een voorbeeld van GCN ISA-code voor een vector-additiekernel:
```asm
.text
.globl vecAdd
.p2align 2
.type vecAdd,@function
vecAdd:
.set DPTR, 0
# Laad kernel-argumenten uit het geheugen
s_load_dwordx4 s[0:3], s[4:5], 0x0
s_load_dword s4, s[4:5], 0x10
s_waitcnt lgkmcnt(0)
# Bereid de geheugenadressering voor
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]
# Voer de vector-additie uit
v_add_f32 v0, v0, v1
flat_store_dword v[3:4], v0
s_endpgmDeze GCN-code laadt de invoervectorelementent met behulp van flat_load_dword, voert de additie uit met v_add_f32 en slaat het resultaat weer op in het geheugen met flat_store_dword. De s_load_dwordx4 en s_load_dword-instructies worden gebruikt om de kernel-argumenten uit het geheugen te laden.
Algoritmen toewijzen aan GPU-architecturen
Het efficiënt toewijzen van algoritmen aan de GPU-architectuur is cruciaal voor het bereiken van hoge prestaties. Belangrijke overwegingen zijn:
Voldoende parallellisme blootleggen
Het algoritme moet worden opgesplitst in veel fijnmazige threads die gelijktijdig kunnen worden uitgevoerd om de parallelle verwerkingscapaciteit van de GPU volledig te benutten. Dit betekent vaak het identificeren van data-parallelle delen van het algoritme die onafhankelijk op verschillende gegevenselementen kunnen worden uitgevoerd.
Divergentie van vertakkingen minimaliseren
Divergente controlestromen binnen een warp/golffront kunnen leiden tot seriële verwerking en een verminderde SIMD-efficiëntie. Algoritmen moeten zodanig worden gestructureerd dat de divergentie van vertakkingen zoveel mogelijk wordt beperkt. Dit kan worden bereikt door het gebruik van datagebaseerde### Exploiteren van geheuërarchie
Het openen van het globale geheugen is duur. Algoritmen moeten het gebruik van gedeeld geheugen en registers maximaliseren om openingen tot het globale geheugen te verminderen. Gegevens moeten ook in het geheugen worden opgeslagen om gecoördineerde geheugenopeninges mogelijk te maken, waarbij threads in een warp aaneengesloten geheugenlocaties openen. Effectief gebruik van de geheuërarchie kan de geheugenlatentie en bandbreedte bottlenecks aanzienlijk verminderen.
Balanceren van berekening en geheugenopenings
Algoritmen moeten een hoge verhouding hebben van rekenkundige bewerkingen tot geheugenoperations om geheugenlatentie effectief te verbergen en een hoge computationele doorvoer te bereiken. Dit kan worden bereikt door gegevens maximaal te hergebruiken, gegevens vooruit te halen en berekeningen te overlappen met geheugenopeninges.
Minimaliseren van host-apparaat gegevensoverdrachten
Het overdragen van gegevens tussen host (CPU) en apparaat (GPU) geheugen is langzaam. Algoritmen moeten dergelijke overdrachten minimaliseren door zoveel mogelijk berekeningen op de GPU uit te voeren. Gegevens moeten in grote batches naar de GPU worden overgebracht en zo lang mogelijk op het apparaat worden gehouden om de overdrachtoverhead af te schrijven.
Verschillende parallelle algoritme-ontwerppatronen worden vaak gebruikt bij het ontwikkelen van GPU-kernels:
- Map: Elke thread voert dezelfde bewerking uit op een verschillend gegevensenlement, waardoor eenvoudige parallelle verwerking van grote datasets mogelijk is.
- Reduce: Parallelle reductie wordt gebruikt om efficiënt één waarde (bv. som, maximum) te berekenen uit een grote invoerdataset. Threads voeren lokale reducties uit, die vervolgens worden gecombineerd om het eindresultaat te produceren.
- Scan: Ook bekend als voorvoegselsom, scan wordt gebruikt om de lopende som van elementen in een array te berekenen. Efficiënte parallelle scan-algoritmen zijn sleutelcomponenten voor veel GPU-versnelde toepassingen.
- Stencil: Elke thread berekent een waarde op basis van naburige gegevenselementen. Stencil-berekeningen komen veel voor in wetenschappelijke simulaties en beeldverwerkingstoepassingen.Hier is de Nederlandse vertaling van de Markdown-file, waarbij alleen de opmerkingen zijn vertaald, niet de code:
Toepassingen.
- Verzamelen/Verspreiden: Threads lezen van (verzamelen) of schrijven naar (verspreiden) willekeurige locaties in het globale geheugen. Zorgvuldige data-indeling en toegangspatronen zijn vereist voor efficiëntie.
Figuur 3.20 illustreert een voorbeeld van het map-patroon, waarbij elke thread een functie (bijv. wortel) toepast op een ander element van de invoerarray.
Invoerarray:
| | | | | | | |
v v v v v v v v
______________________________
Threads: | | | | | | | |
|____|____|____|____|____|____|____|
| | | | | | |
v v v v v v v
Uitvoerarray: Figuur 3.20: Voorbeeld van het map-patroon in GPU-programmering.
Conclusie
GPU-programmeermodellen zoals CUDA en OpenCL stellen ontwikkelaars in staat om de parallelle verwerkingsmogelijkheden van moderne GPU's te benutten, waardoor ze een breed scala aan toepassingen kunnen versnellen. Deze programmeermodellen bieden abstracties waarmee fijnmazige parallelle workloads efficiënt kunnen worden toegewezen aan de GPU-hardware.
Het begrijpen van het uitvoeringsmodel, de geheugenhi??rarchie en de synchronisatieprimitieven die door deze programmeermodellen worden geboden, is essentieel voor het schrijven van GPU-code met hoge prestaties. Ontwikkelaars moeten zorgvuldig rekening houden met factoren zoals threadorganisatie, vertakkingsdivergentie, geheugentoeganspatronen en algoritme-ontwerp om de rekenkracht van GPU's volledig te benutten.
Naarmate GPU-architecturen blijven evolueren, moeten ook programmeermodellen en hulpprogramma's worden verbeterd om ontwikkelaars in staat te stellen om nieuwe hardwarefuncties en -mogelijkheden effectief te benutten. Voortdurend onderzoek op gebieden als ontwerp van programmeertalen, compileroptimalisatie en autotuning zullen cruciaal zijn voor het verbeteren van de productiviteit van programmeurs en de prestatiebaarheid in het tijdperk van heterogene berekeningen.