与CUDA PTX代码和寄存器内存混淆

Question

:)当我试图管理我的内核资源时,我决定调查PTX,但有一些我不明白的事情.这是我写的一个非常简单的内核:

__global__
void foo(float* out, float* in, uint32_t n)
{
    uint32_t idx = blockIdx.x * blockDim.x + threadIdx.x;
    uint32_t one = 5;
    out[idx] = in[idx]+one;
}

然后我用它编译它:nvcc --ptxas-options=-v -keep main.cu我在控制台上得到了这个输出:

ptxas info    : 0 bytes gmem
ptxas info    : Compiling entry function '_Z3fooPfS_j' for 'sm_10'
ptxas info    : Used 2 registers, 36 bytes smem

最终的ptx如下:

    .entry _Z3fooPfS_j (
            .param .u64 __cudaparm__Z3fooPfS_j_out,
            .param .u64 __cudaparm__Z3fooPfS_j_in,
            .param .u32 __cudaparm__Z3fooPfS_j_n)
    {
    .reg .u16 %rh<4>;
    .reg .u32 %r<5>;
    .reg .u64 %rd<8>;
    .reg .f32 %f<5>;
    .loc    15  17  0
$LDWbegin__Z3fooPfS_j:
    .loc    15  21  0
    mov.u16     %rh1, %ctaid.x;
    mov.u16     %rh2, %ntid.x;
    mul.wide.u16    %r1, %rh1, %rh2;
    cvt.u32.u16     %r2, %tid.x;
    add.u32     %r3, %r2, %r1;
    cvt.u64.u32     %rd1, %r3;
    mul.wide.u32    %rd2, %r3, 4;
    ld.param.u64    %rd3, [__cudaparm__Z3fooPfS_j_in];
    add.u64     %rd4, %rd3, %rd2;
    ld.global.f32   %f1, [%rd4+0];
    mov.f32     %f2, 0f40a00000;        // 5
    add.f32     %f3, %f1, %f2;
    ld.param.u64    %rd5, [__cudaparm__Z3fooPfS_j_out];
    add.u64     %rd6, %rd5, %rd2;
    st.global.f32   [%rd6+0], %f3;
    .loc    15  22  0
    exit;
$LDWend__Z3fooPfS_j:
    } // _Z3fooPfS_j

现在有一些我不明白的事情:

• 根据ptx汇编,使用4 + 5 + 8 + 5 = 22个寄存器.那为什么它used 2 registers在编译期间说？
• 看着程序集我意识到threadId,blockId等的数据类型是u16.这是在CUDA规范中定义的吗？或者这可能因CUDA驱动程序的不同版本而异？
• 能否给我有人解释这一行:mul.wide.u16 %r1, %rh1, %rh2;？%r1是u32,为什么wide不u32使用？
• 如何选择寄存器的名称？在我的花瓶里,我理解这%r部分,但我不理解h,(null),d部分.是根据数据类型长度选择的吗？即:h对于16位,对于32位d为null,对于64位？
• 如果我用这个替换我的内核的最后两行out[idx] = in[idx];,那么当我编译程序时它会说3个寄存器被使用!现在如何使用更多寄存器？

请忽略我的测试内核不检查数组索引是否超出范围的事实.

非常感谢你.

Answer 1

PTX是一种中间语言,旨在跨多个GPU架构移植.它由编译器组件PTXAS编译成最终的机器代码,也称为SASS,用于特定的体系结构.nvcc选项-Xptxas -v使PTXAS报告有关生成的机器代码的各种统计信息,包括机器代码中使用的物理寄存器的数量.您可以通过拆卸来检查机器代码cuobjdump --dump-sass.

因此,PTX代码中使用的寄存器数量没有意义,因为它们是虚拟寄存器.CUDA编译器以所谓的SSA形式生成PTX代码(静态单一赋值,请参阅http://en.wikipedia.org/wiki/Static_single_assignment_form).这基本上意味着为每个写入的新结果分配一个新的寄存器.

该指令mul.wide在PTX规范中描述,其当前版本(3.1)可在此处找到:http://docs.nvidia.com/cuda/parallel-thread-execution/index.html.在您的示例代码中,后缀.u16表示它将两个无符号的16位数相乘并返回无符号的32位结果,即它计算源操作数的完整双倍宽度乘积.

PTX中的虚拟寄存器是键入的,但它们的名称可以自由选择,与类型无关.CUDA编译器似乎遵循(据我所知)未记录的某些约定,因为它们是内部实现工件.查看一堆PTX代码,很明显当前生成的寄存器名称是编码类型信息,这可能是为了便于调试:p<num>用于谓词,r<num>32位整数,rd<num>64位整数,f<num>32位浮点数,以及fd<num>64位双打.您可以通过查看.reg创建这些虚拟寄存器的PTX代码中的指令轻松地看到这一点.