模板相关遇到的问题

斑竹,你好。最近写了一个模板相关的程序。就是对于两幅图,各取一个大小为(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):也不知道问题讲清楚没有,如果没有后面在补充吧。还有代码怎么不能加上颜色或者加粗呢?谢谢。

LZ您好:

您的问题分别答复如下:

1:根据您的截图信息,这表示您一个warp的global 访存需要超过5次才能完成。而完全合并访问的情况下,对您使用的4B宽度的数据,一个warp只需一次即可完成。

2:shared6[0M2+tdyblockDim.x+tdx]=Img1[(y+i)Width+x+j]; //f
shared6[1
M2+tdy*blockDim.x+tdx]=Img2[(y+i)*Width+x+j]; //g

这两段代码赋值号右边是global 读取。
其中的i,j是block内线程的二维局部index+固定的偏移量。
Width是您图像的宽度,一般而言是一个较大的值。
y和x是由SubMatrix矩阵查表得到的数据,其中在查表时用到bky和bkx这两个变量,这两个变量是block的index。

因此考虑一个block内的一个warp,warp内各线程的bky和bkx是相同的,所以SubMatrix查表时自变量是相同的,因而y和x是相同的。

因此您的global memory的访存规律化简为Img1[tdy*Width+tdx+C]的形式,C为固定的常数。

对于这个访存规律,以及考虑Width较大,以及不考虑对齐影响的情况下,其合并与否取决于您的block的形状,如果blockDim.x是32或者32的倍数,那么合并情况较好。如果blockDim.x值很小(也就是block形状过窄),您将会遇到严重的非合并访问。

鉴于您并没有给出block的维度信息,本问题讨论截至于此,请您根据实际情况调整您的代码。

3:如果同步是您算法要求的,这个无法避免,请使用即可。
以及,根据您的代码来看,您的同步都是在操作shared memory之后添加的,这个一般无法避免,除非您修改您的算法实现。
__syncthreads()不是特别影响性能,请勿纠结。

4:个人建议在第一次规约的时候做特殊处理,将这次规约后的长度变成2的幂次,之后就都是常规的折半规约实现了。

即:假定有522个数据需要规约求和,那么在第一次规约的时候将第512到第521(下标从0算)的数据规约相加到前面去,您可以启动10个线程来完成这一过程。这样,剩下需要规约的数据就成了最接近但不超过522的最大的2的幂次:512个数据了,再进行常规折半规约即可。

您可以参考这个方法写出您的预处理过程。

5:代码模式似乎无法设置加粗和彩色,也可能是无法直接添加,我没有深究。

大致如此,祝您好运~

谢谢斑竹的详细解答。。

不客气的,祝你好运~