用户
 找回密码
 立即注册
yht1120 该用户已被删除
发表于 2013-11-27 17:02:53
2206715
       有一组char型的实验数据,需要转化到float类型后再做处理,数据长度为10M
       我分别编写了gpu和cpu的代码,发现gpu处理时间是2.1ms,而cpu的处理时间是26ms。当然,可以保证的是,这两段代码的处理结果都是正确的,gpu的处理时间用的cudaEvent获取,cpu端处理时间用的QueryPerformanceCounter获取。  
       gpu端的调用为DataCopyGPU<<<10000,1024>>>。相比之下,gpu处理速度才提高了10几倍,想问下大家怎样才可以提高gpu的处理速度,还有这段代码的瓶颈在哪个地方。

GPU端代码
__global__  void DataCopyGPU(char *pBufChar,float *pBufFloat)
{
        int threadID= threadIdx.x;
        int blockID=  blockIdx.x;
        int IndexThread=blockID*blockDim.x+threadID;

        *(pBufFloat+IndexThread)=*(pBufChar+IndexThread);
}

CPU端代码
void DataCopyCPU(char *pBufChar,float *pBufFloat,int nByte)
{
        int i=0;
        for(i=0;i<nByte;i++)
                *(pBufFloat+i)=*(pBufChar+i)
}
使用道具 举报 回复
发表于 2013-11-27 18:36:14
本帖最后由 yht1120 于 2013-11-28 12:09 编辑

谢版主了

GPU代码的计时如下
        cudaEventRecord(start,0);
                DataCopyGPU<<<10000,1024>>>(pBufChar_D,pBufFloat_D);
        cudaEventRecord(stop,0);
        cudaEventSynchronize(stop);
        cudaEventElapsedTime(&elapsedTime,start,stop);

CPU代码的计时如下
        QueryPerformanceCounter(&Lstart);
                DataCopyCPU(pBufChar,pBufFloat,10000*1024);
        QueryPerformanceCounter(&Lend);
        elapsedTime=(float)(Lend.QuadPart-Lstart.QuadPart)/Frequency.QuadPart*1e3;

这两段计时代码好像没什么问题吧,都是在论坛上搜到的
使用道具 举报 回复 支持 反对
发表于 2013-11-27 17:26:11
LZ您好:

您这样写没问题,但是这是一个访存密集型的操作,最终可能卡在您的显存速度上,而您的显存吞吐量已经定死了,没有太多可以优化的余地。

大致如此,供您参考。

祝您编码顺利~
使用道具 举报 回复 支持 反对
发表于 2013-11-27 17:45:55
谢谢版主,我用cuda-z查到的显卡信息为
Host to Device  5.7GiB/s
Device to Host  6.3GiB/s
Device to Device 106.3GiB/s

按实际情况计算,10M/2.1ms=5GB/s,这个与106GB/s差的还相当远啊,应该没有卡在这个速度上了啊,那这个是怎么回事
使用道具 举报 回复 支持 反对
发表于 2013-11-27 18:14:14
yht1120 发表于 2013-11-27 17:45
谢谢版主,我用cuda-z查到的显卡信息为
Host to Device  5.7GiB/s
Device to Host  6.3GiB/s

LZ您好:

请您提供一下您测时的具体代码,或者profiler的结果也可以。
使用道具 举报 回复 支持 反对
发表于 2013-11-27 20:19:33
yht1120 发表于 2013-11-27 18:36
谢版主了

GPU代码的计时如下

LZ您好:

我不清楚是为何了,您这样一个访存密集型的操作,而且访问是合并的,结果只跑到了您显存带宽的1/20么?

使用道具 举报 回复 支持 反对
发表于 2013-11-28 09:38:18
谢版主
再问下,这个访问速度低会不会和项目配置有关系,会不会有什么设置限制了计算速度。
另外,传递的两个参数指针都是使用cudaMalloc创建的,使用的应该就是显卡的全局内存吧,这个要不要求指针的首地址是128或256的倍数? 是和不是对访问速度有多大的影响?
谢谢
使用道具 举报 回复 支持 反对
发表于 2013-11-28 11:00:25
瓶颈在访存延迟而不在带宽。
需要每线程处理多个数据,注意充分利用寄存器。
使用道具 举报 回复 支持 反对
发表于 2013-11-28 12:06:40
yht1120 发表于 2013-11-27 17:45
谢谢版主,我用cuda-z查到的显卡信息为
Host to Device  5.7GiB/s
Device to Host  6.3GiB/s

你的global函数里用了float,应该用40M计算而不是10M吧

你的统计时间里的代码用start stop后面却蹦出一个end,是手误?

我用的卡用cuda z运行 D to D也正好 106GB/s左右,然后你这段代码kernel时间在0.45ms - 0.5 ms之间

使用道具 举报 回复 支持 反对
发表于 2013-11-28 12:12:29
gpu 发表于 2013-11-28 11:00
瓶颈在访存延迟而不在带宽。
需要每线程处理多个数据,注意充分利用寄存器。
...

谢谢,按照你的说法修改了一下kernel,每个线程处理了128个数据,最后处理时间编程1.8ms了,提高了10%,谢谢
使用道具 举报 回复 支持 反对
12下一页
发新帖
您需要登录后才可以回帖 登录 | 立即注册