减少CUDA内核中使用的寄存器的数量

Question

我有一个使用17个寄存器的内核，将其减少到16会带给我100％的占用率。我的问题是：是否有方法可用于减少使用的数量或寄存器，不包括以不同方式完全重写我的算法。我总是认为编译器比我聪明得多，所以为了清晰起见，我经常使用额外的变量。这个想法我错了吗？

请注意：我不知道有关--max_registers（或任何语法）标志，但使用的本地内存会比降低了25％的入住更不利（我应该测试这个）

Answer 1

真的很难说，在我看来，nvcc编译器不是很聪明。
你可以尝试一些明显的事情，例如使用short而不是int，通过引用传递和使用变量（例如&变量），展开循环，使用模板（如在C++中）。如果你有分裂，先验功能，顺序应用，尝试使它们成为一个循环。尽量摆脱条件，可能用冗余计算替代它们。

如果你发布了一些代码，也许你会得到具体的答案。

Answer 2

入住率可能有点误导，100％的入住率不应该是您的主要目标。如果可以完全合并访问全局内存，那么在高端GPU上占用50％就足以隐藏全局内存的延迟（对于浮点数，甚至更低）。查看去年GTC的Advanced CUDA C演示文稿，了解更多关于此主题的信息。

对于你的情况，你应该测量有无maxrregcount设置为16的性能。假设你没有随机访问本地数组，那么本地内存的延迟应该隐藏起来，因为它有足够的线程。导致非合并访问）。

要回答您关于减少寄存器的具体问题，请发布代码以获取更详细的答案！了解编译器如何在一般情况下工作可能会有帮助，但请记住，nvcc是一个具有大参数空间的优化编译器，因此最大限度地减少寄存器数量必须与总体性能保持平衡。

Answer 3

降低寄存器使用时的指令数增加有一个简单的解释。编译器可以使用寄存器来存储通过代码多次使用的一些操作的结果，以避免重新计算这些值，当被迫使用较少的寄存器时，编译器决定重新计算将存储在寄存器中的那些值除此以外。

Answer 4

这通常不是一种好方法来最小化记录压力。编译器在优化整体计划的内核性能方面做得很好，并且考虑了很多因素，包括注册表。

它是如何工作的时候降低寄存器造成速度较慢

最有可能的编译器必须足够的寄存器中的数据溢出到“本地”的内存，这是基本相同的全局内存，因而非常缓慢

为了优化目的，我会建议在必要时使用像const，volatile等关键字来帮助编译器优化阶段。

无论如何，这不是像寄存器这些微小的问题，经常使CUDA内核运行缓慢。我建议优化全局内存，访问模式，尽可能在纹理内存中缓存，通过PCIe进行交易。

Answer 5

利用共享内存作为高速缓存可能会导致更少的注册使用和防止溢出的寄存器对本地内存...

想想内核计算一些数值与这些计算值是由所有线程的使用，

__global__ void kernel(...) { 
    int idx = threadIdx.x + blockDim.x * blockIdx.x; 
    int id0 = blockDim.x * blockIdx.x; 

    int reg = id0 * ...; 
    int reg0 = reg * a/x + y; 


    ... 

    int val = reg + reg0 + 2 * idx; 

    output[idx] = val > 10; 
} 

所以，与其保持REG和REG0寄存器和使他们possibily溢出到本地内存（全局内存），我们可能会使用共享内存。

__global__ void kernel(...) { 
    __shared__ int cache[10]; 

    int idx = threadIdx.x + blockDim.x * blockIdx.x; 

    if (threadIdx.x == 0) { 
     int id0 = blockDim.x * blockIdx.x; 

     cache[0] = id0 * ...; 
     cache[1] = cache[0] * a/x + y; 
    } 
    __syncthreads(); 


    ... 

    int val = cache[0] + cache[1] + 2 * idx; 

    output[idx] = val > 10; 
} 

看看这个paper了解更多信息..