cuda - exemple minimal, utilisation élevée du registre

Question

Considérons ces 3 noyaux minimes et minimaux. Leur utilisation du registre est beaucoup plus élevé que ce que j'attends. Pourquoi?

A:

__global__ void Kernel_A() 
{ 
//empty 
} 

PTX correspondant:

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>(); 

PTX correspondant:

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 
} 

PTX correspondant:

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] 

---

Question:

Pourquoi est-ce vide Les noyaux A et B utilisent 2 registres? CUDA utilise toujours un registre implicite, mais pourquoi 2 registres explicites sont-ils utilisés?

Le noyau C est encore plus frustrant. 10 registres? Mais il n'y a que 2 pointeurs. Cela donne 2 * 2 = 4 registres pour les pointeurs. Même s'il y a en plus 2 registres mystérieux (suggérés par Kernel A et Kernel B), cela donnerait 6 au total. Toujours beaucoup moins que 10!

---

Si vous êtes intéressé, voici le code ptx pour le noyau A. Le code ptx pour le noyau B est exactement la même, modulo les valeurs entières et les noms de variables.

.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: 
} 

Et pour le noyau 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. Pourquoi ne première Déclarez une variable mémoire locale (.local)?
2. Pourquoi les deux pointeurs (donnés comme arguments de fonction) sont-ils stockés dans des registres? N'y a-t-il pas un espace param spécifique pour eux?
3. Les deux pointeurs d'argument de fonction appartiennent peut-être aux registres, ce qui explique les deux lignes .reg .b64. Mais quelle est la ligne .reg .s64? Pourquoi est-ce là?

---

Il y a pire encore:

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

donne

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

manipulation L'argument (un pointeur) diminue de 10 à 6 registres?

Answer 1

Le premier point à faire est que si vous êtes préoccupé par les registres, ne regardez pas le code PTX, car il ne vous dira rien. PTX utilise un formulaire d'affectation unique statique et le code émis par le compilateur n'inclut aucune des "décorations" requises pour créer un point d'entrée de code machine exécutable.

Avec cela de la manière, regardons le noyau 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; 
     ............................. 

Il y a vos deux registres. Les noyaux vides ne produisent pas d'instructions nulles. Au-delà, je ne peux pas reproduire ce que vous avez montré. Si je regarde votre noyau C comme posté, je reçois ce (CUDA 5 compilateur version):

$ 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; 
    ........................................ 

ie. code de registre 2 identique aux deux premiers noyaux.

et de même pour le noyau 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; 
    ........................................ 

Encore une fois, 2 entrées.

Pour mémoire, la version nvcc J'utilise est:

$ 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