追踪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个寄存器(我相信一个浮点数[N]占用N + 1个寄存器 - 如果我是wronge,请纠正我)基于声明.

但编译时:

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