【问题标题】:CUDA register usageCUDA 寄存器使用
【发布时间】:2013-02-13 17:53:57
【问题描述】:

CUDA 手册指定每个多处理器的 32 位 寄存器的数量。是不是意味着:

  1. 双变量占用两个寄存器?

  2. 指针变量占用两个寄存器? - 6 GB 内存的 Fermi 上必须不止一个寄存器,对吧?

  3. 如果问题 2 的答案是肯定的,那么使用更少的指针变量和更多的int 索引一定会更好。

    E. g.,这个内核代码:

    float* p1;               // two regs
    float* p2 = p1 + 1000;   // two regs
    int i;                   // one reg
    for ( i = 0; i < n; i++ )
    {
        CODE THAT USES p1[i] and p2[i]
    }
    

    理论上需要比这个内核代码更多的寄存器:

    float* p1;               // two regs
    int i;                   // one reg
    int j;                   // one reg
    for ( i = 0, j = 1000; i < n; i++, j++ )
    {
        CODE THAT USES p1[i] and p1[j]
    }
    

【问题讨论】:

    标签: cuda gpu


    【解决方案1】:

    您的三个问题的简短答案是:

    1. 是的。
    2. 是的,如果代码是为 64 位主机操作系统编译的。设备指针大小始终与 CUDA 中的主机应用程序指针大小匹配。
    3. 没有。

    为了扩展第 3 点,请考虑以下两个简单的内存复制内核:

    __global__
    void debunk(float *in, float *out, int n)
    {
        int i = n * (threadIdx.x + blockIdx.x*blockDim.x);
    
        for(int j=0; j<n; j++) {
            out[i+j] = in[i+j];
        }
    }
    
    __global__
    void debunk2(float *in, float *out, int n)
    {
        int i = n * (threadIdx.x + blockIdx.x*blockDim.x);
        float *x = in + i;
        float *y = out + i;
    
        for(int j=0; j<n; j++, x++, y++) {
            *x = *y;
        }
    }
    

    根据您的估计,debunk 必须使用较少的寄存器,因为它只有两个局部整数变量,而 debunk2 有两个额外的指针。然而,当我使用 CUDA 5 发布工具链编译它们时:

    $ nvcc -m64 -arch=sm_20 -c -Xptxas="-v"  pointer_size.cu 
    ptxas info    : 0 bytes gmem
    ptxas info    : Compiling entry function '_Z6debunkPfS_i' for 'sm_20'
    ptxas info    : Function properties for _Z6debunkPfS_i
        0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
    ptxas info    : Used 8 registers, 52 bytes cmem[0]
    ptxas info    : Compiling entry function '_Z7debunk2PfS_i' for 'sm_20'
    ptxas info    : Function properties for _Z7debunk2PfS_i
        0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
    ptxas info    : Used 8 registers, 52 bytes cmem[0]
    

    它们编译为完全相同的寄存器计数。如果你反汇编工具链输出,你会发现除了设置代码之外,最终的指令流几乎是相同的。造成这种情况的原因有很多,但基本上可以归结为两个简单的规则:

    1. 试图从 C 代码(甚至 PTX 汇编程序)中确定寄存器计数大多是徒劳的
    2. 尝试对一个非常复杂的编译器和汇编器进行二次猜测也是徒劳的。

    【讨论】:

    • 能否解释一下,为什么 3. 是“否”?
    • 在简单的情况下,优化器是否有可能产生相同/相似的代码?实际上,我有一个有限差分代码,它处理 3D 数组并迭代 i0i1i2 索引。我通常需要离开当前点p[i],其中i = i0 + i1*stride1 + i2*stride2 在三个方向上。因此,如果我引入指针 px1 = p + 1py1 = p + stride1pz1 = p + stride2(可能还有更多 - px2 = p + 2 等),并与 p[i]px1[i] 等一起使用,代码会更清晰。如果这会增加寄存器使用率吗?优化器无法优化所有这些额外的指针?
    • 所以我想我的问题是:在复杂的内核代码中引入额外的指针而不是整数索引是否安全?也就是说,这是否有可能导致更高的寄存器使用率,大概是因为对于复杂的代码,优化器实际上会使用 2 寄存器(用于指针变量)而不是 1 寄存器(用于整数索引)生成二进制代码?跨度>
    • 我能建议您编译和反汇编演示内核并研究它们,直到您了解它们的作用。这个简单的操作基本上需要一个源地址和目标地址寄存器、一个用于值的数据寄存器(这里没有间接寻址)、一个用于循环行程计数的数据寄存器和一对用于评估循环条件的寄存器。那是8个寄存器。整数索引在这里无关紧要,并且不会神奇地保存寄存器:*p+i 和 p[i] 发出相同的代码。
    猜你喜欢
    • 2023-04-03
    • 1970-01-01
    • 1970-01-01
    • 2012-01-21
    • 2021-08-12
    • 1970-01-01
    • 2012-04-01
    • 2013-06-17
    • 1970-01-01
    相关资源
    最近更新 更多