考虑这 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]
电话:
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 个神秘寄存器(由内核 A 和内核 B 建议),这也将得到 6 个。还远低于10!
如果您有兴趣,这里是
ptx
内核 A 的代码。ptx
内核 B 的代码完全相同,以整数值和变量名称为模。.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
线?为什么会在那里? 情况变得更糟:
电话:
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 个寄存器?
最佳答案
首先要指出的是,如果您担心寄存器,请不要查看 PTX 代码,因为它不会告诉您任何信息。 PTX 使用静态单赋值形式,编译器发出的代码不包含任何创建可运行机器代码入口点所需的“装饰”。
顺便说一下,让我们看看内核 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;
.............................
有你的两个寄存器。空内核不会产生零指令。
除此之外,我无法重现您所展示的内容。如果我查看您发布的内核 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
关于optimization - cuda - 最小的例子,高寄存器使用率,我们在Stack Overflow上找到一个类似的问题: https://stackoverflow.com/questions/17216012/