【发布时间】：2013-06-17 10:47:49
【问题描述】：
考虑这 3 个微不足道的最小内核。他们的寄存器使用率比我预期的要高很多。为什么？
答：
__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]
乙：
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]
问题：
为什么 empty 内核 A 和 B 使用 2 个寄存器？ CUDA 总是使用一个 implicit 寄存器，但是为什么要使用另外 2 个 explicit 寄存器呢？
内核 C 更令人沮丧。 10个寄存器？但是只有2个指针。这为指针提供了 2*2 = 4 个寄存器。即使有另外 2 个神秘的寄存器（由内核 A 和内核 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 个？
【问题讨论】：
标签： optimization assembly cuda gpu ptx