关于寄存器的CUDA内核优化

我对CUDA和GPU编程很陌生 . 我正在尝试为物理学中的应用程序编写内核 . 并行化是在方向的正交上进行的，每个方向导致2D笛卡尔域的扫描 . 这是内核 . 它实际效果很好，效果很好 .

但是，每个块的寄存器数量非常多，导致本地内存溢出，严重降低了代码性能 .

__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

}

如何优化此代码的执行？我运行了8个32个线程的块 . 该内核的占用率非常低，受Visual Profiler的寄存器限制 .

我不知道如何改进它 .

谢谢 ！