Hoofdstuk 6: GPU-prestatiemaatstaven en -analyse

Het analyseren en optimaliseren van de prestaties van GPU-toepassingen is cruciaal voor het bereiken van een hoge efficiëntie en benutting van GPU-hardwarebronnen. In dit hoofdstuk zullen we de belangrijkste GPU-prestatiemaatstaven, profilering- en optimalisatiehulpmiddelen, technieken voor het identificeren van prestatieknelpunten en strategieën voor het verbeteren van GPU-prestaties verkennen.

Doorvoer, latentie en geheugenbandbreedte

Drie fundamentele maatstaven voor het evalueren van GPU-prestaties zijn doorvoer, latentie en geheugenbandbreedte. Het begrijpen van deze maatstaven en hun implicaties is essentieel voor het analyseren en optimaliseren van GPU-toepassingen.

Doorvoer

Doorvoer verwijst naar het aantal bewerkingen of taken dat een GPU in een bepaalde tijd kan voltooien. Het wordt meestal gemeten in drijvende-kommaberekeningen per seconde (FLOPS) of instructies per seconde (IPS). GPU's zijn ontworpen om een hoge doorvoer te bereiken door parallellisme te benutten en een groot aantal threads gelijktijdig uit te voeren.

De theoretische piekdoorvoer van een GPU kan worden berekend met de volgende formule:

Piekdoorvoer (FLOPS) = Aantal CUDA-cores × Klokfrequentie × FLOPS per CUDA-core per cyclus

Een NVIDIA GeForce RTX 2080 Ti GPU heeft bijvoorbeeld 4352 CUDA-cores, een basisklokfrequentie van 1350 MHz en elke CUDA-core kan 2 drijvende-kommaberekeningen per cyclus uitvoeren (FMA - Fused Multiply-Add). Daarom is zijn theoretische piekdoorvoer:

Piekdoorvoer (FLOPS) = 4352 × 1350 MHz × 2 = 11,75 TFLOPS

Het bereiken van de theoretische piekdoorvoer in de praktijk is echter een uitdaging vanwege verschillende factoren zoals geheugentoegangpatronen, vertakkingsdivergentie en bronbeperkingen.

Latentie

Latentie verwijst naar de tijd die het kost voor een enkele bewerking of taak om te voltooien. In de context van GPU's is latentie vaak geassocieerd met geheugentoegangsbewerkingen. GPU's hebben een hiërarchisch geheugensysteem en het toegang krijgen tot gegevens op verschillende niveaus van de geheugenhiTypische latentietijden voor verschillende geheugenlagen in een GPU zijn:

Registers: 0-1 cycli
Gedeeld geheugen: 1-2 cycli
L1-cache: 20-30 cycli
L2-cache: 200-300 cycli
Globaal geheugen (DRAM): 400-800 cycli

Latentie kan een aanzienlijke invloed hebben op de GPU-prestaties, vooral wanneer er afhankelijkheden zijn tussen bewerkingen of wanneer threads wachten op het ophalen van gegevens uit het geheugen. Technieken zoals latentie-hiding, prefetching en caching kunnen helpen om de impact van latentie op de GPU-prestaties te verminderen.

Geheugen bandbreedte

Geheugen bandbreedte verwijst naar de snelheid waarmee gegevens kunnen worden overgedragen tussen de GPU en zijn geheugen subsysteem. Het wordt meestal gemeten in bytes per seconde (B/s) of gigabytes per seconde (GB/s). GPU's hebben high-bandwidth geheugen interfaces, zoals GDDR6 of HBM2, om de data-intensieve aard van graphics- en rekenkundige werkbelastingen te ondersteunen.

De theoretische piek geheugen bandbreedte van een GPU kan worden berekend met de volgende formule:

Piek geheugen bandbreedte (GB/s) = Geheugen klokfrequentie × Geheugen bus breedte ÷ 8

Een NVIDIA GeForce RTX 2080 Ti GPU heeft bijvoorbeeld een geheugen klokfrequentie van 7000 MHz (effectief) en een geheugen bus breedte van 352 bits. Daarom is zijn theoretische piek geheugen bandbreedte:

Piek geheugen bandbreedte (GB/s) = 7000 MHz × 352 bits ÷ 8 = 616 GB/s

