Tracciamento dell'uso del registro del kernel Cuda

Question

Sto cercando di rintracciare l'utilizzo del registro e ho trovato uno scenario interessante. Si consideri il seguente fonte:

#define OL 20 
#define NHS 10 

__global__ void loop_test(float ** out, const float ** in,int3 gdims,int stride){ 

     const int idx = blockIdx.x*blockDim.x + threadIdx.x; 
     const int idy = blockIdx.y*blockDim.y + threadIdx.y; 
     const int idz = blockIdx.z*blockDim.z + threadIdx.z; 

     const int index = stride*gdims.y*idz + idy*stride + idx; 
     int i = 0,j =0; 
     float sum =0.f; 
     float tmp; 
     float lf; 
     float u2, tW; 

     u2 = 1.0; 
     tW = 2.0; 

     float herm[NHS]; 

     for(j=0; j < OL; ++j){ 
       for(i = 0; i < NHS; ++i){ 
         herm[i] += in[j][index]; 
       } 
     } 

     for(j=0; j<OL; ++j){ 
       for(i=0;i<NHS; ++i){ 
         tmp = sum + herm[i]*in[j][index]; 
         sum = tmp; 
       } 
       out[j][index] = sum; 
       sum =0.f; 
     } 

} 

Come nota a margine sulla sorgente - la somma parziale che potevo fare + =, ma stava giocando con il modo che cambia che gli effetti registrano utilizzo (sembra che non si - aggiunge solo un'istruzione extra in movimento). Inoltre questa sorgente è orientata per l'accesso alla memoria mappata allo spazio 3D.

Concludendo i registri sembrerebbe che ci siano 22 registri (credo che un float [N] occupi N + 1 registri - correggimi se sono wronge) in base alle dichiarazioni.

Tuttavia compilazione con:

nvcc -cubin -arch=sm_20 -Xptxas="-v" src/looptest.cu 

rendimenti:

0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 25 registers, 72 bytes cmem[0] 

Ok, quindi il numero è diverso da quello che si 'aspetta'. Inoltre, se compilato con:

nvcc -cubin -arch=sm_13 -Xptxas="-v" src/looptest.cu 

L'utilizzo registro è lontano meno - 8 per l'esattezza (a quanto pare a causa di forte aderenza in sm_20 di sm_13 a standard IEEE matematiche in virgola mobile?):

ptxas info : Compiling entry function '_Z9loop_testPPfPPKfS2_4int3i' for 'sm_13' 
ptxas info : Used 17 registers, 40+16 bytes smem, 8 bytes cmem[1] 

Come nota finale, cambiare l'OL macro a 40, e improvvisamente:

0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 28 registers, 72 bytes cmem[0] 

In conclusione vorrei sapere dove i registri sono in corso mangiato e cosa risulta dalle osservazioni di coppia che ho fatto.

non ho abbastanza esperienza con il gruppo per ottenere attraverso un cuobjdump - la risposta certamente giace sepolto in là - forse qualcuno mi può illuminare su quello che dovrebbe essere alla ricerca o mi mostra una guida su come affrontare il discarica di assemblaggio.

Answer 1

L'utilizzo del registro non ha necessariamente una stretta correlazione con il numero di variabili.

Il compilatore tenta di valutare il vantaggio di velocità di mantenere una variabile in un registro tra due punti di utilizzo nel codice confrontando il guadagno potenziale in un singolo kernel con il costo per tutti i kernel in esecuzione simultaneamente a causa della presenza di meno registri disponibile nel pool di registri. (Un Fermi SM ha 32768 registri). Quindi, non è sorprendente se la modifica del codice causi fluttuazioni impreviste nel numero di registri utilizzati.

Si dovrebbe solo essere preoccupati per l'utilizzo del registro se il profiler dice che la vostra occupazione è limitata dall'uso del registro. In tal caso, è possibile utilizzare l'impostazione --maxrregcount per ridurre il numero di registri utilizzati da un singolo kernel per vedere se migliora la velocità generale di esecuzione.

Per ridurre il numero di registri utilizzati da un kernel, è possibile provare a mantenere l'uso delle variabili il più locale possibile. Ad esempio, se lo fai:

set variable 1 
set variable 2 
use variable 1 
use variable 2 

Ciò potrebbe causare l'utilizzo di 2 registri.Mentre, se:

set variable 1 
use variable 1 
set variable 2 
use variable 2 

Ciò potrebbe causare l'utilizzo di 1 registro.

Answer 2

sm_20 e sm_13 sono architetture molto diverse, con un design ISA (instruction set set) molto diverso. La principale differenza che causa l'aumento dell'utilizzo del registro che vedi è che sm_1x ha registri di indirizzi speciali, mentre sm_2x e successivi no. Invece, gli indirizzi sono memorizzati in registri generici proprio come i valori, il che significa che la maggior parte dei programmi richiede più registri su sm_2x che su sm_1x.

sm_20 ha anche il doppio della dimensione del file di registro di sm_13, per compensare questo effetto.