cuda - minimalny przykład, wysokie użycie rejestru

Question

Rozważ te 3 trywialne, minimalne jądra. Ich użycie rejestru wynosi dużo wyżej niż się spodziewam. Czemu?

A:

__global__ void Kernel_A() 
{ 
//empty 
} 

odpowiadające PTX:

ptxas info : Compiling entry function '_Z8Kernel_Av' for 'sm_20' 
ptxas info : Function properties for _Z8Kernel_Av 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 2 registers, 32 bytes cmem[0] 

B:

template<uchar effective_bank_width> 
__global__ void Kernel_B() 
{ 
//empty 
} 

template 
__global__ void Kernel_B<1>(); 

odpowiadające PTX:

ptxas info : Compiling entry function '_Z8Kernel_BILh1EEvv' for 'sm_20' 
ptxas info : Function properties for _Z8Kernel_BILh1EEvv 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 2 registers, 32 bytes cmem[0] 

C:

template<uchar my_val> 
__global__ void Kernel_C 
     (uchar *const device_prt_in, 
     uchar *const device_prt_out) 
{ 
//empty 
} 

odpowiadające PTX:

ptxas info : Compiling entry function '_Z35 Kernel_CILh1EEvPhS0_' for 'sm_20' 
ptxas info : Function properties for _Z35 Kernel_CILh1EEvPhS0_ 
    16 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 10 registers, 48 bytes cmem[0] 

---

Pytanie:

Dlaczego pusty jądra A i B używają 2 rejestrów? CUDA zawsze używa jednego niejawnego rejestru niejawnego, ale dlaczego są używane 2 dodatkowe rejestry jawne?

Kernel C jest jeszcze bardziej frustrujący. 10 rejestrów? Ale są tylko 2 wskaźniki. Daje to 2 * 2 = 4 rejestry dla wskaźników. Nawet jeśli istnieją dodatkowo 2 tajemnicze rejestry (sugerowane przez Kernel A i Kernel B), dałoby to 6 całkowite. Still znacznie mniej niż 10!

---

W przypadku jesteś zainteresowany, tutaj jest kod ptx dla Kernel A. Kod ptx dla Kernel B jest dokładnie taka sama, modulo wartości całkowite i nazw zmiennych.

.visible .entry _Z8Kernel_Av( 
) 
{   
     .loc 5 19 1 
func_begin0: 
     .loc 5 19 0 

     .loc 5 19 1 

func_exec_begin0: 
     .loc 5 22 2 
     ret; 
tmp0: 
func_end0: 
} 

I Kernel C ...

.weak .entry _Z35Kernel_CILh1EEvPhS0_(
     .param .u64 _Z35Kernel_CILh1EEvPhS0__param_0, 
     .param .u64 _Z35Kernel_CILh1EEvPhS0__param_1 
) 
{ 
     .local .align 8 .b8  __local_depot2[16]; 
     .reg .b64  %SP; 
     .reg .b64  %SPL; 
     .reg .s64  %rd<3>; 


     .loc 5 38 1 
func_begin2: 
     .loc 5 38 0 

     .loc 5 38 1 

     mov.u64   %SPL, __local_depot2; 
     cvta.local.u64 %SP, %SPL; 
     ld.param.u64 %rd1, [_Z35Kernel_CILh1EEvPhS0__param_0]; 
     ld.param.u64 %rd2, [_Z35Kernel_CILh1EEvPhS0__param_1]; 
     st.u64 [%SP+0], %rd1; 
     st.u64 [%SP+8], %rd2; 
func_exec_begin2: 
     .loc 5 836 2 
tmp2: 
     ret; 
tmp3: 
func_end2: 
} 

1. Dlaczego najpierw zadeklarować zmiennej lokalnej pamięci (.local)?
2. Dlaczego te dwa wskaźniki (podane jako argumenty funkcji) są przechowywane w rejestrach? Czy nie ma dla nich specjalnej przestrzeni param?
3. Być może dwa wskaźniki funkcji argumentu należą do rejestrów - to wyjaśnia dwie linie: .reg .b64. Ale jaka jest linia .reg .s64? Dlaczego tam jest?