Geheugen bandbreedte is een kritieke factor in GPU-prestaties, aangezien veel GPU-toepassingen geheugen-gebonden zijn, wat betekent dat hun prestaties worden beperkt door de snelheid waarmee gegevens kunnen worden overgedragen tussen de GPU en het geheugen. Het optimaliseren van geheugen toegangspatronen, het minimaliseren van gegevensoverdrachten en het benutten van de geheugen hiërarchie kunnen helpen om de benutting van de geheugen bandbreedte te verbeteren.

Profilering en prestatie-optimalisatie tools

Profilering en prestatie-optimalisatie tools zijn essentieel voor het analyseren van het gedrag van GPU-toepassingen, het identificeren van prestatiebeperkingen en het begeleiden van optimalisatie-inspanningen. Deze tools bieden inzicht in verschillende aspecten van GPU-prestaties, zoals de uitvoeringstijd van kernels, geheugen toegang patronen en resourcegebruik.Hier is de Nederlandse vertaling van het bestand:

ES-patronen, bezetting en hulpbronnengebruik

Enkele populaire profilering- en prestatie-optimalisatiehulpmiddelen voor GPU's zijn:

NVIDIA Visual Profiler (nvvp): Een grafisch profileringshulpmiddel dat een uitgebreid overzicht biedt van de prestaties van GPU-toepassingen. Het stelt ontwikkelaars in staat om de uitvoering van kernels, geheugenoverdrachten en API-oproepen te analyseren en biedt aanbevelingen voor optimalisatie.
NVIDIA Nsight: Een geïntegreerde ontwikkelomgeving (IDE) die profilering- en foutopsporingscapaciteiten voor GPU-toepassingen bevat. Het ondersteunt verschillende programmeertalen en frameworks, zoals CUDA, OpenCL en OpenACC.
NVIDIA Nsight Compute: Een zelfstandig profileringshulpmiddel dat zich richt op de analyse van GPU-kernelprestaties. Het biedt gedetailleerde prestatiegegevens, zoals instructiedoorvoer, geheugenefficiëntie en bezetting, en helpt bij het identificeren van prestatieknelpunten op broncodeniveau.
AMD Radeon GPU Profiler (RGP): Een profileringshulpmiddel voor AMD-GPU's dat prestatiegegevens vastlegt en visualiseert voor DirectX-, Vulkan- en OpenCL-toepassingen. Het biedt inzicht in GPU-bezetting, geheugengebruik en pipeline-vertragingen.
AMD Radeon GPU Analyzer (RGA): Een statisch analysehulpmiddel dat GPU-shadercode analyseert en prestatievoorspellingen, hulpbronnengebruik en optimalisatievoorstellen biedt.

Deze hulpmiddelen werken meestal door de GPU-toepassingscode te instrumenteren, prestatiegegevens tijdens de uitvoering te verzamelen en de gegevens in een gebruiksvriendelijk formaat weer te geven voor analyse. Ze bieden vaak tijdlijnweergaven, prestatietellers en broncodecorrelatie om ontwikkelaars te helpen bij het identificeren van prestatieknelpunten en het optimaliseren van hun code.

Voorbeeld: Een CUDA-toepassing profileren met NVIDIA Visual Profiler (nvvp)

Bouw de CUDA-toepassing met profilering ingeschakeld:
```
nvcc -o myapp myapp.cu -lineinfo
```
Voer de toepassing uit met profilering:
```
nvprof ./myapp
```
Open de Visual Profiler:
```
nvvp
```
Importeer de gegenereerde profileringgegevensHier is de Nederlandse vertaling van het bestand, waarbij de code-opmerkingen zijn vertaald:
Analyseer de tijdlijnweergave, kernelprestaties, geheugenoverdrachten en API-aanroepen.
Identificeer prestatieknelpunten en optimaliseer de code op basis van de aanbevelingen van de profiler.

Prestatieknelpunten identificeren

Het identificeren van prestatieknelpunten is cruciaal voor het optimaliseren van GPU-toepassingen. Prestatieknelpunten kunnen ontstaan door verschillende factoren, zoals inefficiënte geheugentoegangpatronen, lage bezetting, vertakkingsdivergentie en bronbeperkingen. Enkele veel gebruikte technieken voor het identificeren van prestatieknelpunten zijn:

