Dit is de vierde post in de CUDA Refresher-serie, die tot doel heeft de belangrijkste concepten in CUDA, tools, en optimalisatie voor beginnende of halfgevorderde ontwikkelaars.
Het CUDA-programmeermodel biedt een abstractie van GPU-architectuur die fungeert als een brug tussen een applicatie en de mogelijke implementatie ervan op GPU-hardware. Dit bericht schetst de belangrijkste concepten van het CUDA-programmeermodel door te schetsen hoe ze worden weergegeven in algemene programmeertalen zoals C/C++.
Laat me twee sleutelwoorden introduceren die veel worden gebruikt in het CUDA-programmeermodel: host en apparaat .
De host is de CPU die beschikbaar is in het systeem. Het systeemgeheugen dat bij de CPU hoort, wordt hostgeheugen genoemd. De GPU wordt een apparaat genoemd en het GPU-geheugen wordt ook wel apparaatgeheugen genoemd.
Om een CUDA-programma uit te voeren, zijn er drie hoofdstappen:
- Kopieer de invoergegevens van het hostgeheugen naar apparaatgeheugen, ook wel host-to-device transfer genoemd.
- Laad het GPU-programma en voer het uit, waarbij gegevens op de chip worden opgeslagen voor prestaties.
- Kopieer de resultaten van het apparaatgeheugen naar het hostgeheugen, ook wel apparaat-naar-host-overdracht genoemd.
- CUDA-architectuur beperkt het aantal threads per blok (limiet van 1024 threads per blok).
- De dimensie van het threadblok is toegankelijk binnen de kernel via de ingebouwde blockDim-variabele.
- Alle threads binnen een blok kunnen worden gesynchroniseerd met behulp van een intrinsieke functie __syncthreads. Met __syncthreads moeten alle threads in het blok wachten voordat iemand verder kan gaan.
- Het aantal threads per blok en het aantal blokken per raster gespecificeerd in de <<<…>>> syntaxis kan van het type int of dim3 zijn. Deze drievoudige punthaken markeren een oproep van hostcode naar apparaatcode. Het wordt ook wel een kernellancering genoemd.
- Registers —Deze zijn privé voor elke thread, wat betekent dat registers die aan een thread zijn toegewezen, niet zichtbaar zijn voor andere threads. De compiler neemt beslissingen over het registergebruik.
- L1/Gedeeld geheugen (SMEM)—Elke SM heeft een snel, on-chip kladblokgeheugen dat kan worden gebruikt als L1-cache en gedeeld geheugen. Alle threads in een CUDA-blok kunnen gedeeld geheugen delen, en alle CUDA-blokken die op een bepaalde SM draaien, kunnen de fysieke geheugenbron delen die door de SM wordt geleverd..
- Alleen-lezen geheugen —Elke SM heeft een instructiecache, constant geheugen, textuurgeheugen en RO-cache, die alleen-lezen is voor kernelcode.
- L2-cache—De L2-cache wordt gedeeld door alle SM's, dus elke thread in elk CUDA-blok heeft toegang tot dit geheugen. De NVIDIA A100 GPU heeft de L2-cachegrootte vergroot tot 40 MB in vergelijking met 6 MB in V100 GPU's.
- Globaal geheugen—Dit is de framebuffergrootte van de GPU en het DRAM zittend in de GPU.
ul>
CUDA-kernel en thread-hiërarchie
Afbeelding 1 laat zien dat de CUDA-kernel een functie is die wordt uitgevoerd op GPU. Het parallelle gedeelte van uw toepassingen wordt K keer parallel uitgevoerd door K verschillende CUDA-threads, in tegenstelling tot slechts één keer zoals gewone C/C++-functies.
Figuur 1. De kernel is een functie die wordt uitgevoerd op de GPU.
Elke CUDA-kernel begint met een __global__ declaratiespecificatie. Programmeurs geven elke thread een unieke globale ID door gebruik te maken van ingebouwde variabelen.
Figuur 2. CUDA-kernels zijn onderverdeeld in blokken.
Een groep threads wordt een CUDA-blok genoemd. CUDA-blokken zijn gegroepeerd in een raster. Een kernel wordt uitgevoerd als een raster van blokken threads (Figuur 2).
Elk CUDA-blok wordt uitgevoerd door één streaming multiprocessor (SM) en kan niet worden gemigreerd naar andere SM's in GPU (behalve tijdens preëmptie, foutopsporing of dynamisch CUDA-parallellisme). Eén SM kan meerdere gelijktijdige CUDA-blokken uitvoeren, afhankelijk van de bronnen die nodig zijn voor CUDA-blokken. Elke kernel wordt uitgevoerd op één apparaat en CUDA ondersteunt het tegelijkertijd uitvoeren van meerdere kernels op een apparaat. Afbeelding 3 toont de uitvoering en toewijzing van de kernel op hardwarebronnen die beschikbaar zijn in GPU.
Figuur 3. Kerneluitvoering op GPU.
CUDA definieert ingebouwde 3D-variabelen voor threads en blokken. Threads worden geïndexeerd met behulp van de ingebouwde 3D-variabele threadIdx. Driedimensionale indexering biedt een natuurlijke manier om elementen in vectoren, matrix en volume te indexeren en maakt CUDA-programmering eenvoudiger. Op dezelfde manier worden blokken ook geïndexeerd met behulp van de ingebouwde 3D-variabele genaamd blockIdx.
Hier zijn een paar opvallende punten:
Het CUDA-programma voor het toevoegen van twee matrices hieronder toont multidimensionale blockIdx en threadIdx en andere variabelen zoals blockDim. In het onderstaande voorbeeld is een 2D-blok gekozen om het indexeren te vergemakkelijken en heeft elk blok 256 threads met elk 16 in x- en y-richting. Het totale aantal blokken wordt berekend met behulp van de gegevensgrootte gedeeld door de grootte van elk blok.
//Kernel – Twee matrices toevoegen MatA en MatB __global__ void MatAdd(float MatA[N][N], float MatB[N][N], float MatC[N][N]) { int i = blockIdx.x * blockDim .x + threadIdx.x; int j = blockIdx.y * blockDim.y + threadIdx.y; if (i < N && j < N) MatC[i][j] = MatA[i][j] + MatB[i][j]; } int main() { … //Matrix-toevoeging kernel starten vanaf hostcode dim3 threadsPerBlock(16, 16); dim3 numBlocks((N + threadsPerBlock.x -1)/threadsPerBlock.x, (N+threadsPerBlock.y -1)/threadsPerBlock.y); MatAdd<<<numBlocks, threadsPerBlock>>>(MatA, MatB, MatC); … }
Geheugenhiërarchie
CUDA-compatibele GPU's hebben een geheugenhiërarchie zoals weergegeven in Afbeelding 4.
Figuur 4. Geheugenhiërarchie in GPU's.
De volgende herinneringen worden zichtbaar door de GPU-architectuur:
De NVIDIA CUDA-compiler doet goed werk bij het optimaliseren van geheugenbronnen, maar een ervaren CUDA-ontwikkelaar kan ervoor kiezen om deze geheugenhiërarchie efficiënt te gebruiken om de CUDA-programma's naar behoefte te optimaliseren.
Rekencapaciteit
De rekencapaciteit van een GPU bepaalt de algemene specificaties en beschikbare functies die worden ondersteund door de GPU-hardware. Dit versienummer kan tijdens runtime door applicaties worden gebruikt om te bepalen welke hardwarefuncties of instructies beschikbaar zijn op de huidige GPU.
Elke GPU wordt geleverd met een versienummer dat wordt aangeduid als X.Y, waarbij X een belangrijk revisienummer is en Y een klein revisienummer. Het kleine revisienummer komt overeen met een incrementele verbetering van de architectuur, mogelijk inclusief nieuwe functies.
Zie de CUDA-voorbeeldcode deviceQuery voor meer informatie over de rekencapaciteit van elk CUDA-apparaat. Dit voorbeeld somt de eigenschappen op van de CUDA-apparaten die in het systeem aanwezig zijn
Samenvatting
Het CUDA-programmeermodel biedt een heterogene omgeving waarin de hostcode het C/C++-programma op de CPU uitvoert en de kernel draait op een fysiek gescheiden GPU-apparaat. Het CUDA-programmeermodel gaat er ook van uit dat zowel de host als het apparaat hun eigen afzonderlijke geheugenruimten behouden, respectievelijk hostgeheugen en apparaatgeheugen genoemd. CUDA-code zorgt ook voor gegevensoverdracht tussen host- en apparaatgeheugen, via de PCIe-bus.
CUDA legt ook veel ingebouwde variabelen bloot en biedt de flexibiliteit van multidimensionale indexering om het programmeren te vergemakkelijken. CUDA beheert ook verschillende geheugens, waaronder registers, gedeeld geheugen en L1-cache, L2-cache en globaal geheugen. Geavanceerde ontwikkelaars kunnen sommige van deze geheugens efficiënt gebruiken om het CUDA-programma te optimaliseren.
Pradeep Gupta
Directeur, Solutions Architecture and Engineering Team, NVIDIA
