Mat*_*olo 2 cuda gpu hpc cpu-registers
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文档以及nvcc文档来查看各种编译器选项.
你没有提供你的代码,因此不可能提出任何直接的建议,但是你可以通过减少不断的使用,减少或重构算术使用,切换double到float,以及我确定有许多其他建议也是如此.如果将-G开关传递给编译器,寄存器的使用也会受到影响.
您可以通过将-maxrregcount开关传递给nvcc适当的参数来限制编译器对每个线程的寄存器的使用,例如,-maxrregcount 20这将指示编译器将自己限制为每个线程20个寄存器.但是,这种策略可能无法产生良好的效果,或者您可能需要将参数调整为不会牺牲太多性能的值.然而,您可能会找到一个最佳选择,它不会牺牲太多的基本性能,但可以让您提高入住率.如果您过多地限制编译器,它将开始将其所需的寄存器使用量溢出到本地内存,这通常会降低性能.
你也应该知道,你可以传递-Xptxas -v到nvcc,这将给有关在编译时编译器的寄存器使用和其他相关数据(溢出等)有用的输出.