Registerdruk en spills

Registers zijn de snelste opslag waarover een thread beschikt, maar elke SM heeft er slechts een beperkt aantal. Hoe intensief een kernel ze gebruikt, heet registerdruk.

Elke actieve thread gebruikt registers uit één registerbestand per SM. Hoe meer registers elke thread nodig heeft, hoe minder threads tegelijk actief kunnen blijven.

Intensief registergebruik beperkt rechtstreeks hoeveel warps op een SM passen. Door die daling van de bezettingsgraad kan de hardware te weinig werk hebben om latentie te verbergen.

Wanneer een thread meer registers nodig heeft dan er beschikbaar zijn, verplaatst de compiler extra waarden naar buiten. Deze overloop is een registerspill naar tragere geheugens.

Gespilde waarden gaan naar lokaal geheugen, dat zich in DRAM buiten de chip bevindt. Elke spill vervangt een snelle registertoegang door een trage heen-en-terugbewerking.

Vraag de compiler om het gebruik te rapporteren. Als u --ptxas-options=-v toevoegt aan nvcc, worden het aantal registers per thread en eventuele spillbytes voor elke kernel afgedrukt.

Het rapport vermeldt spill-opslagbewerkingen en -laadbewerkingen. Elk niet-nul spill-aantal is een waarschuwing dat uw kernel betaalt voor traag verkeer naar lokaal geheugen.

U kunt tijdens het compileren een bovengrens instellen. De vlag --maxrregcount beperkt het aantal registers per thread, in ruil voor een hogere bezettingsgraad.

Een aanwijzing per kernel is vaak beter. __launch_bounds__ vertelt de compiler hoe groot uw blok is, zodat deze registers kan begroten voor de gewenste bezettingsgraad.

Vaak kunt u eenvoudig minder waarden tegelijk vasthouden. Een goedkope uitkomst opnieuw berekenen of de scope van een variabele verkleinen, vermindert het aantal registers dat actief blijft.

ILP en unrolling verhogen de registerdruk, terwijl begrenzingen de bezettingsgraad verhogen. De kunst is de balans te vinden waarmee u het snelst werkt; alleen de profiler kan u vertellen waar die ligt.

U hebt geleerd dat een hoge registerdruk de bezettingsgraad verlaagt en spills naar trage DRAM kan veroorzaken. Meet het gebruik, beperk het aantal registers en laat de profiler u sturen. 🎯