关于寄存器的CUDA内核优化

我对CUDA和GPU编程颇为陌生。我正在尝试为物理应用程序编写一个内核。平行化是在方向的正交上进行的，每个方向导致2D笛卡尔域的扫描。这是内核。它实际上运行良好，给出了很好的结果。关于寄存器的CUDA内核优化

然而，一个非常高的每块寄存器导致溢出到本地内存严厉减慢代码的性能。

__global__ void KERNEL (int imax, int jmax, int mmax, int lg, int lgmax, 
         double *x, double *y, double *qd, double *kappa, 
         double *S, double *G, double *qw, double *SkG, 
         double *Ska,double *a, double *Ljm, int *data) 

{ 
int m = 1+blockIdx.x*blockDim.x + threadIdx.x ; 
int tid = threadIdx.x ; 

//Var needed for thread execution 
... 

extern __shared__ double shared[] ; 

//Read some data from Global mem 
mu = qd[  (m-1)]; 
eta = qd[ MSIZE+(m-1)]; 
wm = qd[3*MSIZE+(m-1)]; 
amu = fabs(mu); 
aeta= fabs(eta); 
ista = data[  (m-1)] ; 
iend = data[1*MSIZE+(m-1)] ; 
istp = data[2*MSIZE+(m-1)] ; 
jsta = data[3*MSIZE+(m-1)] ; 
jend = data[4*MSIZE+(m-1)] ; 
jstp = data[5*MSIZE+(m-1)] ; 

j1 = (1-jstp) ; 
j2 = (1+jstp)/2 ; 
i1 = (1-istp) ; 
i2 = (1+istp)/2 ; 

isw = ista-istp ; 
jsw = jsta-jstp ; 

dy = dx = 1.0e-2 ; 

for(i=1 ; i<=imax; i++) Ljm[MSIZE*(i-1)+m] = S[jsw*(imax+2)+i] ; 

//Beginning of the vertical Sweep, can be from left to right, 
// or opposite depending on the thread 

for(j=jsta ; j1*jend + j2*j<=j2*jend + j1*j ; j=j+jstp) { 

Lw = S[j*(imax+2)+isw] ; 

//Beginning of the horizontal Sweep, can be from left to right, 
// or opposite depending on the thread 

    for(i=ista ; i1*iend + i2*i<=i2*iend + i1*i ; i=i+istp) { 

      ax = dy ; 
      Lx = ax*amu/ex ; 
      ay = dx ; 
      Ly = ay*aeta/ey ; 

      dv = ax*ay ; 
      L0 = dv*kappaij ; 
      Sp = S[j*(imax+2)+i]*dv ; 
      Ls = Ljm[MSIZE*(i-1)+m] ; 

      Lp = (Lx*Lw+Ly*Ls+Sp)/(Lx+Ly+L0) ; 

      Lw = Lw+(Lp-Lw)/ex ; 
      Ls = Ls+(Lp-Ls)/ey ; 

      Ljm[MSIZE*(i-1)+m] = Ls ; 

      shared[tid] = wm*Lp ; 
      __syncthreads(); 

      for (s=16; s>0; s>>=1) { 
       if (tid < s) { 
        shared[tid] += shared[tid + s] ; 
       } 
      } 

      if(tid==0) atomicAdd(&SkG[imax*(j-1)+(i-1)],shared[tid]*kappaij); 

    } 
    // End of horizontal sweep 
} 
// End of vertical sweep

}

我怎么能优化这个代码的执行？我运行它的32个线程的8块。该内核的占用率非常低，由Visual Profiler根据寄存器限制。

我对如何改进它不知道。

谢谢！

来源

2013-01-31 user2029298

不要太担心入住率（理论或实现）。它可能会误导为您获得哪种性能的指标。而是检查分析器中的值，如每个时钟指令（IPC）和每个请求事务。 –

首先，您使用的，因为这32个线程块，占用的内核太低。你的显卡运行在256并行线程，但它可以

多少个寄存器，你每使用多处理器运行到1536个线程（计算能力2.x的）？您也可以尝试将您的变量声明为其本地范围，从而帮助设备更好地重用寄存器。

来源

2013-01-31 15:51:49 jgonzac

对，我意识到每块32个线程太少。使用3D版本，我将运行32个线程的16个块。我正在使用63注册。编译时，我得到以下信息：56字节的堆栈帧，64字节的溢出存储，56字节的溢出加载。使用63个寄存器，256 + 0分钟 – user2029298

关于寄存器的CUDA内核优化

回答

相关问题