Titan V做计算真的这么不靠谱么？

今年3月份有一篇文章在高性能计算这个领域算是引起了一个不小的波动英伟达的 Titan V GPU 计算故障：2+2=4，呃=4.1，不，＝4.3.....

这篇文章一出，有几位做科学计算的用户发给过Lady我看，说：Titan V做计算真的这么不靠谱么？

到4月份甚至有Titan V的经销商告诉我这篇文章害的这个卡的销量顿时一落千丈，本来年初此卡发布的时候，还是涌起了一波追新热潮，作为一款还支持双精度计算的卡，加上年初货量紧张的情况下，此卡的销售价格甚至炒到超过3万，文章出来以后，这卡好长一段时间无人问津。

那我们进来一起来看看这篇文章到底说了啥

文章中说：

不知道是哪位计算机科学家说的， 其实Lady要说的是：这种情况很常见.

但不是硬件问题！

如果Titan V总是能10%的计算出错, 那就跪了。

Titan-V很多人和我说过类似问题，主要发生在Titan-V + CUDA 9。

从Titan-V（计算能力7.0）开始, 和以前的卡不一样的，它是更像CPU那样的线程自由更多的执行，不是以前的GPU卡那样warp总是强制同步执行的。

很多老代码上titan-v + 9.2都会出现问题，这就是我们之前说过的兼容性问题（注意：老代码 + 老GPU卡 + CUDA9.2，没事. 但是有警告），是需要改写成CUDA 9.2的风格，否则将来Turing上还得继续出问题，会造成计算错误, 但这个确实是用户的原因。

用户需要注意新卡(计算能力7.0)和老卡是不一样的硬件环境，执行环境不同的。

如果真要像本文说的，能精确的重现的, 一批N次结果必然有10%是错误的，那就这卡不用卖了。可以让用户退回到上一代的卡(Pascal)来临时解决这个问题。

注意在cuda 9+上, 老卡+老写法+9.2, 会出现编译警告, 但运行起来没事，

而cuda 9+上, 新卡+老写法+9.2, 那就有可能要出问题，

或者是用户用新卡+新写法，但是没有正确的写对，例如需要同步warp内部的8个或者16个线程, 用户没有正确的指对目标。

所以我们建议：

用户应当尝试用新写法的，但是能否一定改对，则不一定，特别是以前老项目代码。不过迟早得适应新的，将来不可能说, Turing或者更新的卡，又改成老样子——warp内部再总是强制同步，这个概率较低的。

目前用titan-v的最好方式是当成6.1或者更低计算能力的卡用，也就是编译的时候设定为6.1或者5.0这样的参数。因为这样的参数再运行的时候, runtime或者驱动会发现里面的cubin不能运行的,因为版本不对，会强制从PTX来一次重新动态jit编译，而此时的这种编译会自动注意到兼容问题，例如ballot的时候总是用32个线程的标记要求同步，比用户自己手工改成7.0计算能力, 结果老代码没有写对的做法要强一点。

但这也只是过度用法了。

注意：计算能力7.0的GPU卡上不仅仅是这些函数(_sync结尾的)有问题，中途的很多正常代码部分(用户写的代码, 不是调用的内置函数), 都需要注意不再自动warp内部同步的问题，否则出错的地方可能会很多。

cuda 9起这种不兼容性的改动. 虽然有点很伤，但长远看, 这是进一步提升卡的性能和潜力, 减轻程序员负担必须要经历的。

其实编译过程中，这些错误都会容易发现的，因为有警告，我举个例子：

Warning: function "__ballot" was declared deprecated "__ballot() is not valid on compute_70 and above, and should be replaced with __ballot_sync(). To continue using __ballot(), specify virtual architecture compute_60 when targeting sm_70 and above ptxas warning : Instruction 'vote' without '.sync' may produce unpredictable results on sm_70 and later architectures ptxas warning : Instruction 'vote' without '.sync' is deprecated since PTX ISA version 6.0 and will be discontinued in a future PTX ISA version

我用了1个老式的ballot函数，得到了连续3个警告：

警告1: __ballot()已经过时. 在compute_70或更高上不再有效, 应当被替换成_sync版本。

想继续用它, 请在新卡上设定成compute_60之类的.

警告2: PTX警告, ptx指令vote没有sync结尾将在sm_70+上导致不可预测的结果。

警告3: PTX警告, 指令vote没有sync结尾即将在下个ptx版本被放弃支持。

连续3个警告够明显了，但是这只能对内置函数警告。其他地方需要用户自己发现(用户自己的代码)可能需要插入__syncwarp之类的手工操作。

总结:

1. 内置函数改称新写法.

2.自己的代码需要手工找出来可能会用了以前默认总是warp同步的写法的地方，然后自己加上必须的一处或者多处同步。

3.不用紧张，不是所有代码都需要改。大部分的老代码都不需要改的，可能一个项目有1-2处就改好了，甚至有的项目可能没有用到之前的warp同步的地方，那就可以不改。

希望本文对大家的CUDA编程带来帮助；

另外CUDA 阅读100天的活动，即日起暂时停更7Days，因为我们要出差....其实CUDA9的发展变化有什么功能上的改变，我们在CUDA阅读里都陆陆续续地提到了，建议大家关注我们公众号，点击历史消息，仔细阅读。