程序写出来了，但是发现优化基本上不怎么会啊，问题一大堆。。先下了一个图像相关的程序。代码如下：其中：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）：不知道还有没有其他问题没有想到的，希望版主多多提示。。谢谢