各位好，本人学习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是否花费巨大？