cuda - 最小的例子，註冊表使用率高

考慮這3個微不足道的小內核。他們的註冊使用量是，比我預期的要高很多。爲什麼？cuda - 最小的例子，註冊表使用率高

答：

__global__ void Kernel_A() 
{ 
//empty 
}

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

對應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 
}

相應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]

問：

爲什麼空內核A和B使用2個寄存器？ CUDA總是使用一個隱式的寄存器，但爲什麼有2個額外的顯式的寄存器被使用？內核C更令人沮喪。 10個寄存器？但只有2個指針。這爲指針提供了2 * 2 = 4個寄存器。即使還有另外2個神祕的寄存器（由Kernel A和Kernel B建議），這總共會有6個寄存器。仍然遠遠少於10！

如果你有興趣，這裏是內核A的ptx代碼內核B中的ptx代碼是完全一樣的，模整數值和變量名。

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

而對於內核的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: 
}

爲什麼它首先聲明一個本地內存變量（.local）？
爲什麼兩個指針（作爲函數參數給出）存儲在寄存器中？他們沒有特別的參數空間嗎？
也許這兩個函數參數指針屬於寄存器 - 這解釋了兩條線.reg .b64。但是.reg .s64是什麼？爲什麼在那裏？

它變得更糟的是：

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

給出

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

所以操縱該參數（指針）降低爲10〜6寄存器？

來源

2013-06-20 cmo

要做的第一點是，如果你擔心註冊表，不要看PTX代碼，因爲它不會告訴你任何事情。 PTX使用靜態單一賦值表單，編譯器發出的代碼不包含任何需要製作可運行機器代碼入口點的「裝飾」。

有了這樣的方式，讓我們來看看內核答：

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

有你的兩個寄存器。空內核不會生成零指令。

除此之外，我無法重現您所展示的內容。如果我查看發佈的內核C，我會得到這個（CUDA 5發佈編譯器）：

$ 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。相同的2個寄存器代碼到前兩個內核。

與同爲內核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; 
    ........................................

再者，2個寄存器。

爲了記錄在案，我現在用的NVCC的版本是：

$ 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

來源

2013-06-20 16:09:58 talonmies

我刪除了調試「-G」和「-g」從編譯器標誌...然後我得到了輸出相同你爲內核C. – cmo

我不能相信它。這是真的嗎？ – cmo

它會出現這樣。同樣，PTX不會告訴你你想知道什麼 - 調試器的支持會導致彙編器發出更多的設置代碼。這可能是額外寄存器的來源。 – talonmies

cuda - 最小的例子，註冊表使用率高

回答

相關問題