用户
 找回密码
 立即注册
cherubicwing 该用户已被删除
发表于 2013-11-27 22:10:04
68984
初学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.y*blockDim.y+threadIdx.y;
        int i = blockIdx.x*blockDim.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;
                                        }
                                }               
                        }
                }
        }
}

使用道具 举报 回复
发表于 2013-11-27 22:44:20
LZ您好:

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

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

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

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

祝您编码顺利~
使用道具 举报 回复 支持 反对
发表于 2013-11-28 08:26:27
非常感谢版主的指点,还想请问一下block和thread的大小会影响多少效率啊?
使用道具 举报 回复 支持 反对
发表于 2013-11-28 11:47:56
cherubicwing 发表于 2013-11-28 08:26
非常感谢版主的指点,还想请问一下block和thread的大小会影响多少效率啊?

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大小可能不是主要矛盾,您可以简单测试一下。

大致如此,祝您编码顺利~
使用道具 举报 回复 支持 反对
发表于 2013-11-28 16:10:01
非常感谢版主的指教,我测了一下修改threads per block数值后效率确实没有多大改变。您上面提到的pynamic parallelism
方法,刚刚接触想问一下版主有没有示例可以学习一下,非常感谢
使用道具 举报 回复 支持 反对
发新帖
您需要登录后才可以回帖 登录 | 立即注册