Zamieszanie z kodem i pamięcią CUDA PTX

Question

:) Podczas gdy próbowałem zarządzać zasobami jądra, zdecydowałem się zaglądać do PTX, ale jest kilka rzeczy, których nie rozumiem. Tutaj jest bardzo prosty jądro pisałem:

__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; 
} 

Potem skompilowany przy użyciu: nvcc --ptxas-options=-v -keep main.cu i mam to wyjście na konsoli:

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

i otrzymaną PTX jest następujący:

.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 

Obecnie istnieje kilka rzeczy, które nie rozumieją:

• Według zespołu ptx 4 + 5 + 8 + 5 = 22 rejestry są używane. Dlaczego więc podczas kompilacji jest napisane used 2 registers?
• Patrząc na zespół zdałem sobie sprawę, że typem danych threadId, blockId itp jest u16. Czy jest to zdefiniowane w specyfikacji CUDA? Lub może to być różne dla różnych wersji sterownika CUDA?
• Czy ktoś może mi wytłumaczyć tę linię: mul.wide.u16 %r1, %rh1, %rh2;? %r1 jest u32, dlaczego wide zamiast u32 jest używany?
• W jaki sposób są wybierane nazwy rejestrów? W moim wazonie rozumiem część %r, ale nie rozumiem części: h, (null), d. Czy jest wybierany na podstawie długości typu danych? tj: h dla 16bit, zero dla 32bit, d dla 64bit?
• Jeśli zastąpię ostatnie 2 linie mojego jądra tym out[idx] = in[idx];, to kiedy skompiluję program, zostanie napisane, że są używane 3 rejestry! Jak teraz można korzystać z większej liczby rejestrów?

Proszę zignorować fakt, że moje jądro testowe nie sprawdza, czy indeks tablicy jest poza zakresem.

Dziękuję bardzo.

Answer 1

PTX to język pośredni zaprojektowany do przenoszenia na wiele architektur GPU. Kompilator zostaje skompilowany przez komponent kompilatora PTXAS do końcowego kodu maszynowego, określanego również jako SASS, dla konkretnej architektury. Opcja nvcc -Xptxas -v powoduje, że PTXAS raportuje różne statystyki dotyczące wygenerowanego kodu maszynowego, w tym liczbę fizycznych rejestrów używanych w kodzie maszynowym. Możesz sprawdzić kod urządzenia, rozmontowując go za pomocą cuobjdump --dump-sass.

Liczba rejestrów, które widzi się w kodzie PTX, nie ma znaczenia, ponieważ są to rejestry wirtualne. Kompilator CUDA generuje kod PTX w tak zwanej postaci SSA (statyczne pojedyncze przypisanie, patrz http://en.wikipedia.org/wiki/Static_single_assignment_form). Oznacza to zasadniczo, że każdemu nowemu zapisanemu wynikowi przypisywany jest nowy rejestr.

Instrukcja mul.wide jest opisana w specyfikacji PTX, której aktualna wersja (3.1) znajduje się tutaj: http://docs.nvidia.com/cuda/parallel-thread-execution/index.html. W twoim przykładowym kodzie przyrostek .u16 oznacza, że ​​mnoży on dwie niepodpisane 16-bitowe liczby i zwraca niepodpisany 32-bitowy wynik, tj. Oblicza pełny produkt o podwójnej szerokości operandów źródłowych.

Wirtualne rejestry w PTX są pisane na maszynie, ale ich nazwy można wybierać dowolnie, niezależnie od typu. Wydaje się, że kompilator CUDA podąża za pewnymi konwencjami, które (według mojej wiedzy) nie są dokumentowane, ponieważ są artefaktami implementacji wewnętrznej. Patrząc na mnóstwo kodu PTX, jasne jest, że rejestr generuje aktualnie wygenerowane informacje o typie kodowania, co można zrobić dla łatwego debugowania: p<num> jest używany dla predykatów, r<num> dla 32-bitowych liczb całkowitych, rd<num> dla liczb całkowitych 64-bitowych, f<num> dla 32-bitowych zmiennoprzecinkowych i fd<num> dla 64-bitowych podwójnych. Możesz to łatwo zobaczyć, patrząc na dyrektywy .reg w kodzie PTX, które tworzą te wirtualne rejestry.