Registertrykk og spills

Registre er den raskeste lagringen en tråd har tilgang til, men hver SM har bare et begrenset antall. Hvor intensivt en kernel bruker dem, kalles registertrykk.

Alle aktive tråder henter ressurser fra én registerfil per SM. Jo flere registre hver tråd trenger, desto færre tråder kan være aktive samtidig.

Høy registerbruk begrenser direkte hvor mange warps som får plass på en SM. Denne nedgangen i occupancy kan gi maskinvaren for lite arbeid til å skjule latenstid.

Når en tråd trenger flere registre enn det finnes, flytter kompilatoren ut ekstra verdier. Denne overflyten er en registerspill til tregere minne.

Verdier som spills, havner i lokalt minne, som ligger i DRAM utenfor brikken. Hver spill erstatter et ledig registeroppslag med en langsom tur frem og tilbake.

Be kompilatoren rapportere bruken. Når De legger til --ptxas-options=-v i nvcc, skriver den ut antall registre per tråd og eventuelle spillbyte for hver kernel.

Rapporten viser lagringer og innlastinger for spill. Ethvert spill-antall som ikke er null, er et varselsignal om at kernelen betaler for treg trafikk til lokalt minne.

De kan angi en øvre grense ved kompilering. Flagget --maxrregcount begrenser antallet registre per tråd og bytter litt hastighet mot høyere occupancy.

Et hint per kernel er ofte bedre. __launch_bounds__ forteller kompilatoren blokkstørrelsen, slik at den kan budsjettere registrene for den occupancy-en De ønsker.

Ofte kan De ganske enkelt holde færre verdier samtidig. Ved å beregne et rimelig resultat på nytt eller begrense variabelens scope reduserer De hvor mange registre som forblir aktive.

ILP og utrulling øker trykket, mens begrensninger øker occupancy. Kunsten er å finne balansen som gir høyest hastighet, og bare profilereren kan fortelle Dem hvor den ligger.

De har lært at høyt registertrykk reduserer occupancy og kan føre til spill til treg DRAM. Mål bruken, begrens antallet registre, og la profilereren veilede Dem. 🎯