计算 Cuda 内核中的寄存器/线程答案

【问题标题】：Counting registers/thread in Cuda kernel计算 Cuda 内核中的寄存器/线程
【发布时间】：2013-07-28 22:32:49
【问题描述】：

nSight 分析器告诉我，以下内核每个线程使用 52 个寄存器：

 //Just the first lines of the kernel.
__global__ void voles_kernel(float *params, int *ctrl_params, 
                             float dt, float currTime,
                             float *dev_voles, float *dev_weasels, 
                             curandStateMtgp32 *state) 
{

 __shared__ float dev_params[9];

 __shared__ int BuYeSimStep[4];

 if(threadIdx.x < 4)
 {
   BuYeSimStep[threadIdx.x] = ctrl_params[threadIdx.x];
 }

 if(threadIdx.x < 9){
     dev_params[threadIdx.x] = params[threadIdx.x];
 }

 __syncthreads();

float currVole = curand_uniform(&state[blockIdx.x]) + 3.0;
float currWeas = curand_uniform(&state[blockIdx.x]) + 0.1;
float oldVole = currVole;
float oldWeas = currWeas;

int jj;

if (blockIdx.x * blockDim.x + threadIdx.x < BuYeSimStep[2])
{
int dayIndex = 0;

/* Not declaring any new variable from here on, just doing arithmetics.
   ....... */

如果每个寄存器有 4 个字节，我不明白我们如何得到 52 个寄存器，即使假设数组 params[9] 和 ctrl_params[4] 最终在寄存器中（其中像我一样使用共享内存的情况没有意义）。我会喜欢增加入住率，但我不明白为什么要使用这么多寄存器。有什么想法吗？

【问题讨论】：

标签： cuda gpu cpu-registers hpc

【解决方案1】：

通常很难查看 C 代码并从中预测寄存器的使用情况。编译器可能会通过增加寄存器使用来积极优化代码，也许是为了在这里或那里保存一条指令。您似乎在假设可以从 C 代码变量分配中预测寄存器使用情况，虽然两者之间存在某种联系，但您不能假设可以直接从 C 代码变量分配中计算寄存器使用情况。

由于您尚未提供代码，因此没有人可以真正帮助您使用寄存器。如果您想更好地了解寄存器的用法，则需要直接查看 PTX 代码。为此，请使用 nvcc 和 -ptx 开关编译您的代码，并直接检查生成的 .ptx 文件。为此，您可能希望参考PTX documentation 和nvcc documentation 来查看各种编译器选项。

您还没有提供您的代码，因此实际上不可能提出任何直接建议，但您可以通过减少常量使用、减少或重构算术使用、从 double 切换到 @987654327 来减少寄存器使用@，我相信还有很多其他的建议。如果您将-G 开关传递给编译器，寄存器的使用也会受到影响。

您可以通过将-maxrregcount 开关传递给带有适当参数的nvcc 来限制编译器对每个线程的寄存器使用，例如-maxrregcount 20，它将指示编译器将自身限制为每个线程20 个寄存器。但是，这种策略可能不会产生好的结果，或者您可能需要将参数调整为不会牺牲太多性能的值。但是，您可能会找到一个不会牺牲太多基本性能但可以提高入住率的最佳选择。如果您对编译器进行过多限制，它将开始将所需的寄存器使用溢出到本地内存，这通常会降低性能。

您还应该知道，您可以将-Xptxas -v 传递给nvcc，这将在编译时提供有关编译器寄存器使用情况和其他相关数据（溢出等）的有用输出。

【讨论】：

-maxregcount 不是推荐的方法，因为它为编译单元中的所有函数设置值。您应该使用 ____launch_bounds____ 属性。除非您使用 PTX 写作，否则查看 PTX 往往是无用的。您可以使用 cuobjdump、nvdisasm 或开发工具查看 SASS 代码。可以找到 grep r[0-9]+ 的总寄存器计数并找到最高匹配。如果最高寄存器用于 128 (vec4) 操作，则总寄存器计数最多将高 3。
每个线程的运行时寄存器分配可能高于编译器报告的值，原因是 (a) ABI 合规性，或 (b) 由于每个线程的寄存器分配粒度。每个线程的寄存器分配粒度在开普勒设备上为 8，在 Fermi 上为 2 或 4。
非常感谢您的回答！

【解决方案2】：

如果要增加占用，直接的方法是使用编译器标志：maxregcount 来限制寄存器的使用，但是可能会因为一些寄存器会溢出到本地内存而导致性能损失，这很慢。

【讨论】：

【解决方案3】：

我建议您使用 Eclipse Nsight 调试您的代码。在内核的第一行创建一个断点并单步执行。在 Debug Perspective 中，在 CUDA 线程内，您拥有当前堆栈跟踪。右键单击堆栈，然后单击“指令步进模式”。 “反汇编”窗口将打开您的内核 PTX 程序集。您可以继续单步执行内核以跟踪源代码和程序集的相关性。所以你可以发现使用哪个寄存器。

【讨论】：