0

我对 CUDA 和 GPU 编程很陌生。我正在尝试为物理应用程序编写内核。并行化是在正交方向上进行的，每个方向都会扫描二维笛卡尔域。这里是内核。它实际上效果很好，效果很好。

但是，每个块的大量寄存器会导致溢出到本地内存，从而严重降低代码性能。

__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 受到寄存器的限制。

我不知道如何改进它。

谢谢 ！

4

1 回答 1

1

首先，您使用的是 32 个线程块，因此，占用内核太低。您的 gpu 仅并行运行 256 个线程，但每个多处理器最多可运行 1536 个线程（计算能力 2.x）

你用了多少个寄存器？您还可以尝试将变量声明到它们的本地范围内，帮助设备更好地重用寄存器。

于 2013-01-31T15:51:49.097 回答