Profileren: Het gebruik van profileerhulpmiddelen om de uitvoeringstijd van kernels, geheugenoverdrachtentijd en API-overhead te meten, kan helpen om te identificeren welke delen van de toepassing de meeste tijd en middelen verbruiken.
Bezetting analyseren: Bezetting verwijst naar de verhouding tussen actieve warps en het maximale aantal warps dat door een GPU wordt ondersteund. Een lage bezetting kan duiden op een onderbenutting van GPU-middelen en kan suggereren dat er optimalisaties nodig zijn voor blok- en rasterafmetingen of het verminderen van register- en gedeeld geheugengebruik.
Geheugentoegangpatronen onderzoeken: Inefficiënte geheugentoegangpatronen, zoals niet-gecoördineerde geheugentoegangen of frequente toegangen tot het globale geheugen, kunnen de GPU-prestaties aanzienlijk beïnvloeden. Het analyseren van geheugentoegangpatronen met behulp van profileerhulpmiddelen kan helpen bij het identificeren van optimalisatiemogelijkheden, zoals het gebruik van gedeeld geheugen of het verbeteren van gegevenslocatie.
Vertakkingsdivergentie onderzoeken: Vertakkingsdivergentie treedt op wanneer threads binnen een warp verschillende uitvoerpaden volgen als gevolg van voorwaardelijke instructies. Divergente vertakkingen kunnen leiden tot serialisatie en verminderde prestaties. Het identificeren en minimaliseren van vertakkingsdivergentie kan helpen de GPU-prestaties te verbeteren.
Resourcegebruik bewaken: GPU's hebben beperkte middelen, zoals registers, gedeeld geheugen en thread-blokken. Het bewaken van het resourcegebruik met behulp van profileerhulpmiddelen kan helpen bij het identificeren van resourceknelpunten en het sturen van optimalisatie-inspanningen, zoals het verminderen van registergebruik.Hier is de Nederlandse vertaling van het Markdown-bestand, waarbij de code-opmerkingen zijn vertaald:

Voorbeeld: Het identificeren van een geheugentoegangsfles

Profileer de CUDA-toepassing met behulp van Nsight Compute:
```
ncu -o profile.ncu-rep ./myapp
```
Open het gegenereerde profielrapport in Nsight Compute.
Analyseer de sectie "Memory Workload Analysis" om inefficiënte geheugentoegangpatronen te identificeren, zoals niet-gecoördineerde toegangen of hoog gebruik van globaal geheugen.
Optimaliseer de geheugentoegangpatronen op basis van de inzichten die Nsight Compute biedt, zoals het gebruik van gedeeld geheugen of het verbeteren van de gegevenslocatie.

Strategieën voor het verbeteren van GPU-prestaties

Zodra de prestatiebottlenecks zijn geïdentificeerd, kunnen verschillende strategieën worden toegepast om de GPU-prestaties te verbeteren. Enkele veel voorkomende optimalisatiestrategieën zijn:

Maximaliseren van parallellisme: Zorg ervoor dat de toepassing wordt opgedeeld in voldoende parallelle taken om de GPU-resources volledig te benutten. Dit kan inhouden dat de blok- en rasterafmetingen worden aangepast, dat streams worden gebruikt voor gelijktijdige uitvoering of dat taak-niveau parallellisme wordt benut.
Optimaliseren van geheugentoegangpatronen: Verbeter de efficiëntie van geheugentoegangen door het minimaliseren van globale geheugentoegangen, het gebruik van gedeeld geheugen voor vaak geraadpleegde gegevens en het garanderen van gecoördineerde geheugentoegangen. Technieken zoals geheugentiling, gegevenslay-outtransformaties en caching kunnen helpen bij het optimaliseren van geheugenprestaties.
Verminderen van takdivergentie: Minimaliseer takdivergentie door de code te herstructureren om divergente takken binnen een warp te vermijden. Technieken zoals taktoewijzing, gegevensafhankelijke vertakking en warp-level programmering kunnen helpen om de impact van takdivergentie te verminderen.
Benutten van geheugenhi??rarchie: Benut de GPU-geheugenhi??rarchie effectief door het maximaliseren van het gebruik van registers en gedeeld geheugen voor vaak geraadpleegde gegevens. Gebruik tekstuurgeheugen en constant geheugen voor alleen-lezen gegevens die ruimtelijke lokaliteit vertonen of uniform over de threads worden benaderd.
**Overlappen van berekening en geheugenHier is de Nederlandse vertaling van het Markdown-bestand, waarbij de code-opmerkingen zijn vertaald:

GPU-optimalisatie-tips

