跟踪 cuda 内核寄存器的使用情况

Question

我正在尝试追踪寄存器的使用情况，并遇到了一个有趣的场景。考虑以下来源：

#define OL 20
#define NHS 10

__global__ void loop_test( float ** out, const float ** in,int3 gdims,int stride){

        const int idx = blockIdx.x*blockDim.x + threadIdx.x;
        const int idy = blockIdx.y*blockDim.y + threadIdx.y;
        const int idz = blockIdx.z*blockDim.z + threadIdx.z;

        const int index = stride*gdims.y*idz + idy*stride + idx;
        int i = 0,j =0;
        float sum =0.f;
        float tmp;
        float lf;
        float u2, tW;

        u2 = 1.0;
        tW = 2.0;

        float herm[NHS];

        for(j=0; j < OL; ++j){
                for(i = 0; i < NHS; ++i){
                        herm[i] += in[j][index];
                }
        }

        for(j=0; j<OL; ++j){
                for(i=0;i<NHS; ++i){
                        tmp = sum + herm[i]*in[j][index];
                        sum = tmp;
                }
                out[j][index] = sum;
                sum =0.f;
        }

}

作为对源代码的附注 - 我可以做 += 的运行总和，但正在研究如何改变寄存器使用的影响（似乎没有 - 只是添加了一个额外的 mov 指令）。 此外，此源面向访问映射到 3D 空间的内存。

根据声明计算寄存器，似乎有 22 个寄存器（我相信 float[N] 占用 N+1 个寄存器 - 如果我错了，请纠正我）。

但是编译：

nvcc -cubin -arch=sm_20 -Xptxas="-v" src/looptest.cu

产量：

0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 25 registers, 72 bytes cmem[0]

好的，所以数字与“预期”不同。此外，如果使用 :

编译

nvcc -cubin -arch=sm_13 -Xptxas="-v" src/looptest.cu

寄存器使用量远少 - 准确地说是 8 个（显然是因为 sm_20 比 sm_13 更符合 IEEE 浮点数学标准？）：

ptxas info    : Compiling entry function '_Z9loop_testPPfPPKfS2_4int3i' for 'sm_13'
ptxas info    : Used 17 registers, 40+16 bytes smem, 8 bytes cmem[1]

最后，将宏 OL 更改为 40，然后突然：

0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 28 registers, 72 bytes cmem[0]

总之，我想知道寄存器在哪里被吃掉了，以及我所做的几个观察结果是什么。

我没有足够的组装经验来通过 cuobjdump - 答案肯定隐藏在那里 - 也许有人可以告诉我我应该寻找什么或向我展示如何处理组装的指南转储。

Answer 1

sm_20 和 sm_13 是非常不同的架构，具有非常不同的指令集 (ISA) 设计。您看到的导致寄存器使用量增加的主要区别是 sm_1x 具有专用地址寄存器，而 sm_2x 及更高版本则没有。相反，地址和值一样存储在通用寄存器中，这意味着大多数程序在 sm_2x 上比在 sm_1x 上需要更多的寄存器。

sm_20 的寄存器文件大小也是 sm_13 的两倍，以补偿这种影响。

Answer 2

寄存器的使用不一定与变量的数量密切相关。

编译器尝试通过比较单个内核中的潜在收益与所有同时运行的内核由于可用寄存器较少而导致的成本，来评估在代码中的两个使用点之间将变量保存在寄存器中的速度优势在注册池中。 （一个 Fermi SM 有 32768 个寄存器）。因此，如果更改代码会导致使用的寄存器数量出现意外波动，这并不奇怪。

如果分析器说您的入住率受到寄存器使用量的限制，您真的应该担心寄存器使用情况。在这种情况下，您可以使用--maxrregcount 设置来减少单个内核使用的寄存器数量，看看它是否会提高整体执行速度。

为了帮助减少内核使用的寄存器数量，您可以尝试将变量的使用尽可能保持在本地。例如，如果你这样做：

set variable 1
set variable 2
use variable 1
use variable 2

这可能会导致使用 2 个寄存器。同时，如果您：

set variable 1
use variable 1
set variable 2
use variable 2

这可能会导致使用 1 个寄存器。