Registertryck och spill

Register är den snabbaste lagringen en tråd har tillgång till, men varje SM har bara ett begränsat antal. Hur intensivt en kernel använder dem kallas register pressure.

Varje resident tråd hämtar register från en gemensam register file per SM. Ju fler register varje tråd behöver, desto färre trådar kan vara residenta samtidigt.

Hög registeranvändning begränsar direkt hur många warps som får plats på en SM. Den minskningen av occupancy kan lämna hårdvaran med för lite arbete för att dölja latens.

När en tråd behöver fler register än vad som finns flyttar kompilatorn ut extra värden. Detta spill hamnar som en register spill i långsammare minne.

Spillade värden går till local memory, som finns i DRAM utanför chippet. Varje spill ersätter en snabb registeråtkomst med en långsam tur och retur.

Be kompilatorn rapportera användningen. Om Ni lägger till --ptxas-options=-v i nvcc skrivs antal register per tråd och eventuella spillbyte ut för varje kernel.

Rapporten listar spill stores och loads. Ett spill-antal som inte är noll är ett varningstecken på att Er kernel betalar för långsam trafik till local memory.

Ni kan ange ett tak vid kompilering. Flaggan --maxrregcount begränsar antalet register per tråd och byter lite hastighet mot högre occupancy.

En ledtråd per kernel är ofta bättre. __launch_bounds__ talar om för kompilatorn vilken blockstorlek Ni använder, så att den kan dimensionera registeranvändningen för den occupancy Ni vill ha.

Ofta kan Ni helt enkelt hålla färre värden samtidigt. Om Ni räknar om ett billigt resultat eller begränsar variabelns scope minskar antalet register som förblir aktiva.

ILP och loop-unrolling ökar pressure, medan begränsningar ökar occupancy. Konsten är att hitta den balans som körs snabbast, och bara profileraren kan visa var den finns.

Ni har lärt Er att hög register pressure minskar occupancy och kan orsaka spill till långsamt DRAM. Mät användningen, begränsa registren och låt profileraren vägleda Er. 🎯