Geheugencoalescing: Zorg ervoor dat de geheugentoegangen van threads in een warp gecoalesced zijn om de bandbreedte van het geheugen te maximaliseren.
Gebruik van gedeeld geheugen: Gebruik gedeeld geheugen om de toegang tot het globale geheugen te minimaliseren en de geheugenlatentie te verminderen.
Vermijd divergentie van threads: Vermijd divergentie van threads binnen een warp, omdat dit de prestaties kan verminderen.
Minimaliseer synchronisatie: Minimaliseer het gebruik van synchronisatie tussen threads, omdat dit de prestaties kan verminderen.
Geheugeroverdracht verbergen: Verberg de latentie van geheugeroverdracht door berekeningen te overlappen met geheugeroverdrachten met behulp van CUDA-streams of OpenCL-commando-wachtrijen. Hierdoor kan de GPU berekeningen uitvoeren terwijl gegevens worden overgedragen tussen het host- en apparaatgeheugen.
Afstemmen van kernelstartparameters: Experimenteer met verschillende blok- en rastergroottes om de optimale configuratie voor elke kernel te vinden. De optimale startparameters hangen af van factoren zoals het aantal registers per thread, het gebruik van gedeeld geheugen en de kenmerken van de GPU-architectuur.
Minimaliseren van host-apparaat-gegevensoverdrachten: Verminder de hoeveelheid gegevens die worden overgedragen tussen de host (CPU) en het apparaat (GPU) door zoveel mogelijk berekeningen op de GPU uit te voeren. Groepeer kleine overdrachten in grotere overdrachten om de overhead van elke overdracht te amortiseren.
Gebruik van asynchrone bewerkingen: Maak gebruik van asynchrone bewerkingen, zoals asynchrone geheugenkopiëren en kernelstarts, om berekeningen en communicatie te overlappen. Hierdoor kan de CPU andere taken uitvoeren terwijl de GPU wordt uitgevoerd, waardoor de algehele toepassingsprestaties worden verbeterd.

Voorbeeld: Optimaliseren van geheugeraccesspatronen met behulp van gedeeld geheugen in CUDA

Originele code met inefficiënte globale geheugeraccessen:

__global__ void myKernel(float* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        float result = 0.0f;
        for (int i = 0; i < n; i++) {
            result += data[tid] * data[i];
        }
        data[tid] = result;
    }
}

Geoptimaliseerde code met gebruik van gedeeld geheugen:

