程序优化

程序写出来了,但是发现优化基本上不怎么会啊,问题一大堆。。先下了一个图像相关的程序。代码如下:其中:img1、img2是两幅图像,Height、Width是图像的高和宽、SubMatrix是模板中心、CorCoef_d用于存储相关系数,M是计算模板大小,m_nRow、m_nCol是需要匹配的个数。

__global__ void CorrelationGPU(float *Img1 ,float *Img2,int Height,int Width,int *SubMatrix,float *CorCoef_d,int M,int m_nRow,int m_nCol)
{
   int bky=blockIdx.y;         
   int bkx=blockIdx.x;
   int tdy=threadIdx.y; 
   int tdx=threadIdx.x;
   int i,j,tempv;
   int y=SubMatrix[bky*m_nRow+bkx+m_nRow*m_nCol]; //原图block的中心点
   int x=SubMatrix[bky*m_nRow+bkx];        
   __shared__        float x1,x2,x3;
   __shared__ float shared61[5120];        //大小是(M*M)*5,存储两个模板的值
   int M2=M*M; 
   M=M/2;
   i=tdy-M;
   j=tdx-M; //以块的中心为坐标

   //////////////////////////////////////////////////////////////////////////
   //求以(x,y)为中心的,M*M为模板的的相关系数                                                        
   shared61[0*M2+tdy*blockDim.x+tdx]=Img1[(y+i)*Width+x+j];  //f
   shared61[1*M2+tdy*blockDim.x+tdx]=Img2[(y+i)*Width+x+j];  //g

   shared61[2*M2+tdy*blockDim.x+tdx]=shared61[0*M2+tdy*blockDim.x+tdx]; 
   shared61[3*M2+tdy*blockDim.x+tdx]=shared61[1*M2+tdy*blockDim.x+tdx];
   __syncthreads();

   for (tempv=M2/2;tempv>0;tempv=tempv>>1){
   if (tdy*blockDim.x+tdx<tempv){
   shared61[2*M2+tdy*blockDim.x+tdx]+=shared61[2*M2+tdy*blockDim.x+tdx+tempv];
   shared61[3*M2+tdy*blockDim.x+tdx]+=shared61[3*M2+tdy*blockDim.x+tdx+tempv];
   }
   __syncthreads();
   }
   if (tdy*blockDim.x+tdx==0){ 
   x1=shared61[2*M2+0]; 
   x2=shared61[3*M2+0];
   x1=x1/(M2);
   x2=x2/(M2);
   }
   __syncthreads();

   //用规约求相关系数cor
   shared61[2*M2+tdy*blockDim.x+tdx]=(shared61[0*M2+tdy*blockDim.x+tdx]-x1)*(shared61[0*M2+tdy*blockDim.x+tdx]-x1); //数据,这里传给0,1,2是为了不让规约运算改变原始值
   shared61[3*M2+tdy*blockDim.x+tdx]=(shared61[1*M2+tdy*blockDim.x+tdx]-x2)*(shared61[1*M2+tdy*blockDim.x+tdx]-x2);
   shared61[4*M2+tdy*blockDim.x+tdx]=(shared61[0*M2+tdy*blockDim.x+tdx]-x1)*(shared61[1*M2+tdy*blockDim.x+tdx]-x2); 
   __syncthreads();

   for (tempv=M2/2;tempv>0;tempv=tempv>>1){//规约,得到fsum,gsum,fgsum
   if (tdy*blockDim.x+tdx<tempv){
   shared61[2*M2+tdy*blockDim.x+tdx]+=shared61[2*M2+tdy*blockDim.x+tdx+tempv];
   shared61[3*M2+tdy*blockDim.x+tdx]+=shared61[3*M2+tdy*blockDim.x+tdx+tempv];
   shared61[4*M2+tdy*blockDim.x+tdx]+=shared61[4*M2+tdy*blockDim.x+tdx+tempv];
   }
   __syncthreads();
   }


   if (tdy*blockDim.x+tdx==0){ 
   x1=shared61[2*M2+0];  //x1~fsum/cor,x2~gsum,x3~fgsum
   x2=shared61[3*M2+0];
   x3=shared61[4*M2+0];
   x1=x3/sqrt(x1*x2);
   CorCoef_d[bky*m_nRow+bkx]=x1;
   }
   __syncthreads();
}

其中在CPU上的程序如下,其中bk11为<<40,40>>,td11为<<32,32>>,GPU为GTX 650 TI

