程序调试出问题

万能的斑竹大人,帮我看看这个程序那里有问题。我是在原来那个图像相关的例子上修改的,原来帖子链接如下:http://cudazone.nvidia.cn/forum/ … 7019&extra=page%3D1,原来在斑竹的指导下,运行很正常,后面经过修改如下:

int bky=blockIdx.y;         
   int bkx=blockIdx.x;
   int tdy=threadIdx.y; 
   int tdx=threadIdx.x;
   int y=SubMatrix[bky*m_nRow+bkx+m_nRow*m_nCol]; //原图block的中心点
   int x=SubMatrix[bky*m_nRow+bkx];        
   int i,j,tempv;
   float XPos,YPos;
   __shared__        float result[6],piArray[6];
   __shared__        float corr,iter,fm,gm,fsum,gsum,fgsum;
   __shared__ int M2; 
__shared__ float shared6[5808];

   if (tdy*blockDim.x+tdx==0){
   for (i=0;i<6;i++){        //初始参数
   result[i]=0;
   piArray[i]=InitDeformPr[(bky*gridDim.x+bkx)*8+i];
   }
   }
   M=M/2;
   i=tdy-M;
   j=tdx-M; //以块的中心为坐标
   M2=(2*M+1)*(2*M+1); //(2*M+1)*(2*M+1)常用

   /************************************************************************/
   /*求以(x,y)为中心的,(2M+1)*(2M+1)为模板的的修正值(即解cdpi,cdpj的方程)  */
   /************************************************************************/        
   do 
   {
   //////////////////////////////////////////////////////////////////////////
   //提供计算数据
   if(tdy*blockDim.x+tdx==0){//这里是对share memory的写入,要防止循环写入!
   matrix_subGPU(piArray,6,1,result,piArray);
   }
   if(y+i<0||x+j<0||y+i>=Height||x+j>=Width){ //线程出界取值为0,不要return不然不能同步
   shared6[0*M2+tdy*blockDim.x+tdx]=0;
   }else
   shared6[0*M2+tdy*blockDim.x+tdx]=m_ptrData1[(y+i)*Width+x+j];  //0~f
   
   YPos=y+i+piArray[3]+piArray[4]*j+piArray[5]*i; //点(x,y)通过形函数对应点(g_x,g_y)
   XPos=x+j+piArray[0]+piArray[1]*j+piArray[2]*i;
   
   if(int(XPos)<0||int(YPos)<0||int(XPos)>Width-1||int(YPos)>=Height-1){
   shared6[1*M2+tdy*blockDim.x+tdx]=0;
   } 
   GetAPixelGPU(XPos,YPos,CoefMatrix16,shared6[1*M2+tdy*blockDim.x+tdx],Width); //1~g

   //////////////////////////////////////////////////////////////////////////
   //用规约求fm,gm
   shared6[2*M2+tdy*blockDim.x+tdx]=shared6[0*M2+tdy*blockDim.x+tdx]; //数据,这里传给2、3是为了不让规约运算改变原始值
   shared6[3*M2+tdy*blockDim.x+tdx]=shared6[1*M2+tdy*blockDim.x+tdx];
   __syncthreads();

   for (tempv=M2/2;tempv>0;tempv=tempv>>1){ //奇数做规约
   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];
   }
   __syncthreads();
   }
   if (tdy*blockDim.x+tdx==0){ 
   fm=shared6[2*M2+0];
   gm=shared6[3*M2+0];
   fm+=shared6[2*M2+blockDim.y*blockDim.x-1];
   gm+=shared6[3*M2+blockDim.y*blockDim.x-1];
   fm=fm/(M2);
   gm=gm/(M2);
   }

   //////////////////////////////////////////////////////////////////////////
   //用规约求相关系数cor
   shared6[2*M2+tdy*blockDim.x+tdx]=(shared6[0*M2+tdy*blockDim.x+tdx]-fm)*(shared6[0*M2+tdy*blockDim.x+tdx]-fm); //数据,这里传给3,4,5是为了不让规约运算改变原始值
   shared6[3*M2+tdy*blockDim.x+tdx]=(shared6[1*M2+tdy*blockDim.x+tdx]-gm)*(shared6[1*M2+tdy*blockDim.x+tdx]-gm);
   shared6[4*M2+tdy*blockDim.x+tdx]=(shared6[0*M2+tdy*blockDim.x+tdx]-fm)*(shared6[1*M2+tdy*blockDim.x+tdx]-gm); 
   __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];
   }
   __syncthreads();
   }
   if (tdy*blockDim.x+tdx==0){ 
   fsum=shared6[2*M2+0];
   gsum=shared6[3*M2+0];
   fgsum=shared6[4*M2+0];
   fsum+=shared6[2*M2+blockDim.y*blockDim.x-1];
   gsum+=shared6[3*M2+blockDim.y*blockDim.x-1];
   fgsum+=shared6[4*M2+blockDim.y*blockDim.x-1];
   corr=fgsum/sqrt(fsum*gsum);
   }

在这里前面的帖子是直接坐两个图像的相关,这里加入了一个对应关系:

YPos=y+i+piArray[3]+piArray[4]*j+piArray[5]*i; //点(x,y)通过形函数对应点(g_x,g_y)
XPos=x+j+piArray[0]+piArray[1]*j+piArray[2]*i;

即原图的(y+i,x+j)对应变形后的(XPos,YPos),因为这个地方可能是小数,所以进行插值,插值函数如下:

GetAPixelGPU(XPos,YPos,CoefMatrix16,shared6[1*M2+tdy*blockDim.x+tdx],Width); //1~g

插值函数是没问题的,已经和串行的结果比对过。

