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阅读里都陆陆续续地提到了,建议大家关注我们公众号,点击历史消息,仔细阅读。

原文发布于微信公众号 - 吉浦迅科技(gpusolution)

原文发表时间:2018-08-17

本文参与腾讯云自媒体分享计划,欢迎正在阅读的你也加入,一起分享。

发表于

我来说两句

0 条评论
登录 后参与评论

相关文章

来自专栏跨界架构师

分布式系统关注点——仅需这一篇,吃透「负载均衡」妥妥的

  上一篇《分布式系统关注点——初识「高可用」》我们对「高可用」有了一个初步认识,其中认为「负载均衡」是「高可用」的核心工作。那么,本篇将通过图文并茂的方式,来...

16020
来自专栏人工智能头条

假期充电 | 10大Python开源项目推荐(Github平均star2135)

18820
来自专栏美团技术团队

顶会论文:纠删码存储系统中的投机性部分写技术

本文已被USENIX'17年度技术大会录用,此处为中文简译版。 阅读英文论文完整版请点击:Speculative Partial Writes in Erasu...

704100
来自专栏企鹅号快讯

微信跳一跳可以用 Python 刷分;macOS 再曝重大安全漏洞,或许已存在 15 年

转自:开源中国、solidot、cnBeta、腾讯科技等 ? 如果你每次都能挑到各自的正中间的话,可以 + 2 分,如果连着跳到中间会 + 4、+6、+8、+1...

234100
来自专栏机器之心

教程 | 如何使用Kubernetes GPU集群自动训练和加速深度学习?

选自GitHub 机器之心编译 参与:蒋思源、Smith、吴攀 像 Docker 这样的容器格式和 Kubernetes 之类的容器管理平台正越来越受到人们的欢...

64330
来自专栏码农分享

3.2、苏宁百万级商品爬取 思路讲解 商品爬取

如果我要得到A类别的第B页的商品我应该如何拼接符合条件的地址 我们首先分析地址,地址如下

11830
来自专栏JAVA高级架构

多研究些架构,少谈些框架(2)-- 微服务和充血模型

上篇我们聊了微服务的DDD之间的关系,很多人还是觉得很虚幻,DDD那么复杂的理论,聚合根、值对象、事件溯源,到底我们该怎么入手呢? 实际上DDD和面向对象设计、...

33550
来自专栏吉浦迅科技

DAY26:阅读性能优化策略

10440
来自专栏DHUtoBUAA

基于8211lib库对s57电子海图的解析和存储

  电子海图是为适用航海需要而绘制的包含海域地理信息和航海信息的一种数字化的专题地图,符合国际标准的电子海图数据统称为S-57电子海图。本文主要在S-57电子海...

58640
来自专栏CSDN技术头条

介绍两款大数据清洗工具——DataWrangler、Google Refine

在进行数据分析和可视化之前,经常需要先“清洗”数据。这意味着什么?可能有些词条列表里是“New York City”,而其他人写成“New York,NY”。然...

73490

扫码关注云+社区

领取腾讯云代金券