void FuncionCorrelation(std::vector<CDib8*> m_dib,float* subMatrix,int m_nWidth,int m_nHeight,int m_nRow,int m_nCol,int M)
{
   int i,j,tempij,linebyte; 
   float *m_prData1,*m_prDatak;

   int HWMemB=m_nHeight*m_nWidth*sizeof(float);
   int RCMemBi=m_nRow*m_nCol*sizeof(int);
   int RCMemBf=m_nRow*m_nCol*sizeof(float);

   m_prData1=(float*)malloc(HWMemB);        //第1、2张图(m_prDatak中k是2)
   m_prDatak=(float*)malloc(HWMemB);
   linebyte=(m_nWidth+3)/4*4; //图像的宽,图像数据

   for (i=0;i<m_nHeight;i++){
   tempij=i*linebyte;
   for (j=0;j<m_nWidth;j++){
   m_prData1[i*m_nWidth+j]=float(m_dib[0]->m_pImgData[tempij+j]);
   m_prDatak[i*m_nWidth+j]=float(m_dib[1]->m_pImgData[tempij+j]);
   }
   }

   int *subMatrix_gpu;
   checkCudaErrors(cudaMalloc((void**)&subMatrix_gpu,2*RCMemBi));        //GPU~~~~subMatrix_gpu
   checkCudaErrors(cudaMemcpy(subMatrix_gpu,subMatrix,2*RCMemBi,cudaMemcpyHostToDevice));

   float *m_prData1_gpu,*m_prDatak_gpu; //GPU
   checkCudaErrors(cudaMalloc((void**)&m_prData1_gpu,HWMemB));         //GPU~~~~f
   checkCudaErrors(cudaMalloc((void**)&m_prDatak_gpu,HWMemB));         //GPU~~~~g
   checkCudaErrors(cudaMemcpy(m_prData1_gpu,m_prData1,HWMemB,cudaMemcpyHostToDevice));
   checkCudaErrors(cudaMemcpy(m_prDatak_gpu,m_prDatak,HWMemB,cudaMemcpyHostToDevice));

   dim3 bk11(m_nRow,m_nCol),td11(M,M);
   float *CorCoef_d,*CorCoef_h;
   checkCudaErrors(cudaMalloc((void**)&CorCoef_d,m_nRow*m_nCol*sizeof(float)));
   CorCoef_h=(float*)malloc(m_nRow*m_nCol*sizeof(float));
   CorrelationGPU<<<bk11,td11>>> (m_prData1_gpu,m_prDatak_gpu,m_nHeight,m_nWidth,subMatrix_gpu,CorCoef_d,M,m_nRow,m_nCol);

   cudaDeviceReset();
   ExitProcess(0); 
}

用profiler得到的结果截图如下:

[attach]3400[/attach]
根据提示我认为第一个是有SM是限制的,第二个是数据传输和核函数的执行没有异步执行。我的问题如下:
1):第一个提示是不是应该属于最大化利用率中最大化设备内多处理器的并行执行?这个是怎么导致的呢?怎么改才能避免这个情况?还有在编程指南中在多处理器层次中提到时间的延迟掩盖,这个基本上都没有看懂,不知道版主能不能把这块讲一下,如果太多不方便,能不能推荐一下比较浅显的资料啊?谢谢
2):第二个好像比较明显,就是异步操作没有弄好,但是我感觉好像程序也没法进行异步操作啊,这个本来就应该是这个顺序来的,不知道我这么理解对不对,或者有没有什么好的优化方法。。
3):不知道还有没有其他问题没有想到的,希望版主多多提示。。谢谢

LZ您好:

1:根据您的profiler的截图和提示,第一个问题是您的计算时间再总时间中所占比例甚小。您的计算规模太小以及cudaMalloc时初始化环境有一定时间消耗。

这个和SM层面多个线程的互相掩盖无关,一般地您可以尽量多地安排线程,以方便硬件调度来实现计算和访存的掩盖。详细讨论这个问题超出了本帖的范围,不再展开。

2:您这里只有一个逻辑依赖链,无法使用多流以及多流之间实现计算和数据传输相掩盖。

大致如此,祝您好运~

谢谢版主的解答,但是不知道能不能讲一下延迟的时间掩盖问题,今天看了一下没怎么懂。不知道怎么保持束调度器繁忙可以掩盖延迟呢?
前面讲如果要隐藏L个时钟的延迟,在计算能力为1.x需要L/4条指令,在计算能力为2.0的上面需要L,2.1需要2L,3.x需要8L。但是具体到后面如果是存储器依赖引起的延迟,常见的是22个时钟周期,那么对于编程指南里面讲的1.x需要6个束,2.x需要22个束,3.x需要44个束,但是按照上面的不是应该需要22/4,22,222,228条指令么?这个是怎么对应起来的?
而且在后面由于片下存储器引起的延迟的内置中,假设在3.x设备上延迟了300个周期,这里讲的是需要40个束来掩盖延迟,和上面的也不同啊。。
麻烦斑竹指导一下,谢谢。。

LZ您好:

新问题请新发帖,以及提供详细的教程超出了本版的范围。

一般而言,在保证有足够多resident threads的情况下,访问寄存器,访问shared memory,和常见计算指令所带来的延迟是容易被掩盖的,无须担心。

对于访问global memory这样的长延迟操作,如果延迟有一定影响,那么建议适当将读取位置提前,并在读取和使用之间插入多条无关的其他计算指令。

大致建议如上,由于手册中并无各项指令的具体延迟情况以及SM 调度规律的介绍,不再深入讨论了。

祝您好运~

建议使用CUDA 5.5的nvvp, 里面会先分析程序的瓶颈再给出相应优化建议。比CUDA 5.0智能不少。

谢谢建议。。