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

Question

我有一个使用了17个寄存器的内核，如果将其减少到16个，那么100%的占用率就可以实现。我的问题是：除了完全重写我的算法之外，是否有方法可以减少使用的寄存器数量？我一直认为编译器比我聪明得多，所以例如仅出于清晰起见，我经常使用额外的变量。这种想法是错误的吗？

请注意：我知道--max_registers（或其语法）标志，但是使用本地内存将比降低25％的占用率更具有破坏性（我应该测试一下）。

- zenna

---

很奇怪，我刚刚尝试了maxrregcount=16，实际上它减少了我使用的寄存器数量到15个，并且没有使用本地内存。但是它实际上变慢了！这是怎么回事？ - zenna

尝试对你的应用程序进行分析。编译器可能会引入一些干扰。 - Anycorn

占用率比我预测的高了15个寄存器，除了寄存器数量减少导致指令数量增加，其他都一样。从3.9M到4.3M。 - zenna

运行时分支/分歧的数量如何？ - Anycorn

只晚了8年才加入这个派对，但它变慢的原因可能是编译器开始进行再物化。它避免使用更多的寄存器，而是每次需要时都重新计算值。 - Chris Kitching

Answer 1

占用率可能有点误导性，100％的占用率不应该是您的主要目标。如果您可以获得完全合并到全局内存的访问，则在高端GPU上，50％的占用率足以隐藏对全局内存的延迟（对于浮点数来说，双精度甚至更低）。请查看去年GTC的Advanced CUDA C演示文稿以获取有关此主题的更多信息。

在您的情况下，您应该同时使用和不使用maxrregcount设置为16来测量性能。假设您没有随机访问本地数组（这将导致非合并访问），则本地内存的延迟应该被隐藏。

要回答您关于减少寄存器的具体问题，请发帖子以获取更详细的答案！了解编译器的工作原理可能有所帮助，但请记住，nvcc是一个具有大量参数空间的优化编译器，因此最小化寄存器计数必须与整体性能平衡。

- Tom

---

50% 的占用率会足够吗？你能否详细解释一下呢？非常感谢。 - ZeroCool

Answer 2

很难说，我认为nvcc编译器并不是很智能。

你可以尝试一些显而易见的方法，例如使用short代替int，通过引用传递和使用变量（例如&variable），展开循环，使用模板（如C++）。如果有除法、超越函数被连续应用，请尝试将它们作为一个循环。尽可能摆脱条件语句，可能会用冗余计算来代替它们。

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

- Anycorn

---

由于寄存器是32位的，GPU上的int也是32位的，那么int和short有什么区别吗？ - personne3000

Answer 3

利用共享内存作为缓存可能会减少寄存器使用量并防止寄存器溢出到本地内存...

想象一下内核计算一些值，而这些计算出的值被所有线程使用，

__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作为寄存器并可能将它们溢出到本地内存（全局内存）。

__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;
}

请查看这篇论文了解更多信息。

- phoad

---

每个独立的块需要自己的缓存区，并且每个块的第一个线程应该填充它。因此，每个块都是独立的，不需要同步。if语句之后的__syncthreads同步了块中的线程。尽管如此，串行部分以这种方式增加可能不是一个好的解决方案。 - phoad

已经 threadidx.x=6 不会计算任何东西。它将从缓存中获取计算结果，而缓存将具有计算结果作为同步点通过。不是吗？ - phoad

你是指最后两行吗？从缓存中读取的吗？有没有办法修复它，比如使用thread_fence等？ - phoad

Answer 4

降低寄存器使用量时指令计数会增加，这个现象的原因很简单。编译器可能会使用寄存器来存储某些操作的结果，在您的代码中多次使用，以避免重新计算这些值。当强制使用较少的寄存器时，编译器决定重新计算那些本应存储在寄存器中的值。

- Victor M. Alvarez

Answer 5

通常来说，减少寄存器压力并不是一个好的方法。编译器在优化整个内核性能时表现出色，并且考虑了许多因素，包括寄存器。

当减少寄存器导致速度变慢时，很可能是编译器不得不将不足的寄存器数据溢出到“本地”内存中，这本质上与全局内存相同，因此非常缓慢。

为了优化，建议在必要时使用关键字如const、volatile等，以帮助编译器进行优化阶段。

总之，通常不是像寄存器这样微小的问题使CUDA内核运行缓慢。我建议优化与全局内存的工作、访问模式、如果可能的话，在纹理内存中缓存，通过PCIe进行事务处理。

- Tosha