数组赋值求优化

加速计算 CUDA
1

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

2

LZ您好：

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

大致如此，供您参考。

祝您编码顺利~

3

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

4

LZ您好：

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

5

谢版主了

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.QuadPart1e3;

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

6

LZ您好：

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

7

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

8

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

9

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

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

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

10

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

11

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

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

12

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

13

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

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

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

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

14

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

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

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

15

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

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

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

16

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