斑竹,你好。最近写了一个模板相关的程序。就是对于两幅图,各取一个大小为(2M/2+1)(2*M/2+1)大小的模板,算两者的相关系数。在调试过程中说合并访问有问题。具体的截图如下:
[attach]3396[/attach]
具体代码如下:
其中Img1、Img2代表的是图像1、2,Height、Width代表的他们的高和宽,SubMatrix存储的是各个目标的中心坐标,CorCoef是对应模板的相关系数,m_nRow,m_nCol是把图像划分成m_nRow×m_nCol这么多个模板。
__global__ void TemplateMatch(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;
// extern __shared__ float shared6[]; //大小是(M/2*2+1)*(M/2*2+1)*3,存储两个模板的值
__shared__ float shared6[4847]; //大小是(M/2*2+1)*(M/2*2+1)*3,存储两个模板的值
M=M/2;
i=tdy-M;
j=tdx-M; //以块的中心为坐标
int M2=(2*M+1)*(2*M+1);
//////////////////////////////////////////////////////////////////////////
//求以(x,y)为中心的,(2M+1)*(2M+1)为模板的的相关系数
shared6[0*M2+tdy*blockDim.x+tdx]=Img1[(y+i)*Width+x+j]; //f
shared6[1*M2+tdy*blockDim.x+tdx]=Img2[(y+i)*Width+x+j]; //g
shared6[2*M2+tdy*blockDim.x+tdx]=shared6[0*M2+tdy*blockDim.x+tdx]; //数据,这里传给0,1,2是为了不让规约运算改变原始值
shared6[3*M2+tdy*blockDim.x+tdx]=shared6[1*M2+tdy*blockDim.x+tdx];
__syncthreads();
for (tempv=M2/2;tempv>0;tempv=tempv>>1){ //奇数做规约,这里用i节省空间
if (tdy*blockDim.x+tdx<tempv){
shared6[2*M2+tdy*blockDim.x+tdx]+=shared6[2*M2+tdy*blockDim.x+tdx+tempv];
shared6[3*M2+tdy*blockDim.x+tdx]+=shared6[3*M2+tdy*blockDim.x+tdx+tempv];
}
if (tempv%2!=0&&tdx==0&&tempv!=1){
shared6[2*M2+tdy*blockDim.x+tdx]+=shared6[2*M2+tdy*blockDim.x+tdx+tempv-1];
shared6[3*M2+tdy*blockDim.x+tdx]+=shared6[3*M2+tdy*blockDim.x+tdx+tempv-1];
}
__syncthreads();
}
if (tdy*blockDim.x+tdx==0){
x1=shared6[2*M2+0];
x2=shared6[3*M2+0];
x1+=shared6[2*M2+blockDim.y*blockDim.x-1]; //做的奇数规约x1~fm,x2~gm
x2+=shared6[3*M2+blockDim.y*blockDim.x-1];
x1=x1/(M2);
x2=x2/(M2);
}
__syncthreads();
//用规约求相关系数cor
shared6[2*M2+tdy*blockDim.x+tdx]=(shared6[0*M2+tdy*blockDim.x+tdx]-x1)*(shared6[0*M2+tdy*blockDim.x+tdx]-x1); //数据,这里传给0,1,2是为了不让规约运算改变原始值
shared6[3*M2+tdy*blockDim.x+tdx]=(shared6[1*M2+tdy*blockDim.x+tdx]-x2)*(shared6[1*M2+tdy*blockDim.x+tdx]-x2);
shared6[4*M2+tdy*blockDim.x+tdx]=(shared6[0*M2+tdy*blockDim.x+tdx]-x1)*(shared6[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){
shared6[2*M2+tdy*blockDim.x+tdx]+=shared6[2*M2+tdy*blockDim.x+tdx+tempv];
shared6[3*M2+tdy*blockDim.x+tdx]+=shared6[3*M2+tdy*blockDim.x+tdx+tempv];
shared6[4*M2+tdy*blockDim.x+tdx]+=shared6[4*M2+tdy*blockDim.x+tdx+tempv];
}
if (tempv%2!=0&&tdx==0&&tempv!=1){
shared6[2*M2+tdy*blockDim.x+tdx]+=shared6[2*M2+tdy*blockDim.x+tdx+tempv-1];
shared6[3*M2+tdy*blockDim.x+tdx]+=shared6[3*M2+tdy*blockDim.x+tdx+tempv-1];
shared6[4*M2+tdy*blockDim.x+tdx]+=shared6[4*M2+tdy*blockDim.x+tdx+tempv-1];
}
__syncthreads();
}
if (tdy*blockDim.x+tdx==0){
x1=shared6[2*M2+0]; //x1~fsum/cor,x2~gsum,x3~fgsum
x2=shared6[3*M2+0];
x3=shared6[4*M2+0];
x1+=shared6[2*M2+blockDim.y*blockDim.x-1];
x2+=shared6[3*M2+blockDim.y*blockDim.x-1];
x3+=shared6[4*M2+blockDim.y*blockDim.x-1];
x1=x3/sqrt(x1*x2);
CorCoef_d[bky*m_nRow+bkx]=x1;
}
__syncthreads();
}
这里有以下几个问题:
1):截图中显示的合并访问的影响是有多严重?
2):截图中提示的代码是下面一段,这里是把模板的值存储到share中,这种访问应该怎么设计才能保证合并访问?
shared6[0*M2+tdy*blockDim.x+tdx]=Img1[(y+i)*Width+x+j]; //f
shared6[1*M2+tdy*blockDim.x+tdx]=Img2[(y+i)*Width+x+j]; //g
3):在求相关运算时,感觉要很多次同步,这里是不是对性能会有很大影响?怎么避免?
4):在规约求和的时候,发现数据量不是2^n次方,在网上看了一下讲要么是把数据补齐要么是划一块出来重新做,我在程序中是单独处理每次余下的那一块。请问一下那种方法要好些?
if (tempv%2!=0&&tdx==0&&tempv!=1){
shared6[2*M2+tdy*blockDim.x+tdx]+=shared6[2*M2+tdy*blockDim.x+tdx+tempv-1];
shared6[3*M2+tdy*blockDim.x+tdx]+=shared6[3*M2+tdy*blockDim.x+tdx+tempv-1];
}
5):也不知道问题讲清楚没有,如果没有后面在补充吧。还有代码怎么不能加上颜色或者加粗呢?谢谢。