Confusione con il codice PTX CUDA e memoria di registro

Question

:) Mentre stavo cercando di gestire le risorse del mio kernel ho deciso di guardare al PTX ma ci sono un paio di cose che non capisco. Ecco un semplice kernel ho scritto:

__global__ 
void foo(float* out, float* in, uint32_t n) 
{ 
    uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x; 
    uint32_t one = 5; 
    out[idx] = in[idx]+one; 
} 

Poi ho compilato utilizzando: nvcc --ptxas-options=-v -keep main.cu e ho ottenuto questo output sulla console:

ptxas info : 0 bytes gmem 
ptxas info : Compiling entry function '_Z3fooPfS_j' for 'sm_10' 
ptxas info : Used 2 registers, 36 bytes smem 

E il PTX risultante è il seguente:

.entry _Z3fooPfS_j (
      .param .u64 __cudaparm__Z3fooPfS_j_out, 
      .param .u64 __cudaparm__Z3fooPfS_j_in, 
      .param .u32 __cudaparm__Z3fooPfS_j_n) 
    { 
    .reg .u16 %rh<4>; 
    .reg .u32 %r<5>; 
    .reg .u64 %rd<8>; 
    .reg .f32 %f<5>; 
    .loc 15 17 0 
$LDWbegin__Z3fooPfS_j: 
    .loc 15 21 0 
    mov.u16  %rh1, %ctaid.x; 
    mov.u16  %rh2, %ntid.x; 
    mul.wide.u16 %r1, %rh1, %rh2; 
    cvt.u32.u16  %r2, %tid.x; 
    add.u32  %r3, %r2, %r1; 
    cvt.u64.u32  %rd1, %r3; 
    mul.wide.u32 %rd2, %r3, 4; 
    ld.param.u64 %rd3, [__cudaparm__Z3fooPfS_j_in]; 
    add.u64  %rd4, %rd3, %rd2; 
    ld.global.f32 %f1, [%rd4+0]; 
    mov.f32  %f2, 0f40a00000;  // 5 
    add.f32  %f3, %f1, %f2; 
    ld.param.u64 %rd5, [__cudaparm__Z3fooPfS_j_out]; 
    add.u64  %rd6, %rd5, %rd2; 
    st.global.f32 [%rd6+0], %f3; 
    .loc 15 22 0 
    exit; 
$LDWend__Z3fooPfS_j: 
    } // _Z3fooPfS_j 

Ora ci sono alcune cose che non capisco:

• In base all'assieme ptx vengono utilizzati 4 + 5 + 8 + 5 = 22 registri. Quindi perché dice used 2 registers durante la compilazione?
• Guardando l'assemblea mi sono reso conto che il tipo di dati di threadId, blockId ecc è u16. È definito nella specifica CUDA? O questo può variare tra le diverse versioni del driver CUDA?
• Qualcuno può spiegarmi questa riga: mul.wide.u16 %r1, %rh1, %rh2;? %r1 è u32, perché viene utilizzato wide anziché u32?
• Come vengono scelti i nomi dei registri? Nel mio vaso capisco la parte %r ma non capisco la parte h, (null), d. Viene scelto in base alla lunghezza del tipo di dati? vale a dire: h per 16 bit, null per 32 bit, d per 64 bit?
• Se sostituisco le ultime 2 linee del mio kernel con questo out[idx] = in[idx];, allora quando compilo il programma dice che sono usati 3 registri! Com'è possibile usare più registri adesso?

Si prega di ignorare il fatto che il mio kernel di test non controlla se l'indice dell'array è fuori limite.

Grazie mille.

Answer 1

PTX è un linguaggio intermedio progettato per essere trasferito su più architetture GPU. Viene compilato dal componente del compilatore PTXAS nel codice macchina finale, indicato anche come SASS, per una particolare architettura. L'opzione nvcc -Xptxas -v consente a PTXAS di riportare varie statistiche sul codice macchina generato, incluso il numero di registri fisici utilizzati nel codice macchina. È possibile ispezionare il codice macchina smontandolo con cuobjdump --dump-sass.

Quindi il numero di registri che si vedono utilizzati nel codice PTX non ha alcun significato, poiché si tratta di registri virtuali. Il compilatore CUDA genera il codice PTX in quello che è noto come modulo SSA (assegnazione singola statica, vedere http://en.wikipedia.org/wiki/Static_single_assignment_form). Ciò significa fondamentalmente che a ogni nuovo risultato scritto viene assegnato un nuovo registro.

L'istruzione mul.wide è descritta nella specifica PTX, la versione corrente di cui (3.1) è disponibile qui: http://docs.nvidia.com/cuda/parallel-thread-execution/index.html. Nel codice di esempio, il suffisso .u16 significa che moltiplica due quantità a 16 bit senza segno e restituisce un risultato a 32 bit senza segno, vale a dire calcola il prodotto completo a doppia larghezza degli operandi di origine.

I registri virtuali in PTX sono digitati, ma i loro nomi possono essere scelti liberamente, indipendentemente dal tipo. Il compilatore CUDA sembra seguire alcune convenzioni che sono (a mia conoscenza) non documentate poiché sono artefatti di implementazione interna. Guardando un gruppo di codice PTX è chiaro che i nomi dei registri generano attualmente informazioni sul tipo di codifica, questo può essere fatto per semplificare il debug: p<num> viene utilizzato per i predicati, r<num> per numeri interi a 32 bit, rd<num> per numeri interi a 64 bit, f<num> per i float a 32 bit e fd<num> per i doppi a 64 bit. Puoi facilmente vederlo da solo osservando le direttive .reg nel codice PTX che creano questi registri virtuali.