4

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] 最终在寄存器中（在这种情况下，像我一样使用共享内存不会使感觉）。我想增加入住率，但我不明白为什么要使用这么多寄存器。有任何想法吗？

4

3 回答 3

9

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

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

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

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

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

于 2013-07-28T22:56:03.570 回答
1

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

于 2013-08-02T14:58:15.777 回答
1

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

于 2014-08-21T21:19:17.440 回答