各位好,本人学习CUDA编程时,遇到的奇怪问题如下,在CUDA编程手册3.2.2节有介绍矩阵相乘以及优化的相关内容。本人按照手册中的程序将其完成后,采用cudaEvent计时,在GTX680显卡上运行,分别运行16001600大小的全部元素初始化为1的两矩阵相乘统计时间,校验计算结果正确后,得出的结果是未经shared memory优化的矩阵相乘程序快过优化后时间分别为390.904633ms与458.45260ms。之后又在GTX560Ti以及GT630M上均是未经过优化的程序更快得出结果。
本人以为是自己实现错误,于是在下述链接找到附件中文章一篇
有实验数据辅助,粘贴其第三章,第四章代码同样运行16001600大小全部元素为1的矩阵乘法,在GTX680,GTX560Ti以及GT630M上得到的结果依然是未经优化的程序更快得出结果。
本人十分疑惑,不知道中间是哪里出了问题,请各位不吝赐教。
PS:时间测试代码段如下,其余代码与附件中代码相同
cudaEvent_t start,stop;
float time;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(start,0);
MatMulKernel<<<dimGrid, dimBlock>>>(d_A, d_B, d_C);
cudaThreadSynchronize();
cudaEventRecord(stop,0);
cudaEventSynchronize(stop);
cudaEventElapsedTime(&time,start,stop);
printf(“it cost time %f ms\n”,time);
cudaEventDestroy(start);
cudaEventDestroy(stop);[attach]3393[/attach]
[attach]3393[/attach]
LZ您好:
大致看过您提供的文献以及您的测时代码,您的测时是正确的。
至于您提出的问题简要表述我的看法如下:
1:无论programming guide给出的矩阵乘法算法还是您引用的文献,其实都是一种示例性的方法,如果您需要真正快速计算矩阵乘法,推荐使用GPU上的线性代数库,这些库进行了更多更为细致的优化,一般来说要比自行实现的计算更快。(注意只是一般来说,库因为考虑通用性,有时候比某个特定的实现要慢)
2:一种优化算法的实现,一般来说并不能保证在任何情况下都是优化的。您给出的参考文献的计算结果表明,在2.0的计算能力的GPU上(这个测试和您使用的GPU特性比较接近),在矩阵尺寸为8000*8000的时候,使用shared memory有良好的效果,而且1000~2000的时候,两者差距不大,在1000以下的时候,结果耗时结果经常有反复和交错的现象。
我认为这个结果对您有一定参考价值,您也不妨测试一下更大计算量的情况。
同时您引用的文献对不同计算能力版本的GPU的测试结果还表明,对于越早的GPU,shared memory在一开始就可以取得较好的优化效果,但是越往后,需要有较大的计算规模,shared memory的优势才能得以显现。这大概是因为,global memory直接访问的能力一直在增强。
3:实际上对于不同的计算数据规模,不同的计算速度的GPU,不同的计算能力版本的硬件特性,往往需要根据实际情况具体分析当前的瓶颈在哪里并予以优化。这就说明了不存在一个统一的方法适合于所有情况的优化。
大致如此,供您参考。
祝您编码顺利~
感谢版主热心回复,祝版主身体健康,工作生活顺利。另外还有一个小问题叨扰版主,我在本机通过修改注册表中TdrLevel = 0后,依然出现运算时黑屏,显卡驱动无响应,使用Nsight monitor关闭TDR一样没有效果,不知道版主有没有方法或一些参考资料,再次感谢。
LZ您好:
关闭TDR之后您的程序在死循环的时候会一直算下去,以及如果有错误还是可能会其他问题的。
仅仅是不再检测超时而已。
以及关闭TDR后需要重启一下电脑。
大致有这几个注意事项。
祝您好运~
不知道版主您在使用CUDA的时候是否试过这段代码?如果有那您当时的结果是怎样的?
我刚刚在Tesla K20X,以及普通的GT650M上好像都无法重现这个问题。shared memory版的总是快得多?不知道Shi哥是在什么环境下测出来的结果?
我的环境:CUDA 5.0 built on VS2010 (Win 7 64 bits)
感谢佛倾情回复,我的环境CUDA 5.5 built on VS2010(Win 7 64 bits),难道是因为游戏卡的原因么?另外,佛是自己另写的代码还是用的附件中的代码呢?跪拜
yll您好:
我没有自己测试和评估过这段代码,上面的叙述是就LZ和LZ给出的文献数据按照一般规律说的。
欢迎您给出您的测试结果。
以及,不知LZ编译的时候是否都是按照release编译后测试的。
欢迎您继续参与讨论。
谢谢版主提醒,方才按照release编译后测试,确实是使用shared memory优化后的程序远远快于未优化程序。
能否麻烦版主大人与yll佛告知debug模式下出现这种情况的缘由,不甚感激。
我觉得你的结果略显不正常……你用GTX680需要花0.4s的时间,而我的GT650M甚至都不用这么久;此外你给的资料里M2070 @ 1000也远远比你的GTX680耗时少?而我印象中M2070的单精度运算是远远比不过GTX680的?
正如版主所说,你确定你是Release版的?
system
10
LZ您好:
debug模式下测试效能不具备任何参考价值,因为debug模式下有很多冗余的代码以支持调试功能,以及其实现是非优化状态的。
所以debug模式仅供调试使用。
祝您好运~
system
11
感谢版主以及yll大神的回复,祝二位工作生活顺利。
system
12
yll您好:
您说的很正确,我之前也觉得这个时间有些偏长,但是没有详细评估这一点。看到文献中小尺寸矩阵,两个实现的时间相差不多,就没有继续仔细考虑了,这个有失严谨。
GTX680的单精度效能要比M2070强不少,不过同时680的双精度效能比较差。
感谢您积极参与回帖讨论,使得LZ真正的问题得以尽快确定。
system
13
谢谢版主大人以及yll大神的回复,祝二位工作学习顺利。
system
15
正如版主大人说的,Debug版几乎没有测试意义~~
另外我斗胆地猜测下原因:
印象中从global到shared的copy在debug模式下是极慢无比的(加了一大堆校验指令,此外还会安插array head和array tail以检查你是否搞出类似下标溢出这样的操作具体不大记得了,如果版主了解的话还麻烦您补充),而release模式下这些指令会被移除所以我猜测导致Debug模式下shared更慢和这个memory copy有莫大联系~~
不过,事实上我也没亲自试过。如果Shi哥有兴趣倒是可以检查下kernel各部分的执行时间,看看这个copy是否花费巨大?