现在自己发现的问题是在这里,有时候代码执行到最后时,发现

shared6[2*M2+tdy*blockDim.x+tdx]=(shared6[0*M2+tdy*blockDim.x+tdx]-fm)*(shared6[0*M2+tdy*blockDim.x+tdx]-fm); //数据,这里传给3,4,5是为了不让规约运算改变原始值
shared6[3*M2+tdy*blockDim.x+tdx]=(shared6[1*M2+tdy*blockDim.x+tdx]-gm)*(shared6[1*M2+tdy*blockDim.x+tdx]-gm);
shared6[4*M2+tdy*blockDim.x+tdx]=(shared6[0*M2+tdy*blockDim.x+tdx]-fm)*(shared6[1*M2+tdy*blockDim.x+tdx]-gm); 

这里面根本没有执行,数据都为xxxe+10,xxxe+11这种巨大的数据,有时候又是对的,有时候每次求出的均值有不同,这个到底是怎么回事啊?一直怀疑是同步的时候有问题,但是一直不知道到底是不是。。万能的斑竹。。求指导。。

楼主你的do{}后面没while结束么?发完整点?

这个真心看不出神马。不过直觉的说(虽然你的循环发了一半),在do{}体里的第一行加个__syncthreads();看看(不一定有效)

斑竹我找到地方了。。其实和后面代码没有关系。。

嗯嗯。我看到楼主您在补充完整代码后,最后有__syncthreads()的。那么上文给出的推测性建议将不再有效,请无视上文建议。

麻烦斑竹费心给仔细想想,都憋了一天了。。

楼主您在:
GetAPixelGPU(XPos,YPos,CoefMatrix16,shared6[1M2+tdyblockDim.x+tdx],Width); //1~g

使用了shared6[1M2+tdyblockDim.x+tdx]

然后您上文有shared6[1M2+tdyblockDim.x+tdx]的写入

换句话说,您的线程编号相差M2的2个线程,将分别进行写入和读取,但是您没有执行同步,您看这样安全吗?

还有个问题,就是有时候相关系数会大于1,到1.00几。。

嗯,谢谢斑竹的提醒。斑竹的意思是在GetAPixelGPU(XPos,YPos,CoefMatrix16,shared6[1M2+tdyblockDim.x+tdx],Width); //1~g后面加个同步操作么?我原来是把同步操作加载赋值过后的,因为感觉是在规约前面需要同步。

GetAPixelGPU(XPos,YPos,CoefMatrix16,shared6[1*M2+tdy*blockDim.x+tdx],Width); //1~g
	//////////////////////////////////////////////////////////////////////////
	//用规约求fm,gm
		shared6[2*M2+tdy*blockDim.x+tdx]=shared6[0*M2+tdy*blockDim.x+tdx]; 
		shared6[3*M2+tdy*blockDim.x+tdx]=shared6[1*M2+tdy*blockDim.x+tdx];
		__syncthreads();

后面我在函数GetAPixelGPU后面也加了同步操作,但是还是有问题啊。这里有个很明显的错误

if (tdy*blockDim.x+tdx==0){ //见correlation的3条注解
			fm=shared6[2*M2+0];
			gm=shared6[3*M2+0];
			fm+=shared6[2*M2+blockDim.y*blockDim.x-1];
			gm+=shared6[3*M2+blockDim.y*blockDim.x-1];
			fm=fm/(M2);
			gm=gm/(M2);
		}

在这里面shared6[2M2+blockDim.yblockDim.x-1]、shared6[3M2+blockDim.yblockDim.x-1]的数据分别为:-4116024,3.26e+09,但是前面的规约值两个都才为57239,好像是在做规约的时候,后面部分的数据被改变了。。但是这个错误有时有有时又没有。。不知道怎么回事。。

斑竹还有个明显的问题就是如果我一步步调试还好,如果我直接一下把断点加在

if (tdy*blockDim.x+tdx==0){ 
   fsum=shared6[2*M2+0];
   gsum=shared6[3*M2+0];
   fgsum=shared6[4*M2+0];
   fsum+=shared6[2*M2+blockDim.y*blockDim.x-1];
   gsum+=shared6[3*M2+blockDim.y*blockDim.x-1];
   fgsum+=shared6[4*M2+blockDim.y*blockDim.x-1];
   corr=fgsum/sqrt(fsum*gsum);
   }

这里,这里会明显出错,shared6[2M2+tdyblockDim.x+tdx]、shared6[2M2+tdyblockDim.x+tdx]、shared6[2M2+tdyblockDim.x+tdx]里面的数据全部错误。。都为xxx*e+10这种形式。。好像前面赋值根本没进行。

这个真心不知道怎么回事了。。代码快200行。又不懂您的算法,真心爱莫能助。也许其他会员、版主、NVIDIA原厂支持以及总版主能为您人肉调试下。

没事,谢谢斑竹,有可能是这里要加上同步:

__syncthreads(); //新加的
   shared6[2*M2+tdy*blockDim.x+tdx]=(shared6[0*M2+tdy*blockDim.x+tdx]-fm)*(shared6[0*M2+tdy*blockDim.x+tdx]-fm); //数据,这里传给3,4,5是为了不让规约运算改变原始值
   shared6[3*M2+tdy*blockDim.x+tdx]=(shared6[1*M2+tdy*blockDim.x+tdx]-gm)*(shared6[1*M2+tdy*blockDim.x+tdx]-gm);
   shared6[4*M2+tdy*blockDim.x+tdx]=(shared6[0*M2+tdy*blockDim.x+tdx]-fm)*(shared6[1*M2+tdy*blockDim.x+tdx]-gm); 
   __syncthreads();

我在调试一下,谢谢斑竹了。。