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

发布于 2024-08-21 22:03:03 字数 221 浏览 7 评论 0原文

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

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

原文

分享到QQ

分享到微博

如果你对这篇内容有疑问，欢迎到本站社区发帖提问参与讨论，获取更多帮助，或者扫码二维码加入 Web 技术交流群。

发布评论

需要登录才能够评论，你可以免费注册一个本站的账号。

绮烟 2024-08-28 22:03:03

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

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

回复收藏 0 原文

愚人国度 2024-08-28 22:03:03

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

在您的情况下，您应该测量 maxrregcount 设置为 16 和不设置 maxrregcount 的性能。由于拥有足够的线程，本地内存的延迟应该被隐藏，假设您不随机访问本地数组（这将导致非-合并访问）。

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

回复收藏 0 原文

極樂鬼 2024-08-28 22:03:03

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

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

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

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

Utilizing shared memory as cache may lead less register usage and prevent register spilling to local memory...

Think that the kernel calculates some values and these calculated values are used by all of the threads,

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

So, instead of keeping reg and reg0 as registers and making them possibily spill out to local memory (global memory), we may use shared memory.

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

Take a look at this paper for further information..

回复收藏 0 原文