CUDA程序优化问题

我的算法是做了个计算,程序如下:
global void SquaresKernel(const float Img, const float A,int width,int height,floatdest)
{
int x=threadIdx.x+blockIdx.x
blockDim.x;
int y=threadIdx.y+blockIdx.y*blockDim.y;

    int tid=y*width+x;//就是本线程处理的那个像素

    if(x<width&&y<height)
    {
            dest[tid]=(float)((A[0]-Img[tid*3])*(A[0]-Img[tid*3])+
                    (A[1]-Img[tid*3+1])*(A[1]-Img[tid*3+1])+
                              (A[2]-Img[tid*3+2])*(A[2]-Img[tid*3+2]));
    }

}
为什么这么做,还没有用C快,我的思路是对其进行优化,已经知道那个长度为三的数组在计算能力2.0以上的版本中自动放入常数存储器中,下一步使用纹理存储器来进行处理

LZ您好:

1:您的叙述有误:“已经知道那个长度为三的数组在计算能力2.0以上的版本中自动放入常数存储器中”。您这里是将global memory中的数组以指针形式传进kernel,这样做的话,使用这三个常数的时候,依然需要访问global memory的。

而之前给您的建议为“您的3个常数作为kernel参数就行了,在2.x及以上的硬件上,他们就在常数存储器中”,这指的是直接将这三个常数自身作为kernel的参数,而不是将他们放在global memory上,而将指针作为参数。
按照您的写法,指针A是在constant cache里面的,而不是3个具体的参数。

(待续)

2:为何这个代码还没有C快(理解为为何没有跑在CPU上的C语言的版本快)。这个有很多相关因素的,比如您的计算规模大小,比如您的GPU好坏,比如您的kernel启动规模是否合理,比如您是否估算了算上来回copy的时间用GPU计算是否值得,甚至请注意下您的代码是否是由debug模式编译的,论坛最近很多这个问题。

而您可以跑一下profiler,看看参数统计以及有什么建议,然后再做判断和修改。

大致如此,祝您好运~

“已经知道那个长度为三的数组在计算能力2.0以上的版本中自动放入常数存储器中”
–你从哪里已知的?

你需要写成(…, float a0, float a1, float a2…)方能自动进去。
这是实际上ICE说过了,但是你没看仔细。

关于您的第二个问题,C是啥?CPU么?
显然这个不会没有CPU快的,您这代码,基本卡死在访存上了(虽然你的访存的写法很不好),
但是比CPU慢基本无可能。您这个大量无关的像素处理在并行呢。不可能“比CPU还慢”。
(如果您的卡是GT210, 那另当别论,请无视此说法,直接换卡吧)

谢谢版主,
我的显卡是GT660,计算能力3.0,应该不是显卡的问题,像将height和width一样不写入全局存储器,将长度为3的数组,写成三个常量做为内核函数的传入参数

谢谢版主
我的显卡是GT660,计算能力3.0,应该不是显卡的问题,我也认为是卡死在访存上了,我是按照正常的写入全局存储器,那个长度为3的长量数组是频繁访问的数据,应该是这的问题

"谢谢版主,
我的显卡是GT660,计算能力3.0,应该不是显卡的问题,像将height和width一样不写入全局存储器,将长度为3的数组,写成三个常量做为内核函数的传入参数 "
LZ您好:

您5#的全部内容如上,请您将叙述补充完整,感谢合作。

祝您好运~

LZ您好:

1:昨晚我们内部简单评测了您的代码,发现即便您将长度为3的小数组直接放置在global memory里面,由于kepler L2cache的强大作用,直接就能获得不错的效果。在将这个数组直接作为参数传入kernel以后,测试结果得到了进一步的提升。

2:根据昨晚的测试结果,在某款较您的GPU规格更低的显卡上,实测获得了非常好的显存吞吐率,该吞吐率已经超过了目前主流桌面级别高端CPU内存吞吐率的数倍。尽管该算法访存十分密集,但是对于CPU而言也是访存密集型的,此时GPU有大量的并行线程在大量的并行计算单元上干活,并达成了数倍于CPU的访存带宽,因此无法理解您为何一直抱怨这没有CPU跑得快,您的叙述与我们实测的结果不符。

3:如果LZ继续坚持您的观点,请拿出令人信服的测试结果,而不是凭主观判断,并下结论。

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

谢谢版主
我的内核函数是这么编的,按照您的说法是将这个长度为三的数组写成三个常量的形式,不写入全局存储器,修改成__global__ void SquaresKernel(const float Img, float a1,float a2,float a3,int width,int height,floatdest)的形式吗

原来的内核函数如下:
global void SquaresKernel(const float Img, const float A,int width,int height,floatdest)
{
int x=threadIdx.x+blockIdx.x
blockDim.x;
int y=threadIdx.y+blockIdx.y*blockDim.y;

int tid=y*width+x;//就是本线程处理的那个像素

if(x<width&&y<height)
{
dest[tid]=(float)((A[0]-Img[tid3])(A[0]-Img[tid3])+
(A[1]-Img[tid
3+1])(A[1]-Img[tid3+1])+
(A[2]-Img[tid3+2])(A[2]-Img[tid*3+2]));
}
}