__global__ void myKernel(float* data, int n) {
    __shared__ float sharedData[256];
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int localIdx = threadIdx.x;
 
    if (tid < n) {
        sharedData[localIdx] = data[tid];
    }
    __syncthreads();
 
    if (tid < n) {
        float result = 0.0f;
        for (int i = 0; i < blockDim.x; i++) {
            result += sharedData[localIdx] * sharedData[i];
        }
        data[tid] = result;
    }
}
```Hier is de Nederlandse vertaling van het bestand, waarbij de code-opmerkingen zijn vertaald:
 
a[tid] = result;
    }
}

In de geoptimaliseerde code wordt de invoergegevens eerst geladen in het gedeelde geheugen, dat veel lagere latentie heeft in vergelijking met het globale geheugen. De berekening wordt vervolgens uitgevoerd met behulp van het gedeelde geheugen, waardoor het aantal toegangen tot het globale geheugen wordt verminderd en de prestaties worden verbeterd.

Conclusie

Het analyseren en optimaliseren van GPU-prestaties is essentieel voor het ontwikkelen van efficiënte en hoogwaardige GPU-toepassingen. Door inzicht te krijgen in belangrijke prestatiemaatstaven zoals doorvoer, latentie en geheugenbandbreedte, kunnen ontwikkelaars geïnformeerde beslissingen nemen over het optimaliseren van hun code.

Profilering- en prestatie-optimalisatiehulpmiddelen spelen een cruciale rol bij het identificeren van prestatiebeperkingen en het begeleiden van optimalisatie-inspanningen. Deze tools bieden waardevolle inzichten in de uitvoering van kernels, geheugentoegangpatronen, bezetting en hulpbronnengebruik, waardoor ontwikkelaars hun optimalisatie-inspanningen kunnen richten op de meest kritieke gebieden.

Veel voorkomende optimalisatiestrategieën zijn het maximaliseren van parallellisme, het optimaliseren van geheugentoegangpatronen, het verminderen van takdivergentie, enz.

Hier zijn enkele veel voorkomende strategieën voor het optimaliseren van GPU-prestaties, vervolgd in Markdown-formaat:

Verminderen van takdivergentie: Divergente controlestromen binnen een warp/wavefront kunnen leiden tot serialisatie en verminderde SIMD-efficiëntie. Algoritmen moeten zodanig worden gestructureerd dat takdivergentie zoveel mogelijk wordt vermeden. Technieken zoals taktoewijzing, datagebaseerde vertakking en warp-level programmering kunnen helpen bij het verminderen van de impact van takdivergentie.
Gebruik maken van de geheugenhi??rarchie: Maak effectief gebruik van de GPU-geheugenhi??rarchie door het gebruik van registers en gedeeld geheugen voor vaak geraadpleegde gegevens te maximaliseren. Gebruik tekstuurgeheugen en constant geheugen voor alleen-lezen gegevens die ruimtelijke lokaliteit vertonen of uniform worden benaderd door threads.
Overlappen van berekeningen en geheugenovertochten: Verberg de latentie van geheugenovertochten door berekeningen te overlappen met geheugenovertochten met behulp van CUDA-streams of OpenCL-commando-wachtrijen. Hierdoor kunnen GPU-kernen blijven werken terwijl geheugenovertochten worden uitgevoerd.Hier is de Nederlandse vertaling van het Markdown-bestand, waarbij de code-opmerkingen zijn vertaald:
Afstemmen van Kernel Launch Parameters: Experimenteer met verschillende block- en grid-groottes om de optimale configuratie voor elke kernel te vinden. De optimale launch-parameters hangen af van factoren zoals het aantal registers gebruikt per thread, het gebruik van gedeeld geheugen en de kenmerken van de GPU-architectuur.
Minimaliseren van Host-Device Gegevensoverdrachten: Verminder de hoeveelheid gegevens die worden overgedragen tussen de host (CPU) en het apparaat (GPU) door zoveel mogelijk berekeningen op de GPU uit te voeren. Groepeer kleine overdrachten in grotere overdrachten om de overhead van elke overdracht te spreiden.
Gebruik maken van Asynchrone Bewerkingen: Maak gebruik van asynchrone bewerkingen, zoals asynchrone geheugenkopiëren en kernel-starts, om berekeningen en communicatie te overlappen. Dit stelt de CPU in staat om andere taken uit te voeren terwijl de GPU wordt uitgevoerd, waardoor de algehele applicatieprestaties worden verbeterd.

Voorbeeld: Optimaliseren van geheugenadrespatronen met behulp van gedeeld geheugen in CUDA

Originele code met inefficiënte globale geheugentoegangen:

__global__ void myKernel(float* data, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        float result = 0.0f;
        for (int i = 0; i < n; i++) {
            result += data[tid] * data[i];
        }
        data[tid] = result;
    }
}

Geoptimaliseerde code met gebruik van gedeeld geheugen:

__global__ void myKernel(float* data, int n) {
    __shared__ float sharedData[256];
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    int localIdx = threadIdx.x;
 
    if (tid < n) {
        sharedData[localIdx] = data[tid];
    }
    __syncthreads();
 
    if (tid < n) {
        float result = 0.0f;
        for (int i = 0; i < blockDim.x; i++) {
            result += sharedData[localIdx] * sharedData[i];
        }
        data[tid] = result;
    }
}

In de geoptimaliseerde code wordt de invoergegevens eerst geladen in het gedeelde geheugen, dat een veel lagere latentie heeft in vergelijking metHere is the Dutch translation for the provided markdown file, with the code comments translated:

Globaal geheugen

De berekening wordt vervolgens uitgevoerd met behulp van het gedeelde geheugen, waardoor het aantal toegangen tot het globale geheugen wordt verminderd en de prestaties worden verbeterd.

// Alloceer geheugen op het apparaat
cudaMalloc(&d_a, size);
cudaMalloc(&d_b, size);
cudaMalloc(&d_c, size);
 
// Kopieer gegevens van host naar apparaat
cudaMemcpy(d_a, h_a, size, cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b, size, cudaMemcpyHostToDevice);
 
// Voer de kernel uit
kernel<<<gridSize, blockSize>>>(d_a, d_b, d_c);
 
// Kopieer resultaten van apparaat naar host
cudaMemcpy(h_c, d_c, size, cudaMemcpyDeviceToHost);
 
// Maak geheugen op het apparaat vrij
cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);

Chapter 5 Gpu Memory System Design Chapter 7 Streaming Multiprocessor Design