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] 在寄存器中结束(其中
像我一样使用共享内存的情况没有意义)。我会
想增加入住率，但我不明白为什么要使用这么多寄存器。
有任何想法吗？
最佳答案
通常很难查看 C 代码并从中预测寄存器使用情况。编译器可能会通过增加寄存器使用量来积极优化代码，也许是为了在此处或此处保存指令。您似乎假设可以从 C 代码变量分配中预测寄存器使用情况，虽然两者之间存在某种联系，但您不能假设可以直接从 C 代码变量分配中计算寄存器使用情况。
由于您尚未提供代码，因此实际上没有人可以帮助使用寄存器。如果你想更好地理解寄存器的用法，你需要直接查看 PTX 代码。为此，请使用 nvcc 和 -ptx 开关编译您的代码，并直接检查生成的 .ptx 文件。为此，您可能希望引用 PTX documentation 和 nvcc documentation 以查看各种编译器选项。
您尚未提供代码，因此无法提出任何直接建议，但您可以通过减少常量使用、减少或重构算术使用、从 double 切换到 float 来减少寄存器使用，我敢肯定还有许多其他建议。如果您将 -G 开关传递给编译器，则寄存器使用也会受到影响。
您可以通过使用适当的参数将 -maxrregcount 开关传递给 nvcc 来限制编译器对每个线程寄存器的使用，例如 -maxrregcount 20 它将指示编译器将自身限制为每个线程 20 个寄存器。然而，这种策略可能不会产生好的结果，或者您可能需要将参数调整为一个不会牺牲太多性能的值。然而，您可能会找到一个最佳选择，它不会牺牲太多基本性能，但可以提高入住率。如果您过多地限制编译器，它将开始将其所需的寄存器使用溢出到本地内存，这通常会降低性能。
您还应该知道，您可以将 -Xptxas -v 传递给 nvcc，这将在编译时提供有关编译器寄存器使用情况和其他相关数据(溢出等)的有用输出。
关于cuda - 计算 Cuda 内核中的寄存器/线程，我们在Stack Overflow上找到一个类似的问题：https://stackoverflow.com/questions/17913630/