Hoofdstuk 3: Parallelle programmeringsmodellen in GPU-ontwerp
Grafische verwerkingseenheden (GPU's) zijn geëvolueerd van vaste-functie grafische versnellers tot zeer parallelle, programmeerbare rekenmachines die in staat zijn een breed scala aan toepassingen te versnellen. Om programmeurs in staat te stellen de enorme parallellisme in GPU's effectief te benutten, zijn er verschillende parallelle programmeringsmodellen en API's ontwikkeld, zoals NVIDIA CUDA, OpenCL en DirectCompute. Deze programmeringsmodellen bieden abstracties waarmee programmeurs parallellisme in hun toepassingen kunnen uitdrukken, terwijl de low-level 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 een focus op het SIMT (Single Instruction, Multiple Thread) uitvoeringsmodel, het CUDA-programmeringsmodel en -API's, en het OpenCL-framework. We zullen ook technieken bespreken voor het in kaart brengen van algoritmen naar GPU-architecturen om hoge prestaties en efficiëntie te bereiken.
SIMT (Single Instruction, Multiple Thread) uitvoeringsmodel
Het SIMT-uitvoeringsmodel is het fundamentele paradigma dat door moderne GPU's wordt gebruikt om massieve parallellisme te bereiken. In het SIMT-model voeren een groot aantal threads hetzelfde programma (een kernel genoemd) parallel uit, maar heeft elke thread zijn eigen programmateller en kan op basis van zijn thread-ID en de gegevens waarop het werkt, verschillende uitvoeringsroutes volgen.
Kernels en thread-hiërarchie
Een GPU-kernel is een functie die parallel wordt uitgevoerd door een groot aantal threads. Bij het lanceren van een kernel geeft de programmeur het aantal threads op dat moet worden gemaakt en hoe ze worden georganiseerd in een hiërarchie van grids, blokken (of "cooperative thread arrays" - CTAs) en individuele threads.
- Een grid vertegenwoordigt de volledige probleemruimte en bestaat uit een of meer blokken.
- Een blok is een groep threads die kunnen samenwerken en synchroniseren via gedeeld geheugen en barrières. Threads binnen een blok worden uitgevoerd op dezelfde GPU-core (een "streaming multiprocessor" genoemd).Dit is de Nederlandse vertaling van de Markdown-bestanden, waarbij de code-opmerkingen zijn vertaald, maar de code zelf niet:
sor of rekenunit).
- Elke thread heeft een unieke ID binnen zijn blok en raster, die gebruikt kan worden om geheugenaddressen te berekenen en beslissingen over de controleflow te nemen.
Deze hiërarchische organisatie stelt programmeurs in staat zowel data-parallelisme (waar dezelfde bewerking wordt toegepast op meerdere gegevenselementen) als taak-parallelisme (waar verschillende taken parallel worden uitgevoerd) uit te drukken.
Figuur 3.1 illustreert de thread-hiërarchie in het SIMT-uitvoeringsmodel.
Raster
________________
/ / / / /
/ / / / /
/ / / / /
/ / / / /
/__/__/__/__/__/
| | | |
| | Blok |
| | | |
Thread Thread ...Figuur 3.1: Thread-hiërarchie in het SIMT-uitvoeringsmodel.
SIMT-uitvoering
In het SIMT-uitvoeringsmodel voert elke thread dezelfde instructie uit, maar werkt op verschillende gegevens. In tegenstelling tot SIMD (Single Instruction, Multiple Data) waar alle verwerkingselementen in lockstep werken, staat SIMT threads toe onafhankelijke uitvoeringspaden te hebben en af te wijken bij vertakkinginstructies.
Wanneer een warp (een groep van 32 threads in NVIDIA GPU's of 64 threads in AMD GPU's) een vertakkinginstructie tegenkomt, evalueert de GPU-hardware de vertakkingsvoorwaarde voor elke thread in de warp. Als alle threads hetzelfde pad volgen (geconvergeerd), gaat de warp normaal verder met de uitvoering. Als echter sommige threads verschillende paden volgen (divergeren), wordt de warp gesplitst in twee of meer subwarps, waarbij elk een ander pad volgt. De GPU-hardware serialiseert de uitvoering van de uiteenlopende paden, waarbij de inactieve threads in elke subwarp worden gemaskeerd. Wanneer alle paden zijn voltooid, convergeren de subwarps opnieuw en gaan ze verder in lockstep.
Figuur 3.2 illustreert SIMT-uitvoering met afwijkende controleflow.
Warp
________________
/ / / / /
/ / / / /
/ / / / /
| | |
| Vertakking |
| | |
/ \ / \ / \
/ X \ \
/ / \ \ \
/ \ \
/ \ \
/ \ \
/ \ \
\
```Reconvergentie
Figuur 3.2: SIMT-uitvoering met divergente controle stroom.
Dit mechanisme voor het afhandelen van divergentie stelt SIMT in staat om flexibeler controlestromen te ondersteunen dan SIMD, maar dit gaat ten koste van een verminderde SIMD-efficiëntie wanneer divergentie optreedt. Programmeurs moeten ernaar streven om divergentie binnen een warp te minimaliseren om een optimale prestatie te bereiken.
Geheugenopbouw
GPU's hebben een complexe geheugenopbouw om te voldoen aan de hoge bandbreedte- en lage latentie-eisen van parallelle werkbelastingen. De geheugenopbouw bestaat meestal uit:
- Globaal geheugen: De grootste maar traagste geheugenruimte, toegankelijk voor alle threads in een kernel. Globaal geheugen wordt typisch geïmplementeerd met behulp van high-bandwidth GDDR- of HBM-geheugen.
- Gedeeld geheugen: Een snel, on-chip geheugenruimte die gedeeld wordt door alle threads in een blok. Gedeeld geheugen wordt gebruikt voor communicatie tussen threads en gegevensuitwisseling binnen een blok.
- Constant geheugen: Een alleen-lezen geheugenruimte gebruikt voor het uitzenden van alleen-lezen gegevens naar alle threads.
- Textuurgeheugen: Een alleen-lezen geheugenruimte geoptimaliseerd voor ruimtelijke lokaliteit en toegankelijk via textuurcaches. Textuurgeheugen wordt vaker gebruikt in grafische werkbelastingen.
- Lokaal geheugen: Een privé-geheugenruimte voor elke thread, gebruikt voor registerafgifte en grote gegevensstructuren. Lokaal geheugen wordt meestal toegewezen aan globaal geheugen.
Effectief gebruik van de geheugenopbouw is cruciaal voor het bereiken 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 3.3 illustreert de GPU-geheugenopbouw.
| Gedeeld |
| Geheugen |
____________
|
____________
| |
| Lokaal |
| Geheugen |
____________Figuur 3.3: GPU geheugenhi??rarchie.
CUDA Programmeermodel en API's
CUDA (Compute Unified Device Architecture) is een parallel computing platform en programmeermodel ontwikkeld door NVIDIA voor algemeen gebruik op GPU's. CUDA biedt een set uitbreidingen op standaardprogrammeertalen, zoals C, C++ en Fortran, die programmeurs in staat stellen om parallellisme uit te drukken en gebruik te maken van de rekenkracht van NVIDIA GPU's.
CUDA Programmeermodel
Het CUDA-programmeermodel is gebaseerd op het concept van kernels, die functies zijn die parallel worden uitgevoerd door een groot aantal threads op de GPU. De programmeur specificeert het aantal threads dat moet worden gelanceerd en hun organisatie in een grid van thread blocks.
CUDA introduceert enkele belangrijke abstracties om parallel programmeren te vergemakkelijken:
- Thread: De basiseenheid van uitvoering in CUDA. Elke thread heeft zijn eigen programmateller, registers en lokaal geheugen.
- Block: Een groep threads die kunnen samenwerken en synchroniseren met elkaar. Threads binnen een block worden uitgevoerd op dezelfde streaming multiprocessor en kunnen communiceren via shared memory.
- Grid: Een verzameling thread blocks die dezelfde kernel uitvoeren. Het grid vertegenwoordigt de hele probleemruimte en kan een-, twee- of driedimensionaal zijn.
CUDA biedt ook ingebouwde variabelen (bijvoorbeeld threadIdx, blockIdx, blockDim, gridDim) waarmee threads zichzelf kunnen identificeren en geheugenaddressen kunnen berekenen op basis van hun positie in de thread-hi??rarchie.
Figuur 3.4 illustreert het CUDA-programmeermodel.
Grid
________________
/ / / / /
/ / / / /
/ / / / /
/ / / / /
/__/__/__/__/__/
| | | |
| | Block |
| | | |
Thread Thread ...Figuur 3.4: CUDA-programmeermodel.
CUDA GeheugenCuda exposeert de GPU-geheugenhi??rarchie aan de programmeur, waardoor expliciete controle over gegevensplaatsing en -beweging mogelijk is. De belangrijkste geheugenruimtes in Cuda zijn:
- Globaal geheugen: Toegankelijk voor alle threads in een kernel en blijft bestaan tussen kernel-starts. Het globale geheugen heeft de hoogste latentie en wordt typisch gebruikt voor grote datastructuren.
- Gedeeld geheugen: Een snel, on-chip geheugen dat wordt gedeeld door alle threads in een blok. 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. Constant geheugen wordt gecacht en biedt lage latentie toegang.
- Textuurgeheugen: Een alleen-lezen geheugenruimte die is geoptimaliseerd voor ruimtelijke lokaliteit en wordt benaderd via textuurcaches. Textuurgeheugen wordt meer gebruikt in grafische werkbelastingen.
- Lokaal geheugen: Een privaat geheugenruimte voor elke thread, gebruikt voor registerspilling en grote datastructuren. Lokaal geheugen wordt doorgaans toegewezen aan het globale geheugen.
Programmeurs kunnen gegevens toewijzen en overdragen tussen het host- (CPU-) en apparaat- (GPU-) geheugen met behulp van CUDA-runtime-API's, zoals cudaMalloc, cudaMemcpy en cudaFree.
Figuur 3.5 illustreert de CUDA-geheugenhi??rarchie.
____________
| |
| Global |
| Memory |
____________
|
____________
| |
| Constant |
| Memory |
____________
|
____________
| |
| Texture |
| Memory |
____________
|
|
____________
| |
| Shared |
| Memory |
____________
|
____________
| |
| Local |
| Memory |
____________Figuur 3.5: CUDA-geheugenhi??rarchie.
CUDA-synchronisatie en -co??rdinatie
CUDA biedt synchronisatie- en co??rdinatieprimities om samenwerking en communicatie tussen threads mogelijk te maken:
- Barrière-synchronisatie: De __syncthreadHere is the Dutch translation of the markdown file, with the comments in the code translated:
De s() -functie fungeert als een barrière die ervoor zorgt dat alle threads in een blok hetzelfde punt hebben bereikt voordat ze verder gaan.
- Atomaire bewerkingen: CUDA ondersteunt atomaire bewerkingen (bijvoorbeeld atomicAdd, atomicExch) waarmee threads lees-wijzig-schrijfbewerkingen kunnen uitvoeren op gedeeld of globaal geheugen zonder interferentie van andere threads.
- Warp-level primitieven: CUDA biedt warp-level intrinsieke functies (bijvoorbeeld __shfl, __ballot) die efficiënte communicatie en synchronisatie binnen een warp mogelijk maken.
Het juiste gebruik van synchronisatie- en coördinatieprimitivenen is essentieel voor het schrijven van correcte en efficiënte parallelle programma's in CUDA.
Voorbeeld 3.1 toont een eenvoudige CUDA-kernel die vectoroptelling uitvoert.
__global__ void vectorAdd(int *a, int *b, int *c, int n) {
// De index van de huidige thread berekenen
int i = blockIdx.x * blockDim.x + threadIdx.x;
// Controleren of de thread binnen het bereik van de vector valt
if (i < n) {
// De som van de elementen op index i berekenen en opslaan
c[i] = a[i] + b[i];
}
}
int main() {
int *a, *b, *c;
int n = 1024;
// Geheugen op de host alloceren
a = (int*)malloc(n * sizeof(int));
b = (int*)malloc(n * sizeof(int));
c = (int*)malloc(n * sizeof(int));
// Inputvectoren initialiseren
for (int i = 0; i < n; i++) {
a[i] = i;
b[i] = i * 2;
}
// Geheugen op het apparaat alloceren
int *d_a, *d_b, *d_c;
cudaMalloc(&d_a, n * sizeof(int));
cudaMalloc(&d_b, n * sizeof(int));
cudaMalloc(&d_c, n * sizeof(int));
// Inputvectoren van host naar apparaat kopiëren
cudaMemcpy(d_a, a, n * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, b, n * sizeof(int), cudaMemcpyHostToDevice);
// De kernel starten
int blockSize = 256;
int numBlocks = (n + blockSize - 1) / blockSize;
vectorAdd<<<numBlocks,blockSize>>>(d_a, d_b, d_c, n);
// Resultaatvector van apparaat naar host kopiëren
cudaMemcpy(c, d_c, n * sizeof(int), cudaMemcpyDeviceToHost);
// Geheugen op het apparaat vrijgeven
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);
// Geheugen op de host vrijgeven
free(a);
free(b);
free(c);
return 0;
}
```Hier is de Nederlandse vertaling van de opgegeven markdown-bestand, waarbij de code-opmerkingen zijn vertaald:
n 0;
}Deze CUDA-code lanceert de vectorAdd-kernel met numBlocks-blokken en blockSize-threads per blok. De kernel voert een element-voor-element-optelling uit van de invoervectoren a en b en slaat het resultaat op in vector c. De <<<...>>>-syntax wordt gebruikt om de grid- en blokmaten op te geven bij het lanceren van een kernel.
CUDA-streams en -events
CUDA-streams en -events bieden een mechanisme voor gelijktijdige uitvoering en synchronisatie van kernels en geheugenoperaties:
- Streams: Een reeks bewerkingen (kernel-lanceringen, geheugenkopiëringen) die in volgorde worden uitgevoerd. Verschillende streams kunnen gelijktijdig worden uitgevoerd, waardoor de overlapping van berekening en geheugenoverdrachten mogelijk is.
- Events: Markeringen die in een stream kunnen worden ingevoegd om de voltooiing van specifieke bewerkingen op te nemen. Events kunnen worden gebruikt voor synchronisatie- en tijdmetingsdoeleinden.
Streams en events stellen programmeurs in staat om de prestaties van hun CUDA-toepassingen te optimaliseren door berekening en geheugenoverdrachtingen te overlappen en de volledige capaciteiten van de GPU-hardware te benutten.
Voorbeeld 3.2 toont het gebruik van CUDA-streams om de kerneluitvoering en geheugenoverdrachten te overlappen.
// Maak twee streams aan
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
// Kopieer invoergegevens asynchroon naar het apparaat
cudaMemcpyAsync(d_a, a, size, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_b, b, size, cudaMemcpyHostToDevice, stream2);
// Start kernels in verschillende streams
kernelA<<<blocks, threads, 0, stream1>>>(d_a);
kernelB<<<blocks, threads, 0, stream2>>>(d_b);
// Kopieer resultaten asynchroon terug naar de host
cudaMemcpyAsync(a, d_a, size, cudaMemcpyDeviceToHost, stream1);
cudaMemcpyAsync(b, d_b, size, cudaMemcpyDeviceToHost, stream2);
// Synchroniseer streams
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);In dit voorbeeld worden twee CUDA-streams gemaakt. Invoergegevens worden asynchroon naar het apparaat gekopieerd met behulp van elke stream. Vervolgens worden kernels in de verschillende streams gestart, waardoor de kerneluitvoering en geheugenoverdrachten kunnen worden overlappend.Hieronder vindt u de Nederlandse vertaling van de markdown-bestand, waarbij de code-opmerkingen zijn vertaald, maar de code zelf onvertaald is gebleven.
OpenCL-framework
OpenCL (Open Computing Language) is een open, royalty-vrije standaard voor parallelle programmering op heterogene platforms, waaronder CPU's, GPU's, FPGA's en andere acceleratoren. OpenCL biedt een uniforme programmeringsmodel en een set API's waarmee ontwikkelaars draagbare en efficiënte parallelle code kunnen schrijven.
OpenCL-programmeringsmodel
Het OpenCL-programmeringsmodel is vergelijkbaar met CUDA, met een paar belangrijke verschillen in terminologie en abstracties:
- Kernel: Een functie die parallel wordt uitgevoerd door een groot aantal work-items (threads) op een OpenCL-apparaat.
- Work-item: De basiseenheid van uitvoering in OpenCL, vergelijkbaar met een thread in CUDA.
- Work-group: Een verzameling work-items die kunnen synchroniseren en data kunnen delen via lokaal geheugen. Work-groups zijn vergelijkbaar met thread blocks in CUDA.
- NDRange: Definieert de indexruimte en work-item-organisatie voor een kernel-uitvoering. Het kan één-, twee- of driedimensionaal zijn.
OpenCL definieert ook een hiërarchisch geheugenmodel dat vergelijkbaar is met CUDA:
- Globaal geheugen: Toegankelijk voor alle work-items in alle work-groups, vergelijkbaar met globaal geheugen in CUDA.
- Lokaal geheugen: Gedeeld door alle work-items in een work-group, vergelijkbaar met gedeeld geheugen in CUDA.
- Privé geheugen: Privé voor een enkele work-item, vergelijkbaar met registers in CUDA.
- Constant geheugen: Alleen-lezen geheugen toegankelijk voor alle work-items.
OpenCL-kernels worden runtime-compilatie door de OpenCL-runtime. Het host-programma kan de beschikbare OpenCL-apparaten opvragen, een geschikt apparaat selecteren, een context maken en de kernel voor dat specifieke apparaat bouwen. Hierdoor kunnen OpenCL-toepassingen zeer draagbaar zijn op verschillende hardware-platforms.
Voorbeeld 3.3 toont een OpenCL-kernel die vectoroptelling uitvoert, vergelijkbaar met het CUDA-voorbeeld in Voorbeeld 3.1.
__kernel void vectorAdd(__global const int *a, __global const int *b, __global int *c) {
// Krijg de index van de huidige work-item
int index = get_global_id(0);
// Voer vectoroptelling uit
c[index] = a[index] + b[index];
}
```Here is the Dutch translation of the provided Markdown file, with the code comments translated:
```c
__kernel void bal(const int *a, const int *b, __global int *c, int n) {
int i = get_global_id(0);
if (i < n) {
c[i] = a[i] + b[i];
}
}De __kernel-sleutelwoord definieert een OpenCL-kernelfunctie. Het __global-sleutelwoord geeft aan dat een pointer naar het globale geheugen wijst. De get_global_id-functie retourneert de globale index van het huidige work-item, die gebruikt wordt om de geheugenaddressen voor de invoer- en uitvoervectoren te berekenen.
Algoritmen toewijzen aan GPU-architecturen
Efficiënt algoritmen toewijzen aan de GPU-architectuur is cruciaal voor het bereiken van hoge prestaties. Belangrijke overwegingen zijn:
-
Voldoende parallelisme 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.
-
Divergentie van vertakkingen minimaliseren: Afwijkende controlestromen binnen een warp/wavefront kunnen leiden tot seriële verwerking en verminderde SIMD-efficiëntie. Algoritmen moeten zo worden gestructureerd dat divergentie van vertakkingen zoveel mogelijk wordt vermeden.
-
Geheuërarchie benutten: Toegang tot het globale geheugen is duur. Algoritmen moeten het gebruik van gedeeld geheugen en registers maximaliseren om het aantal toegangen tot het globale geheugen te verminderen. Gegevens moeten ook zodanig in het geheugen worden gerangschikt dat gecoalesceerde geheugenaccess mogelijk is.
-
Balans vinden tussen berekeningen en geheugenaccesses: Algoritmen moeten een hoge verhouding hebben tussen rekenkundige bewerkingen en geheugenoperaties om de geheugenlatentie effectief te kunnen verbergen en een hoge rekenkracht te bereiken.
-
Host-apparaat datatransfers minimaliseren: Het overdragen van gegevens tussen het host- en apparaatgeheugen is traag. Algoritmen moeten dergelijke transfers minimaliseren door zoveel mogelijk berekeningen op de GPU uit te voeren.
Er worden verschillende parallelle algoritmepatronen veel gebruikt bij het ontwikkelen van GPU-kernels:
-
Map: Elke thread voert dezelfde bewerking uit op een ander gegevenspunt, waardoor eenvoudige parallelle verwerking van grote datasets mogelijk is.
-
Reduce: Parallelle reductie wordt gebruikt om efficiënt één waarde (bijv. som, maximum) te berekenen uit een grote invoerdataset.Threads voeren lokale reducties uit, die vervolgens worden gecombineerd om het uiteindelijke resultaat te produceren.
-
Scan: Ook bekend als prefix sum, wordt scan gebruikt om de lopende som van elementen in een array te berekenen. Efficiënte parallelle scan-algoritmen zijn belangrijke bouwstenen voor veel GPU-versnelde toepassingen.
-
Stencil: Elke thread berekent een waarde op basis van naburige gegevenselementen. Stencilberekeningen zijn gebruikelijk in wetenschappelijke simulaties en beeldverwerkingstoepassingen.
-
Gather/Scatter: Threads lezen van (verzamelen) of schrijven naar (verstrooien) willekeurige locaties in het globale geheugen. Zorgvuldige gegevensindeling en toegangspatronen zijn nodig voor efficiëntie.
Conclusie
GPU-programmeringsmodellen zoals CUDA en OpenCL onthullen de parallelle verwerkingsmogelijkheden van moderne GPU's aan ontwikkelaars, waardoor ze in staat zijn om een breed scala aan toepassingen te versnellen. Deze programmeringsmodellen bieden abstracties die het mogelijk maken om fijnmazige parallelle werkbelastingen efficiënt op de GPU-hardware af te beelden.
Het begrijpen van het uitvoeringsmodel, de geheugenhiërarchie en de synchronisatie primitieven die door deze programmeringsmodellen worden geboden, is essentieel voor het schrijven van high-performance GPU-code. Ontwikkelaars moeten zorgvuldig rekening houden met factoren zoals threadorganisatie, vertakking, geheugenadrespatronen en algoritme-ontwerp om de rekenkracht van GPU's volledig te benutten.
Naarmate GPU-architecturen blijven evolueren, moeten ook programmeringsmodellen en tools vooruitgaan om ontwikkelaars in staat te stellen nieuwe hardwarefuncties en -mogelijkheden effectief te benutten. Voortdurend onderzoek op gebieden als programmeertaalontwerp, compileroptimalisatie en autotuning zal cruciaal zijn voor het verbeteren van ontwikkelaarsproductiviteit en prestatiebaarheid in het tijdperk van heterogene computing.