我要处理的问题类似矩阵乘法。问题描述如下：内核是要计算B(i)*A(i)=C(i)
其中A B C都是矩阵（但是规模不是很大，i也不是很大）
A B C的存储结构为：A[n][m] B[p][n] C[p][m]
对于不同的i ，m n p 都是一样的。但是对于不同的i，B C不是满矩阵——第二维取值范围是 dp>=i&&dp<p(其他位置为空位）。
我写了一个内核，考虑了对齐、应用shared memory等问题，但是nvvp跑出来的结果显示global/DRAM load/store throughtput都比较低，不知道还能如何优化？
下面是我程序代码，以及一组可取的i m n p 值。
其中：
WFin同上面的A
DMt同上面的B
WFout同上面的C
NQA1同上面的 m，取128
NQA2同上面的 n ，取159
(J2max+1)同上面的p ,取159（J2max=158)
(Mmax+1)同上面的 i，取64(Mmax=63)
内核代码如下：
void __global__ kernel(int ldNQA1,int ldNQA2,int J2max,int Mmax,int NQA1,int NQA2,float *WFin,float *DMt, float *WFout){
__shared__ float tmpWFin[32][32];
__shared__ float tmpDMt[8][33];
int tidx,tidy,bidx,bidy,bidz;
int dblock,dNQA1,dNQA2,dj2,offset;
float tmp=0.0f;
tidx=threadIdx.x;
tidy=threadIdx.y;
bidx=blockIdx.x;
bidy=blockIdx.y;
bidz=blockIdx.z;
if(bidz<=Mmax&&J2max-bidy*blockDim.y>=bidz){
dNQA1=bidx*blockDim.x+tidx;
dj2=J2max-bidy*blockDim.y-tidy;
for(dblock=0;dblock<=ldNQA2-32;dblock+=32){
if(dNQA1<NQA1){
#pragma unroll 4
for(offset=tidy;offset<=tidy+24;offset+=8){
dNQA2=dblock+offset;
if(dNQA2<NQA2)
tmpWFin[offset][tidx]=WFin[(bidz*NQA2+dNQA2)*ldNQA1+dNQA1];
}
}
if(dj2>=bidz&&dj2<=J2max){
dNQA2=dblock+tidx;
if(dNQA2<NQA2)
tmpDMt[tidy][tidx]=DMt[(bidz*(J2max+1)+dj2)*ldNQA2+dNQA2];
}
__syncthreads();
if(dNQA1<NQA1&&dj2>=bidz&&dj2<=J2max){
#pragma unroll 32
for(offset=0;offset<32;offset++){
if(dblock+offset<NQA2)
tmp+=tmpWFin[offset][tidx]*tmpDMt[tidy][offset];
}
}
__syncthreads();
}
if(dNQA1<NQA1&&dj2>=bidz&&dj2<=J2max)
WFout[(bidz*(J2max+1)+dj2)*ldNQA1+dNQA1]=tmp;
}
}
调用函数代码如下：
int main(){
int ldNQA1=((64-1)/32+1)*32*2;
int ldNQA2=((159-1)/32+1)*32;
int J2max=158;
int NQA1=64*2;
int NQA2=159;
int Mmax=63;
float *dWFin,*dDMt,*dWFout;
float a;
cudaMalloc((void **)&dWFin,sizeof(float)*ldNQA1*NQA2*(Mmax+1));
cudaMalloc((void **)&dDMt,sizeof(float)*ldNQA2*(J2max+1)*(Mmax+1));
cudaMalloc((void **)&dWFout,sizeof(float)*ldNQA1*NQA2*(Mmax+1));
dim3 block,grid;
grid=dim3((NQA1-1)/32+1,J2max/8+1,Mmax+1);
block=dim3(32,8,1);
kernel<<<grid,block>>>(ldNQA1,ldNQA2,J2max,Mmax,NQA1,NQA2,dWFin,dDMt,dWFout);
cudaDeviceSynchronize();
printf("%s\n",cudaGetErrorString(cudaGetLastError()));
cudaMemcpy(dWFout,&a,sizeof(float),cudaMemcpyDeviceToHost);
printf("%f\n",a);
return 0;
}
------------------------------------------------------------------------------更改了下字体格式