---

Robi się jeszcze gorzej:

D:

template<uchar my_val> 
__global__ void Kernel_D 
     (uchar * device_prt_in, 
     uchar *const device_prt_out) 
{ 
    device_prt_in = device_prt_in + blockIdx.x*blockDim.x + threadIdx.x; 
} 

daje

ptxas info : Used 6 registers, 48 bytes cmem[0] 

więc manipulowanie argument (wskaźnik) zmniejsza się od 10 do 6 rejestry?

Answer 1

Po pierwsze, jeśli martwisz się o rejestry, nie patrz na kod PTX, ponieważ nic ci nie powie. PTX używa statycznego formularza z pojedynczym przydziałem, a kod emitowany przez kompilator nie zawiera żadnej "dekoracji" wymaganej do utworzenia punktu wejścia do kodu maszynowego.

Mając to na uboczu, spójrzmy na jądra A:

$ nvcc -arch=sm_20 -m64 -cubin -Xptxas=-v null.cu 
ptxas info : 0 bytes gmem 
ptxas info : Compiling entry function '_Z8Kernel_Av' for 'sm_20' 
ptxas info : Function properties for _Z8Kernel_Av 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 2 registers, 32 bytes cmem[0] 

$ cuobjdump -sass null.cubin 

    code for sm_20 
     Function : _Z8Kernel_Av 
    /*0000*/  /*0x00005de428004404*/  MOV R1, c [0x1] [0x100]; 
    /*0008*/  /*0x00001de780000000*/  EXIT; 
     ............................. 

Są twoi dwa rejestry. Puste jądra nie generują zerowych instrukcji.

Poza tym nie mogę odtworzyć tego, co pokazałeś. Jeśli spojrzeć na swoje jądra C, jak pisał, otrzymuję ten (CUDA 5 zwalniający kompilatora):

$ nvcc -arch=sm_20 -m64 -cubin -Xptxas=-v null.cu 
ptxas info : 0 bytes gmem 
ptxas info : Compiling entry function '_Z8Kernel_CILh1EEvPhS0_' for 'sm_20' 
ptxas info : Function properties for _Z8Kernel_CILh1EEvPhS0_ 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 2 registers, 48 bytes cmem[0] 


$ cuobjdump -sass null.cubin 

code for sm_20 
    Function : _Z8Kernel_CILh1EEvPhS0_ 
/*0000*/  /*0x00005de428004404*/  MOV R1, c [0x1] [0x100]; 
/*0008*/  /*0x00001de780000000*/  EXIT; 
    ........................................ 

tj. identyczny 2 kod rejestracyjny do pierwszych dwóch jąder.

to samo jądro D:

$ nvcc -arch=sm_20 -m64 -cubin -Xptxas=-v null.cu 
ptxas info : 0 bytes gmem 
ptxas info : Compiling entry function '_Z8Kernel_DILh1EEvPhS0_' for 'sm_20' 
ptxas info : Function properties for _Z8Kernel_DILh1EEvPhS0_ 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 2 registers, 48 bytes cmem[0] 

$ cuobjdump -sass null.cubin 
code for sm_20 
    Function : _Z8Kernel_DILh1EEvPhS0_ 
/*0000*/  /*0x00005de428004404*/  MOV R1, c [0x1] [0x100]; 
/*0008*/  /*0x00001de780000000*/  EXIT; 
    ........................................ 

Ponownie, 2 rejestrów.

Dla przypomnienia, wersja nvcc używam jest:

$ nvcc --version 
nvcc: NVIDIA (R) Cuda compiler driver 
Copyright (c) 2005-2012 NVIDIA Corporation 
Built on Fri_Sep_28_16:10:16_PDT_2012 
Cuda compilation tools, release 5.0, V0.2.1221