8

:) 当我试图管理我的内核资源时，我决定研究 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对于 64 位？
  • 如果我用 this 替换内核的最后两行out[idx] = in[idx];，那么当我编译程序时，它会说使用了 3 个寄存器！现在怎么可能使用更多的寄存器？

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

非常感谢。

4

1 回答 1

12

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 代码中的指令，您可以很容易地看到这一点。

于 2013-06-07T05:24:50.623 回答