追踪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 - 答案肯定埋在其中 - 也许有人可以告诉我该寻找什么或向我展示如何处理汇编转储的指南。

- Marm0t

---

你的循环在 OL 值为 20 时被编译器展开了，但在 OL 值为 40 时没有展开，这可能是原因吗？ - Ashwin Nanjappa

我非常有信心，寄存器计数的差异与浮点数、循环展开或先前提到的任何其他因素都没有关系。请记住，sm_20内部是一个64位架构，而sm_13是一个32位架构。这意味着指针在为sm_20编译时具有比sm_12编译时多两倍的寄存器占用空间。 - talonmies

64位指针（或任何64位值）每个需要2个寄存器，因为寄存器是32位。但@talonmies，指针大小是否取决于是否指定了“-m32”或“-m64”？我记不清默认值是哪个；可能默认匹配当前操作系统。 - harrism

还有其他原因，但总的来说，sm_20需要比sm_1x更高的寄存器使用率。首先，sm_20是一种纯粹的加载存储体系结构，而sm_1x具有一些非加载存储指令，可以将内存（例如共享内存）用作操作数。为什么不在两个二进制文件上使用cuobjdump查看实际生成的机器代码-然后您就可以自己看到发生了什么。 - harrism

实际上，我记得 sm_20 设备寄存器计数上升的主要原因是没有地址寄存器，这也是我认为可能是这个问题的答案，所以我会将其发布为答案。 - harrism

Answer 1

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

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

- harrism

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个寄存器。

- Roger Dahl

---

嗯，编译器可能会将你的两个例子都视为第二个例子。 - harrism

编译器如何能够在第一个示例中仅使用一个寄存器？ - Roger Dahl

谢谢您的纠正。您知道那个额外的“r”代表什么吗？ - Roger Dahl

不确定，但我认为它可能是“实际的”（如浮点数）。这是为了区分其他类型的寄存器，例如地址寄存器，在sm_13之后的NVIDIA GPU上不存在。 - harrism