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个寄存器。那么为什么编译时显示只用了2个寄存器？
• 看到汇编后我意识到threadId、blockId的数据类型是。这是CUDA规范中定义的吗？还是可能因CUDA驱动版本不同而有所变化？
• 谁能解释一下这行代码：mul.wide.u16 %r1, %rh1, %rh2;？%r1是u32，为什么要使用wide而不是u32？
• 寄存器的名称是如何选择的？在我的情况下，我理解了%r，但我不理解h、(null)、d等部分。它是基于数据类型长度选择的吗？例如：16位使用h，32位使用(null)，64位使用d？
• 如果我将内核的最后两行替换为out[idx] = in[idx];，那么当我编译程序时，它说使用了3个寄存器！现在怎么可能使用更多的寄存器呢？

请忽略我的测试内核没有检查数组索引是否越界的事实。

非常感谢。

- AstrOne

---

(1) PTXAS是将PTX翻译成机器码的编译器组件。 因此，来自-Xptxas -v的寄存器计数涉及机器码中使用的物理寄存器（您可以使用cuobjdump --dump-sass进行检查）。 PTX是一种使用虚拟寄存器的中间语言。 由于PTX代码以SSA（单静态赋值）形式生成，因此每个新写入的结果都被分配一个新的虚拟寄存器号码。 (2) mul.wide在PTX规范中有描述（该规范是CUDA文档的一部分）。 在此情况下，它将两个u16操作数相乘，给出一个u32结果（即完整的乘积）。 - njuffa

Answer 1

PTX是一种中间语言，旨在跨多个GPU架构进行移植。它由编译器组件PTXAS编译成最终的机器代码，也称为特定架构的SASS。nvcc选项-Xptxas -v会导致PTXAS报告有关生成的机器代码的各种统计信息，包括在机器代码中使用的物理寄存器数量。您可以通过使用cuobjdump --dump-sass来反汇编机器代码。

因此，在PTX代码中使用的寄存器数目没有意义，因为这些是虚拟寄存器。CUDA编译器以所谓的SSA形式（静态单赋值，参见http://en.wikipedia.org/wiki/Static_single_assignment_form）生成PTX代码。这基本上意味着每个新写入的结果被分配一个新寄存器。

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位双精度浮点数。您可以通过查看创建这些虚拟寄存器的PTX代码中的 .reg 指令来轻松查看这一点。

- njuffa