Registerpres og spills
Afvej genbrug mod occupancy
Registerpres og spills er en gratis CUDA Academy-lektion på CoddyKit. Dette er lektion 4 af 4. Du kan læse hele lektionen gratis nedenfor — og derefter øve dig praktisk i browseren med en indbygget kodeeditor og en AI-vejleder, der er tilgængelig døgnet rundt. Den er en del af læringsforløbet i CUDA Academy, og dine fremskridt synkroniseres på tværs af nettet og CoddyKit-appen. CUDA Academy-kurset indeholder 4 lektioner i alt.
Registre er værdifulde
Registre er den hurtigste lagerplads, en tråd har, men hver SM har kun et begrænset antal. Hvor intensivt en kerne bruger dem, kaldes registerpres.
En fælles pulje med fast størrelse
Alle aktive tråde trækker fra én registerfil pr. SM. Jo flere registre hver tråd har brug for, desto færre tråde kan forblive aktive samtidig.
Pres sænker belægningsgraden
Høj registerbrug begrænser direkte, hvor mange warps der kan være på en SM. Dette fald i belægningsgrad kan efterlade hardwaren med for lidt arbejde til at skjule latenstid.
Hvad er et registerspild
Når en tråd har brug for flere registre, end der findes, flytter compileren ekstra værdier ud. Denne overbelastning er et registerspild til langsommere hukommelse.
Spild placeres i lokal hukommelse
Værdier, der er flyttet ud, placeres i lokal hukommelse, som ligger i DRAM uden for chippen. Hvert spild erstatter en fri registeradgang med en langsom tur frem og tilbage.
Se dit registerforbrug
Bed compileren om at rapportere forbruget. Hvis du føjer --ptxas-options=-v til nvcc, udskriver den antal registre pr. tråd samt eventuelle spildbytes for hver kerne.
nvcc --ptxas-options=-v kernel.cuLæs tallene for spild
Rapporten viser spildskrivninger og -indlæsninger. Et spild med en værdi større end nul er et advarselstegn på, at din kerne betaler for langsom trafik til lokal hukommelse.
Begræns registre pr. tråd
Du kan angive en grænse ved kompilering. Flaget --maxrregcount begrænser antallet af registre pr. tråd og bytter en smule hastighed for højere belægningsgrad.
nvcc --maxrregcount=32 kernel.cuGiv et hint med __launch_bounds__
Et hint pr. kerne er ofte bedre. __launch_bounds__ fortæller compileren din blokstørrelse, så den kan budgettere registre til den belægningsgrad, du ønsker.
__global__ void __launch_bounds__(256)
myKernel() { /* ... */ }Formindsk det aktive værdisæt
Ofte kan du ganske enkelt holde færre værdier på én gang. Hvis du genberegner et billigt resultat eller indsnævrer en variables scope, reducerer du antallet af registre, der forbliver aktive.
Afvej genbrug og belægningsgrad
ILP og udfoldning øger presset, mens begrænsninger øger belægningsgraden. Kunsten er at finde den balance, der kører hurtigst, og kun profileringsværktøjet kan fortælle dig, hvor den findes.
Hurtigt tjek
Hvor ender værdierne, når en kerne løber tør for registre?
Opsummering: Hold presset under kontrol
Du har lært, at højt registerpres reducerer belægningsgraden og kan forårsage spild til langsom DRAM. Mål forbruget, begræns antallet af registre, og lad profileringsværktøjet guide dig. 🎯
Lær C++ med en AI-underviser — gratis
Skriv og kør rigtig kode i din browser, få øjeblikkelig hjælp fra en AI-underviser døgnet rundt, og fortsæt, hvor du slap, på web eller i appen.
- Kurser
- 30
- Lektioner
- 120
Ofte stillede spørgsmål
Er lektionen “Registerpres og spills” gratis?
Ja — hele teksten til “Registerpres og spills” kan læses gratis her på nettet. Hvis du vil øve dig interaktivt med en indbygget kodeeditor og en AI-vejleder døgnet rundt og få adgang til resten af CUDA Academy-kurset, skal du opgradere til CoddyKit PRO. CUDA Academy-kurset indeholder 4 lektioner i alt.
Hvad lærer jeg i “Registerpres og spills”?
Afvej genbrug mod occupancy Du øver dig i CUDA Academy med praktisk kode, som du kører direkte i browseren, og en AI-vejleder døgnet rundt besvarer dine spørgsmål, mens du arbejder dig gennem lektionen.
Skal jeg have erfaring for at begynde på CUDA Academy?
Der kræves ingen tidligere erfaring. CUDA Academy på CoddyKit er tilrettelagt for både begyndere og øvede, så du kan starte her eller fra begyndelsen og lære i dit eget tempo. Dette er lektion 4 af 4.
Hvor lang tid tager lektionen “Registerpres og spills”?
De fleste CoddyKit-lektioner tager cirka 5–10 minutter. Hver lektion er kort og interaktiv, så du gør løbende fremskridt og kan fortsætte, hvor du slap – på både web og app.
Kan jeg skrive og køre kode i denne CUDA Academy-lektion?
Ja. Alle CUDA Academy-lektioner har en indbygget kodeeditor, så du kan skrive og køre rigtig kode direkte i din browser og få øjeblikkelig feedback fra AI – uden lokal opsætning.
Alle lektioner i dette kursus
- Parallelitet på instruktionsniveau
- Loop-unrolling med #pragma unroll
- Vektoriserede indlæsninger med float4
- Registerpres og spills