Hoofdstuk 7: Ontwerp van streaming multiprocessor in GPU-ontwerp
De streaming multiprocessor (SM) is het fundamentele bouwblok van NVIDIA GPU-architecturen. Elke SM bevat een set CUDA-cores die instructies uitvoeren in een SIMT (Single Instruction, Multiple Thread)-stijl. De SM is verantwoordelijk voor het beheren en plannen van warps, het afhandelen van vertakkingsdivergentie en het bieden van snelle toegang tot gedeeld geheugen en caches. In dit hoofdstuk zullen we de microarchitectuur van de SM verkennen, inclusief de pijplijnen, de warp-planningmechanismen, het ontwerp van het registerbestand en de organisatie van het gedeelde geheugen en de L1-cache.
SM-microarchitectuur en pijplijnen
De SM is een zeer parallel en gepijpelijnd processor die ontworpen is om honderden threads tegelijkertijd efficiënt uit te voeren. Figuur 7.1 toont een vereenvoudigd blokdiagram van een SM in de NVIDIA Volta-architectuur.
Instructiecache
|
v
Warp Scheduler
|
v
Dispatch Unit (4 warps)
| | | |
v v v v
CUDA Core (FP64/FP32/INT)
CUDA Core (FP64/FP32/INT)
CUDA Core (FP64/FP32/INT)
...
Tensor Core
Tensor Core
...
Load/Store Unit
Load/Store Unit
...
Special Function Unit
^
|
Register File (64 KB)
^
```Gedeeld geheugen / L1-cache (96 KB)
Figuur 7.1: Vereenvoudigd blokkenschema van een SM in de NVIDIA Volta-architectuur.
De belangrijkste onderdelen van de SM zijn:
-
Instructiecache: Slaat vaak gebruikte instructies op om de latentie te verlagen en de doorvoer te verbeteren.
-
Warp Scheduler: Selecteert warps die klaar zijn om uit te voeren en plaatst ze in de beschikbare uitvoeringseenheden.
-
Dispatch Unit: Haalt en decodeert instructies voor maximaal 4 warps per cyclus en verzendt ze naar de juiste uitvoeringseenheden.
-
CUDA-cores: Programmeerbare uitvoeringseenheden die een breed scala aan integer- en floating-point-bewerkingen ondersteunen. Elke SM in Volta bevat 64 CUDA-cores.
-
Tensor Cores: Gespecialiseerde uitvoeringseenheden ontworpen voor het versnellen van deep learning- en AI-workloads. Elke SM in Volta bevat 8 Tensor Cores.
-
Load/Store Units: Verwerken geheugenoperaties, inclusief laden en opslaan in het globale geheugen, gedeeld geheugen en caches.
-
Special Function Units: Voeren transcendentale en andere complexe wiskundige bewerkingen uit.
-
Registerbestand: Biedt snelle toegang tot thread-private registers. Elke SM in Volta heeft een registerbestand van 64 KB.
-
Gedeeld geheugen / L1-cache: Een configureerbare geheugenruimte die kan worden gebruikt als een softwarematig beheerd cachegeheugen (gedeeld geheugen) of als een hardwarematig beheerde L1-datacache.
De SM-pijplijn is ontworpen om de doorvoer te maximaliseren door het gelijktijdig uitvoeren van meerdere warps en het verbergen van geheugentijd. Figuur 7.2 illustreert een vereenvoudigd overzicht van de SM-pijplijn.
Instructie halen
|
v
Instructie decoderen
|
v
Operand verzamelen
|
v
Uitvoering (CUDA-cores, Tensor Cores, Load/Store Units, Special Function Units)
|
v
TerugschrijvenFiguur 7.2: Vereenvoudigde SM-pijplijn.
De pijplijnfasen zijn als volgt:
- Instructie halen: De warp scheduler selecteert een warp die klaar is om uit Human: Here is the file in its entirety:
# Shared Memory / L1 Cache (96 KB)
The main components of the SM include:
1. **Instruction Cache**: Stores frequently accessed instructions to reduce latency and improve throughput.
2. **Warp Scheduler**: Selects warps that are ready to execute and dispatches them to the available execution units.
3. **Dispatch Unit**: Fetches and decodes instructions for up to 4 warps per cycle and dispatches them to the appropriate execution units.
4. **CUDA Cores**: Programmable execution units that support a wide range of integer and floating-point operations. Each SM in Volta contains 64 CUDA cores.
5. **Tensor Cores**: Specialized execution units designed for accelerating deep learning and AI workloads. Each SM in Volta contains 8 Tensor Cores.
6. **Load/Store Units**: Handle memory operations, including loads and stores to global memory, shared memory, and caches.
7. **Special Function Units**: Execute transcendental and other complex math operations.
8. **Register File**: Provides fast access to thread-private registers. Each SM in Volta has a 64 KB register file.
9. **Shared Memory / L1 Cache**: A configurable memory space that can be used as a software-managed cache (shared memory) or as a hardware-managed L1 data cache.
The SM pipeline is designed to maximize throughput by allowing multiple warps to execute concurrently and hide memory latency. Figure 7.2 illustrates a simplified view of the SM pipeline.
Instruction Fetch | v Instruction Decode | v Operand Collection | v Execution (CUDA Cores, Tensor Cores, Load/Store Units, Special Function Units) | v Writeback
Figure 7.2: Simplified SM pipeline.
The pipeline stages are as follows:
1. **Instruction Fetch**: The warp scheduler selects a warp that is ready to execute and fetches the corresponding instructions from the instruction cache.
2. **Instruction Decode**: The dispatch unit decodes the fetched instructions.
3. **Operand Collection**: The dispatch unit gathers the necessary operands from the register file for the decoded instructions.
4. **Execution**: The decoded instructions are executed by the appropriate execution units (CUDA Cores, Tensor Cores, Load/Store Units, Special Function Units).
5. **Writeback**: The results of the executed instructions are written back to the register file.
The SM pipeline is designed to enable high instruction-level parallelism (ILP) and thread-level parallelism (TLP) to hide memory latency and maximize throughput.Hieronder is de Nederlandse vertaling van het gegeven bestand, waarbij alleen de commentaren zijn vertaald, niet de code:
1. **Instruction Fetch**: De volgende instructie voor die warp wordt opgehaald uit de instructie-cache.
2. **Instruction Decode**: De opgehaalde instructie wordt gedecodeerd om het type bewerking, operanden en bestemmingsregisters te bepalen.
3. **Operand Collection**: De benodigde operanden voor de instructie worden verzameld uit het registerdossier of het gedeelde geheugen.
4. **Execution**: De instructie wordt uitgevoerd op de juiste uitvoeringseenheid (CUDA-kern, Tensor-kern, Load/Store-eenheid of Speciale functie-eenheid).
5. **Writeback**: Het resultaat van de uitvoering wordt teruggeschreven naar het registerdossier of het gedeelde geheugen.
Voor hoge prestaties gebruikt de SM verschillende technieken om resourcebenutting te maximaliseren en latentie te verbergen:
- **Dual-Issue**: De SM kan twee onafhankelijke instructies per warp in één cyclus uitgeven, waardoor de instructie-niveau-parallellisatie toeneemt.
- **Pipelined Execution Units**: De uitvoeringseenheden zijn gepipelined, waardoor de SM een nieuwe bewerking op een eenheid kan starten voordat de vorige bewerking is voltooid.
- **Latency Hiding**: De SM kan tussen warps schakelen op basis van cycli, waardoor hij de latentie van geheugentoegangen en langdurende bewerkingen kan verbergen door instructies van andere warps uit te voeren.
Voorbeeld 7.1 toont een eenvoudige CUDA-kernel die elementgewijs optelling van twee vectoren uitvoert.
```cpp
__global__ void vectorAdd(int *a, int *b, int *c, int n) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (tid < n) {
c[tid] = a[tid] + b[tid];
}
}Voorbeeld 7.1: CUDA-kernel voor vectoroptelling.
In dit voorbeeld berekent elke thread in de kernel de som van de overeenkomstige elementen uit de invoervectoren a en b en slaat het resultaat op in de uitvoervector c. De SM voert deze kernel uit door elke thread toe te wijzen aan een CUDA-kern en warps van threads te plannen voor uitvoering op de beschikbare kernen. De load/store-eenheden worden gebruikt om de invoergegevens uit het globale geheugen op te halen en de resultaten terug te schrijven.
Warp Scheduling en Divergentie-afhandeling
EfHier is de Nederlandse vertaling van de gegeven Markdown-bestanden, waarbij de code-opmerkingen zijn vertaald:
Efficiënte warp-scheduling is cruciaal voor het maximaliseren van de prestaties van de SM. De warp-scheduler is verantwoordelijk voor het selecteren van warps die klaar zijn om uit te voeren en deze af te handelen naar de beschikbare uitvoeringseenheden. Het primaire doel van de warp-scheduler is om de uitvoeringseenheden bezig te houden door ervoor te zorgen dat er altijd warps beschikbaar zijn om uit te voeren.
De SM gebruikt een twee-niveau warp-scheduling mechanisme:
-
Warp-scheduling: De warp-scheduler selecteert warps die klaar zijn om uit te voeren op basis van een scheduling-beleid, zoals round-robin of oudste-eerst. De geselecteerde warps worden vervolgens afgehandeld naar de beschikbare uitvoeringseenheden.
-
Instructie-scheduling: Binnen elke warp plant de SM instructies in op basis van hun afhankelijkheden en de beschikbaarheid van uitvoeringseenheden. De SM kan meerdere onafhankelijke instructies uit dezelfde warp in één cyclus uitgeven om instructie-niveau parallelisme te maximaliseren.
Figuur 7.3 illustreert het twee-niveau warp-scheduling mechanisme.
Warp-pool
Warp 1 (Gereed)
Warp 2 (Wachtend)
Warp 3 (Gereed)
...
Warp N (Gereed)
|
v
Warp-scheduler
|
v
Dispatch-eenheid
|
v
UitvoeringseenhedenFiguur 7.3: Twee-niveau warp-scheduling mechanisme.
Één van de belangrijkste uitdagingen in warp-scheduling is het omgaan met vertakkingsdivergentie. In het SIMT-uitvoeringsmodel voeren alle threads in een warp dezelfde instructie in lockstep uit. Wanneer een warp echter een vertakking (bijvoorbeeld een if-else-instructie) tegenkomt, kunnen sommige threads het if-pad nemen, terwijl anderen het else-pad nemen. Deze situatie wordt vertakkingsdivergentie genoemd.
Om vertakkingsdivergentie te behandelen, gebruikt de SM een techniek genaamd predicatie. Wanneer een warp een divergente vertakking tegenkomt, voert de SM beide paden van de vertakking sequentieel uit, waarbij de threads die elk pad niet volgen, worden gemaskeerd. De resultaten worden vervolgens gecombineerd met behulp van predicaatregisters om ervoor te zorgen dat elke thread het juiste resultaat ontvangt.
Voorbeeld 7.2 toont een CUDA-kernel met een divergente vertakking.Hier is de Nederlandse vertaling van het Markdown-bestand:
__global__ void divergentKernel(int *data, int *result) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
if (data[tid] > 0) {
result[tid] = data[tid] * 2;
} else {
result[tid] = data[tid] * 3;
}
}Voorbeeld 7.2: CUDA-kernel met een divergerende tak.
In dit voorbeeld kan de takconditie data[tid] > 0 ertoe leiden dat sommige threads in een warp het if-pad volgen, terwijl anderen het else-pad volgen. De SM verwerkt deze divergentie door beide paden sequentieel uit te voeren en de inactieve threads in elk pad te maskeren.
Figuur 7.4 illustreert het predicatie-proces voor een warp met divergerende threads.
Warp (32 threads)
Thread 1: data[1] = 5, result[1] = 10
Thread 2: data[2] = -3, result[2] = -9
...
Thread 32: data[32] = 7, result[32] = 14
Divergente tak:
if (data[tid] > 0) {
result[tid] = data[tid] * 2;
} else {
result[tid] = data[tid] * 3;
}
Predicatie:
Stap 1: Voer het if-pad uit met masker
Thread 1: result[1] = 10
Thread 2: (gemaskeerd)
...
Thread 32: result[32] = 14
Stap 2: Voer het else-pad uit met masker
Thread 1: (gemaskeerd)
Thread 2: result[2] = -9
...
Thread 32: (gemaskeerd)
Uiteindelijk resultaat:
Thread 1: result[1] = 10
Thread 2: result[2] = -9
...
Thread 32: result[32] = 14Figuur 7.4: Predicatie-proces voor een warp met divergerende threads.
Door gebruik te maken van predicatie kan de SM omgaan met taktivergentie zonder de noodzaak van expliciete taktinstructies of controle-stroom-divergentie. Divergente takken kunnen echter nog steeds de prestaties beïnvloeden, omdat de SM beide paden sequentieel moet uitvoeren, wat de effectieve parallelliteit vermindert.
Register Bestand en Operand Collectors
Het register bestand is een kritische component van de SM, die snelle toegang biedt tot thread-privé registers. Elke SM heeft een groot register bestand om de vele actieve threads te ondersteunen en efficiënte context-switching tussen warps mogelijk te maken.Hier is de Nederlandse vertaling van het gegeven Markdown-bestand, waarbij de code-opmerkingen zijn vertaald:
In de NVIDIA Volta-architectuur heeft elk SM een 64 KB register-bestand, georganiseerd in 32 banken van 2 KB elk. Het register-bestand is ontworpen om hoge bandbreedte en lage latentie toegang te bieden ter ondersteuning van het grote aantal gelijktijdige threads.
Om conflicten tussen banken te minimaliseren en de prestaties te verbeteren, gebruikt de SM een techniek genaamd "operand collection". Operand collectors zijn gespecialiseerde units die operanden verzamelen uit de register-bestandsbanken en ze leveren aan de uitvoerings-units. Door gebruik te maken van operand collectors, kan de SM de impact van bankenconflicten verminderen en de benutting van de uitvoerings-units verbeteren.
Figuur 7.5 toont een vereenvoudigd diagram van het register-bestand en de operand collectors in een SM.
Register-bestand (64 KB)
Bank 1 (2 KB)
Bank 2 (2 KB)
...
Bank 32 (2 KB)
|
v
Operand Collectors
|
v
Uitvoerings-unitsFiguur 7.5: Register-bestand en operand collectors in een SM.
De operand collectors verzamelen operanden uit meerdere instructies en meerdere warps, waardoor de SM instructies uit verschillende warps in één cyclus naar de uitvoerings-units kan sturen. Dit helpt de latentie van toegangen tot het register-bestand te verbergen en verbetert de algehele doorvoer van de SM.
Voorbeeld 7.3 toont een CUDA-kernel die het dot-product van twee vectoren berekent.
__global__ void dotProduct(float *a, float *b, float *result, int n) {
// Gedeelde geheugen voor partiële sommen
__shared__ float partialSum[256];
int tid = threadIdx.x;
int i = blockIdx.x * blockDim.x + threadIdx.x;
// Initialiseer partiële som naar 0
partialSum[tid] = 0;
// Bereken partiële sommen
while (i < n) {
partialSum[tid] += a[i] * b[i];
i += blockDim.x * gridDim.x;
}
// Wacht tot alle threads klaar zijn
__syncthreads();
// Reduceer partiële sommen
for (int s = blockDim.x / 2; s > 0; s >>= 1) {
if (tid < s) {
partialSum[tid] += partialSum[tid + s];
}
__syncthreads();
}
// Schrijf uiteindelijke som naar globale geheugen
if (tid == 0) {
result[blockIdx.x] = partialSum[0];
}
}In dit voorbeeld berekent elke thread een partiële som van het dot-product met behulp van de toegewezenHere is the Dutch translation of the provided markdown file, with the code comments translated:
elementen uit de ingangsvectoren. De gedeeltelijke sommen worden opgeslagen in het gedeelde geheugenarray partialSum. Nadat alle threads hun gedeeltelijke sommen hebben berekend, wordt er een parallelle reductie uitgevoerd om de gedeeltelijke sommen op te tellen en het uiteindelijke dot product-resultaat te verkrijgen.
De operand collector speelt een cruciale rol in dit voorbeeld door efficiënt de operanden voor de gedeelde geheugens toegangen en de rekenkundige bewerkingen te verzamelen. Het helpt conflicten tussen banken te vermijden en verbetert het gebruik van de uitvoeringseenheden.
Conclusie
De streaming multiprocessor is de belangrijkste rekenkundige eenheid in moderne GPU-architecturen. Het ontwerp ervan richt zich op het maximaliseren van de doorvoer en het verbergen van geheugenlatentie door een combinatie van fijnmazige multithreading, SIMT-uitvoering en efficiënte operandverzameling.
Belangrijke componenten van de SM zijn de warp-scheduler, die warps selecteert voor uitvoering; de SIMT stack, die takdivergentie en -convergentie afhandelt; het register bestand en de operand collectors, die snelle toegang bieden tot thread-privé registers; en het gedeelde geheugen en de L1-cache, die low-latency gegevensuitwisseling en hergebruik mogelijk maken.
Naarmate GPU-architecturen zich blijven ontwikkelen, zal onderzoek op gebieden als takdivergentie-afhandeling, warp-planning en register bestandontwerp cruciaal zijn voor het verbeteren van de prestaties en efficiëntie van toekomstige GPU's. Nieuwe technieken zoals dynamische warp-formatie, thread block-compactie en operand reuse-caches hebben de potentie om de mogelijkheden van de SM aanzienlijk te verbeteren en nieuwe prestatieniveaus in parallelle computerwerkladen mogelijk te maken.