Cuda -最小的例子，高寄存器使用率

考虑这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总是使用一个隐式寄存器，但为什么要使用两个额外的显式寄存器?

内核C更令人沮丧。10寄存器?但是只有两个指针。这为指针提供了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线是什么呢?它为什么在那里?

情况更糟:

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。与前两个内核相同的寄存器码。

内核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

相关内容

最新更新

热门标签：