数组赋值求优化

有一组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)
}

LZ您好:

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

大致如此,供您参考。

祝您编码顺利~

谢谢版主,我用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差的还相当远啊,应该没有卡在这个速度上了啊,那这个是怎么回事

LZ您好:

请您提供一下您测时的具体代码,或者profiler的结果也可以。

谢版主了

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,100001024);
QueryPerformanceCounter(&Lend);
elapsedTime=(float)(Lend.QuadPart-Lstart.QuadPart)/Frequency.QuadPart
1e3;

这两段计时代码好像没什么问题吧,都是在论坛上搜到的

LZ您好:

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

谢版主
再问下,这个访问速度低会不会和项目配置有关系,会不会有什么设置限制了计算速度。
另外,传递的两个参数指针都是使用cudaMalloc创建的,使用的应该就是显卡的全局内存吧,这个要不要求指针的首地址是128或256的倍数? 是和不是对访问速度有多大的影响?
谢谢

瓶颈在访存延迟而不在带宽。
需要每线程处理多个数据,注意充分利用寄存器。

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

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

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

谢谢,按照你的说法修改了一下kernel,每个线程处理了128个数据,最后处理时间编程1.8ms了,提高了10%,谢谢

谢谢
复制到论坛粘贴的时候搞错了,源代码没有问题
数据是从char类型复制到float类型

我的运行时间确实是2ms左右,按照8#的修改了下之后时间变为1.8ms了。。
但是确实奇怪了,怎么你那儿的运行时间那么短,而在我机子上运行起来时间却这么长,我确定应该就是我这儿出问题了,但是到底是怎么回事呢?

NO,如果你做的好,应该不是10%。10%跟没提高差不多啊,目标是接近理论上限。

难道楼主测试的是第一次运行的时间?

如果是这样,请从第二次开始测试。

或者,请使用profiler报告时间而不要自己测试。

请您尝试下此建议是否适用。

以及,如果您的项目是debug配置编译的,请立刻改用Release。

Debug下只是为了调试清晰方便的,如果需要性能,请考虑使用Release.
(Debug下的性能评估无任何意义的)

请您考虑上条建议和本建议。
感谢来访。

以及,LZ不会是debug下测试的吧,如果是的话,请release编译。

如果已经是release编译,请无视这个建议。

以及请LZ详细贴一下您的GPU的信息。

如果是在开普勒上,作如下调整:
1 DataCopyGPU<<<10000,1024>>>改为:DataCopyGPU<<<10000*8,128>>>。
2 每线程处理16个数据。
应该轻松达到带宽理论上限。
如果不是开普勒,由于硬件不同,请适当做出调整。