这是我调用内核函数的host端代码,您看看这么编时间是不是会卡在访存上,按照您的说法将长度为三的向量写成a1,a2,a3的形式,此时不需要这句,cudaStatus=cudaMalloc((void**)&dev_colorM,3*sizeof(float));
这么做就是将这个长度为三的数组自动写入上述存储器了吗

源代码如下:
/******************************************************************************/
cudaError_t Square(const float
hos_Icols,const float hos_colorM,int width,int height,float hos_dest)
{
float
dev_Icols=NULL;
float
dev_colorM=NULL;
float
dev_dest=NULL;
//测时
cudaEvent_t start, stop;
float time;
//检测状态
cudaError_t cudaStatus;
//设置为0号设备
cudaStatus=cudaSetDevice(0);
if(cudaStatus!=cudaSuccess)
{
fprintf(stderr,“cudaMalloc failed”);
goto Error;
}
//在device端分配内存(two inpot ,one output)
cudaStatus=cudaMalloc((void
)&dev_Icols,width
height
3
sizeof(float));
if(cudaStatus!=cudaSuccess)
{
fprintf(stderr,“cudaMalloc failed”);
goto Error;
}

cudaStatus=cudaMalloc((void**)&dev_colorM,3*sizeof(float));
if(cudaStatus!=cudaSuccess)
{
	fprintf(stderr,"cudaMalloc failed");
	goto Error;
}

cudaStatus=cudaMalloc((void**)&dev_dest,width*height*sizeof(float));
if(cudaStatus!=cudaSuccess)
{
	fprintf(stderr,"cudaMalloc failed");
	goto Error;
}
//将数据从host端复制到device端
cudaStatus=cudaMemcpy(dev_Icols,hos_Icols,width*height*3*sizeof(float),cudaMemcpyHostToDevice);
if(cudaStatus!=cudaSuccess)
{
	fprintf(stderr,"cudaMemcpy failed");
	goto Error;
}
cudaStatus=cudaMemcpy(dev_colorM,hos_colorM,3*sizeof(float),cudaMemcpyHostToDevice);
if(cudaStatus!=cudaSuccess)
{
	fprintf(stderr,"cudaMemcpy failed");
	goto Error;
}
//保证设备同步
cudaDeviceSynchronize();

//创建时间
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start,0);

//调用kernel函数
dim3 DimBlock(16,16,1);
dim3 DimGrid((width+DimBlock.x-1)/DimBlock.x,(height+DimBlock.y-1)/DimBlock.y,1);
SquaresKernel<<<DimGrid,DimBlock>>>(dev_Icols,dev_colorM,width,height,dev_dest);

//结束计时开始统计时间
cudaEventRecord(stop,0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&time,start,stop);

printf("kernel take time:%fms\n",time);
//将总的时间统计出来也就是所有视频帧在此块处所花费的时间
time_total+=time;

printf("total time=%fms\n",time_total);

cudaMemcpy(hos_dest,dev_dest,width*height*sizeof(float),cudaMemcpyDeviceToHost);
if(cudaStatus!=cudaSuccess)
{
	fprintf(stderr,"cudaMemcpy failed");
	goto Error;
}

Error:
cudaFree(dev_Icols);
cudaFree(dev_colorM);
cudaFree(dev_dest);

return cudaStatus;

}

谢谢版主
我用C语言编写时处理的速度是每秒23帧,而用GPU加速后居然变成了16帧每秒,速度下降了很多,您看看是不是我的这种Host段代码的数据访存的问题

LZ您好:

1:将三个常量直接作为kernel函数的参数即可,您给出的函数声明的形式是正确的。

2:您之前的kernel写法也是可以的,会比用参数传递这3个常量慢一点,但是也有还不错的效果。

3:如果您将这三个常数作为函数参数传递了,那么在启动kernel的时候会自动压入的,您无需再cudaMalloc申请空间并cudaMemcpy复制数值。

4:您只给出了一个函数的写法,目测并无明显问题。其中注释为“保证设备同步”的cudaDeviceSynchronize()可以不加。您的DimBlock大小是典型的,Dimgrid写法是正确的,(但不知规模如何)。

注意到您有time_total+=time;这样一句,以及您10#指出您实际是处理多帧的。

那么如果您就在您的代码中直接将您9#给出的函数予以循环,恐难得到好的效果。
因为您的函数没循环调用一次,就需要申请一次空间,用完再释放掉,这些是比较耗费时间的。而您申请的空间,实际上每次计算都是可以重复使用的。所以每次copy即可,而无需释放后再申请。

此外,考虑到一帧视频文件一般并不会很大,您也可以使用多stream的方法,实现异步传输,使得计算和传输互相掩盖,这样会进一步提升您的计算速度。

大致建议如上,祝您编码顺利~

谢谢楼主
我按照您的说法去修改,版主辛苦了

不客气的,欢迎您常来论坛讨论~

祝您编码顺利~