11

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

请注意：我确实知道 --max_registers （或任何语法）标志，但使用本地内存比降低 25% 的占用率更有害（我应该对此进行测试）

4

5 回答 5

8

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

在您的情况下，您应该在 maxrregcount 设置为 16 和不设置为 16 的情况下测量性能。假设您不随机访问本地数组（这将导致非合并访问）。

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

于 2010-02-18T10:37:13.630 回答
6

真的很难说，在我看来，nvcc 编译器不是很聪明。
您可以尝试一些显而易见的事情，例如使用 short 而不是 int、通过引用传递和使用变量（例如&变量）、展开循环、使用模板（如在 C++ 中）。如果你有除法，超越函数，按顺序应用，试着把它们做成一个循环。尝试摆脱条件，可能用冗余计算代替它们。

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

于 2010-02-17T19:12:39.123 回答
4

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

认为内核计算了一些值，并且这些计算值被所有线程使用，

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

看看这篇论文以获取更多信息。

于 2013-01-22T13:17:05.613 回答
2

这通常不是最小化套准压力的好方法。编译器在优化整体预计内核性能方面做得很好，它考虑了很多因素，包括寄存器。

当减少寄存器导致速度变慢时它是如何工作的

很可能编译器不得不将不足的寄存器数据溢出到“本地”内存中，这与全局内存基本相同，因此非常慢

出于优化目的，我建议在必要时使用 const、volatile 等关键字，以帮助编译器进行优化阶段。

无论如何，并不是这些像寄存器这样的小问题经常使 CUDA 内核运行缓慢。我建议优化全局内存、访问模式、纹理内存中的缓存（如果可能）以及通过 PCIe 的事务。

于 2010-08-05T11:57:37.550 回答
1

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

于 2010-02-21T13:41:14.233 回答