求图像相关邻域计算优化

加速计算 CUDA
1

初学CUDA的开发，实现了一个功能，主要是对图像进行运算，对行列值为6560*5792的一个数据进行运算，
改写后的CUDA代码如下，想请楼主帮忙看看为什么与cpu比加速比并不大？希望可以提供点帮助

主代码如下：
int blocksize = 32;
long nx_g=6560,ny_g=5792;
dim3 numthreads(blocksize,blocksize);
dim3 numblocks((nx_g+blocksize-1)/blocksize,(ny_g+blocksize-1)/blocksize);

//循环处理
while(1)
{
*cpustop = 1;
cudaMemcpy(gpustop, cpustop, sizeof(int), cudaMemcpyHostToDevice);
GPUDemDeal<<<numblocks,numthreads>>>(pGpuDemData,pGpuFillData,gpustop,gpunx,gpuny);
cudaMemcpy(cpustop, gpustop, sizeof(int), cudaMemcpyDeviceToHost);

if(*cpustop) break;

}

GPU实现代码
global void GPUDemDeal (float *pDemData,float *pFillData,int *gpuStop,long nx,long ny)
{
int j = blockIdx.yblockDim.y+threadIdx.y;
int i = blockIdx.xblockDim.x+threadIdx.x;

 float				mindata = 0;
 float              datpol[8];
long				nxvalue = *nx;
long				nyvalue = *ny;

if(i<nxvalue && j<nyvalue)
{
	if(pFillData[j*nxvalue+i]>pDemData[j*nxvalue+i])
	{
		//8领域处理
		datpol[0]=pFillData[(j-1)*nxvalue+i];
		datpol[1]=pFillData[(j-1)*nxvalue+i-1];
		datpol[2]=pFillData[j*nxvalue+i-1];
		datpol[3]=pFillData[(j+1)*nxvalue+i-1];
		datpol[4]=pFillData[(j+1)*nxvalue+i];
		datpol[5]=pFillData[(j+1)*nxvalue+i+1];
		datpol[6]=pFillData[j*nxvalue+i+1];
		datpol[7]=pFillData[(j-1)*nxvalue+i+1];

		for(int k=0;k<8;k++)
		{
			mindata=(pDemData[j*nxvalue+i]-datpol[k])/10;
			if(mindata>0)
			{
				if(pDemData[j*nxvalue+i]>=datpol[k]+mindata)
				{
					pFillData[j*nxvalue+i]=pDemData[j*nxvalue+i];
					*gpuStop=0;	break;
				}
			}
			mindata=(pFillData[j*nxvalue+i]-datpol[k])/1000;
			if(mindata>0)
			{
				if(pFillData[j*nxvalue+i]>datpol[k]+mindata)
				{
					pFillData[j*nxvalue+i]=datpol[k]+mindata;
					*gpuStop=0;
				}
			}		
		}
	}
}

}

2

LZ您好：

我大致查看了您的代码，比较明显的问题有：

1：您的算法中有大量的分支，而且这些分支是依赖于数据的，无法预先设计为分支按warp对齐，这会对GPU的执行效率产生较大的影响。

2：您的算法是某种迭代实现的，每次kernel计算完成，都需要将gpustop这个flag传回host端，然后判断是否进行下一步的计算，这样来回与host端通信，也将极大地影响您的GPU效率。不过这一点可以通过使用最新的SM3.5的kepler显卡提供的DP（Dynamic Parallelism）功能，将判断是否迭代的逻辑转移到device端来解决。

初步的意见如上，供您参考，欢迎其他网友/斑竹/总版主/NV原厂支持人员参与讨论，加以补充。

祝您编码顺利~

3

非常感谢版主的指点，还想请问一下block和thread的大小会影响多少效率啊？

4

LZ您好：

这里讨论的实际上是blocks per grid的数量，和threads per block的数量。

而实际上，前者一般不用讨论，因为您的算法决定之后，每个thread的工作量也就决定了。然后您选定适当的threads per block的大小，用总的工作量除以您每个block的工作量，实际上就是blocks per grid的数量了，也就是说这个是根据您的总工作量和算法实现确定的。

至于threads per block，一般建议选择192,256,512这样适中的典型值。您在1#中使用的1024个thread的大型block，这种block在fermi架构上会影响occupancy，因为fermi上一个SM最大resident threads数量是1536，如果一个block有1024个threads的话，那么这个SM上只能resident 1个block，从而达不到resident threads的上限（因为不能拆分block）。

此外，每个block在SM上计算完成的时候，是整体让出资源而供其他block加载的。此时如果一个SM上的resident blocks数量较少，比如说只有一个，那么当这个block里面还有少量线程没有结束的时候，其他block是无法加载的，而这个block里面的少量剩余线程又不足以充分利用SM的计算资源，会造成一些硬件闲置的现象，从而影响总体效率。

所以，一般都建议选择适中大小的典型值作为threads per block的数量。

以及，您这个问题中，block大小可能不是主要矛盾，您可以简单测试一下。

大致如此，祝您编码顺利~

5

非常感谢版主的指教，我测了一下修改threads per block数值后效率确实没有多大改变。您上面提到的pynamic parallelism
方法，刚刚接触想问一下版主有没有示例可以学习一下，非常感谢