Pressione sui registri e spill

I registri sono la memoria più veloce a disposizione di un thread, ma ogni SM ne contiene un numero limitato. L'intensità con cui un kernel li utilizza si chiama pressione sui registri.

Ogni thread residente attinge a un unico file di registri per SM. Più registri richiede ogni thread, meno thread possono rimanere residenti contemporaneamente.

Un uso elevato dei registri limita direttamente il numero di warp che possono entrare in un SM. Questa riduzione dell'occupancy può lasciare all'hardware troppo poco lavoro per nascondere la latenza.

Quando un thread necessita di più registri di quanti ne siano disponibili, il compilatore sposta all'esterno i valori aggiuntivi. Questo overflow è un register spill nella memoria più lenta.

I valori soggetti a spill vengono trasferiti nella memoria locale, che risiede nella DRAM off-chip. Ogni spill sostituisce l'accesso a un registro libero con un lento viaggio di andata e ritorno.

Chieda al compilatore di riportare l'utilizzo. Aggiungendo --ptxas-options=-v a nvcc vengono stampati i registri per thread e gli eventuali byte soggetti a spill per ogni kernel.

Il report elenca gli store e i load degli spill. Qualsiasi conteggio spill diverso da zero indica che il kernel sta pagando il costo del traffico lento verso la memoria locale.

Può impostare un limite in fase di compilazione. Il flag --maxrregcount limita i registri per thread, scambiando un po' di velocità con una maggiore occupancy.

Spesso è preferibile un suggerimento specifico per kernel. __launch_bounds__ comunica al compilatore la dimensione del blocco, così può allocare i registri in funzione dell'occupancy desiderata.

Spesso può semplicemente mantenere meno valori contemporaneamente. Ricalcolare un risultato poco costoso o restringere lo scope di una variabile riduce il numero di registri che rimangono attivi.

ILP e unrolling aumentano la pressione, mentre i limiti aumentano l'occupancy. La difficoltà consiste nel trovare il bilanciamento che offre le prestazioni migliori, e solo il profiler può indicare qual è.

Dove finiscono i valori quando un kernel esaurisce i registri?

Ha imparato che un'elevata pressione sui registri riduce l'occupancy e può causare spill verso la lenta DRAM. Misuri l'utilizzo, limiti i registri e lasci che sia il profiler a guidarLa. 🎯