同一程序在不同GPU上计算结果不同

求教各位,我同样的程序在不同的GPU上计算却得到不同的结果。
我程序本身是在一个装有Quadro FX 580显卡的机器上编写的,这块显卡的计算能力是1.x, 可以得到正确的计算结果;
我将程序放到装有Quadro 4000显卡的机器上(计算能力是2.0)运行后,得到的结果显示值都溢出了。
两台机器使用的都是CUDA 5.0,本身驱动也是最新的。
我程序本身也只是使用了shared memory而已,使用的是VS2010编程。

请问,在不同设备之间的程序移植需要改动什么吗?或者是需要在VS2010里改动什么参数?

谢谢各位~

楼主您升级到2.x后,得到了对subnormal(很小的接近0的一些特殊的float)数的支持,以及得到了融合乘加(FMA)的支持。

在新卡上编译运行的这2点将您的代码提高了精度。但为何却得不到“正确的结果”呢?

可能您原来的算法、公式本来需要一定的“误差”从而使得能得到您期盼的结果,而“提高的精度”却无法在您的算法下保证这点。

这是其一。

其二,您重新编译的时候,选择了compute_20,sm_20了么?您可以尝试下这个。以及还可以尝试下打开-use_fast_math, 看看能否得到类似1.x的结果。

(可以在您的cu属性的command line里填入-user_fast_math, 请试试看)

根据您的文字描述,只能给出这些建议。但不妨试试看。

谢谢你,确实是你说的这个问题。

需要在VS下Properties->CUDA C/C++ → Device → Code Generation里选择compute_20,sm_20。
但是我发现只是使用compute_20还是不行,需要将compute_30和compute_35都加进去才可以。

希望我的这个问题能给其他遇到类似问题的朋友有帮助,再次感谢版主的回答~

额。。。。

显然您的部分回答是错误的,您选择compute_20,sm_20即可。您的compute_30和35毫无意义(在贵卡上完全没有用到的机会的)。

不过恭喜您解决了。虽然您不知道您的哪一步解决了。

LZ您好,在计算能力2.0及以上的GPU上,使用shared memory有时需要添加volatile关键字的,或者添加额外的__syncthreads()同步。

具体可以参考CUDA C Programming Guide中规约那个例子的代码,以及Fermi compatibility Guide中的说明。
同一warp内部连续进行写入和读取shared memory的时候,避免因为shared memory被寄存器缓冲而没有真正写入造成读取错误。
在1.x计算能力的硬件中无此要求。

您可以参考如下链接的讨论:
http://cudazone.nvidia.cn/forum/forum.php?mod=viewthread&tid=6788

请您检查一下您的代码,看看是否是这个问题。

此外,添加compute_30,compute_35的编译参数,应该无法对您的结果产生影响,这个参数添加的PTX代码只针对kepler的GPU,而您是fermi的GPU。所以无法理解您添加了这些参数之后,“解决了问题”。请您也检查下此情况。

祝您编码顺利,早日彻底解决问题。

呃,我现在也想这么说了……
之前确实修改添加30和35能成功,而只选择20无法成功……
现在的情况是,无论怎样都失败……
我移植过去的程序,只跑成功了一次,然后就一直失败得不到正确的结果

谢谢你,关于compute_30的问题,我确实是错误的,关于你说的shared memory问题我会去检查一下的~

现在一个诡异的情况是,我跑一次CUDA samples里的程序一次,在运行自己的程序就可以得到正确的结果了……
费解……
????

LZ您好,这个我也不知道是何原因了,建议您淡定地检查下您的代码。

祝您好运~