![]() |
一文详解英伟达刚发布的 Tesla V100 究竟牛在哪? |
2017年5月18日
【
转载
】 编辑:
浏览次数:
|
|
编者按:5 月 11 日,在加州圣何塞举办的的 2017 年度 GPU 技术大会上,英伟达公布了 Tesla V100,号称史上最强的 GPU 加快器 。公布之后,英伟达第一
工夫在官方开发者博客放出一篇博文,
详尽
分析了包括 Tesla V100,GV100 GPU,Tensor Core,以及 Volta 架构等在内的各项新
特点/新产品的技术内涵,雷锋网编译如下 。
家喻户晓,当前无论是语音 鉴别,还是 虚构个人助理的训练;路线探测,还是自动驾驶系统的研发,在这些人工智能领域,数据科学家们正在面对越来越复杂的 AI 挑战 。而为了更好地实现这些颇具 将来感的 壮大 性能,就必须在 实际中引入一些指数级的更加复杂的深度学习模型 。
另一方面,HPC(高性能计算)在现代科学探究中向来起着至关主要的作用 。无论是预测天气,新药物的探究,或是探究 将来能源,科研人员天天都需要利用大型计算系统对 事实世界做各种各样的仿真和预测 。而通过引入 AI 技术,HPC 就 可以卓著 晋升科研人员进行大数据 综合的效率,并得到一些此前通过传统的仿真和预测 步骤 无奈得到新 论断 。
为了进一步推进 HPC 和 AI 领域的 有关进展,英伟达最近公布了新一代 Tesla V100 GPU 加快器 。它基于最新的 NVIDIA Volta GV100 GPU 平台和各种 打破性技术创新, 可 认为各种超级计算系统提供一个 壮大的运算平台, 不管在以科学仿真为主要 目的的计算科学领域,还是在以洞悉数据神秘为 指标的数据科学领域,Tesla V100 都能为 有关 利用提供 壮大的算力 支撑 。
下面,我们会通过这篇博客对 Tesla V100 的核心:Volta 架构做一个深度 分析,同时协助开发者了解它在实际开发中具体带来了哪些优势 。
Tesla V100:AI 计算和 HPC 的源动力
NVIDIA Tesla V100 是当前世界上最高性能的并行 解决器,专门用于 解决需要 壮大计算 威力 支撑的密集型 HPC、AI、和图形 解决 使命 。
Tesla V100 加快器的核心是 GV100 GPU 解决器 。基于台积电专门为 NVIDIA 设计的最新 12nm FFN 高精度制程封装技术,GV100 在 815 平方毫米的芯片尺寸中,内部集成了高达 211 亿个晶体管 构造 。相较于上一代产品,也便是 Pascal 系列 GPU,GV100 不单在计算性能上有了长足的 普及,同时还添加了许多令人眼前一亮的新 特点 。包括进一步精简的 GPU 编程和 利用部署流程,以及针对 GPU 资源利用状况的深度优化 。其 后果是,GV100 在提供 壮大计算性能的同时还十分省电,下图显示了 Tesla V100 加快器和上代产品 Tesla P100 加快器在 ResNet-50 模型训练和推理中的性能对照, 可以看到最新的 V100 要远超上一代 P100 。
Tesla V100 的 要害 特点总结如下:
● 针对深度学习优化的流式多 解决器(SM)架构 。作为 GPU 解决器的核心组件,在 Volta 架构中 NVIDIA 再一次设计了 SM,相比之前的 Pascal 架构而言,这一代 SM 普及了约 50% 的能效,在同样的功率 规模内 可以大幅 晋升 FP32(单精度浮点)和 FP64(双精度浮点)的运算性能 。专为深度学习设计的崭新 Tensor Core 在模型训练场景中,最高 可以达到 12 倍速的 TFLOP(每秒万亿次浮点运算) 。另外,由于崭新的 SM 架构对整型和浮点型数据采取了 彼此独立且并行的数据通路, 因此在一般计算和寻址计算等混合场景下也能输出不错的效率 。Volta 架构新的独立线程调度 性能还 可以实现并行线程中间的细粒度同步和 合作 。最终,一个新组合的 L1 高速数据缓存和共享内存子系统也卓著 普及了性能,同时大大简化了开发者的编程步骤 。
● 第二代 NVLink 。第二代 NVIDIA NVLink 高速互连技术为多 GPU 和多 GPU/CPU系统配置提供了更高的带宽,更多的衔接和更强的可 扩大性 。GV100 GPU 最多 支撑 6 个 NVLink 链路,每个 25 GB/s,总共 300 GB/s 。NVLink 还 支撑基于 IBM Power 9 CPU 服务器的 CPU 操纵和高速缓存 统一性 性能 。另外,新公布的 NVIDIA DGX-1V 超级 AI 计算机也 使用了 NVLink 技术为超 快捷的深度学习模型训练提供了更强的 扩大性 。
● HBM2 内存:更快,更高效 。Volta 高度优化的 16GB HBM2 内存子系统可提供高达 900 GB/s 的峰值内存带宽 。相比上一代 Pascal GP100,来自三星的新一代 HBM2 内存与 Volta 的新一代内存操纵器相 联合,带宽 晋升 1.5 倍,而且在性能 体现上也超过了 95% 的工作负载 。
● Volta 多 解决器服务(Multi-Process Service,MPS) 。Volta MPS 是 Volta GV100 架构的一项新 特点, 可以提供 CUDA MPS 服务器 要害组件的硬件加快 性能,从而在共享 GPU 的多计算 使命场景中卓著 晋升计算性能、隔离性和服务 品质(QoS) 。Volta MPS 还将 MPS 支撑的客户端最大数量从 Pascal 时代的 16 个添加到 48 个 。
● 加强的统一内存和地址转换服务 。Volta GV100 中的 GV100 统一内存技术实现了一个新的 拜访计数器,该计数器 可以依据每个 解决器的 拜访频率准确调整内存页的寻址,从而大大 晋升了 解决器中间共享内存的 使用效率 。另外,在 IBM Power 平台上,新的地址转换服务(Address Translation Services,ATS)还同意 GPU 直接 拜访 CPU 的存储页表 。
● Cooperative Groups( 合作组)和新的 Cooperative Launch API( 合作启动 API) 。Cooperative Groups 是在 CUDA 9 中引入的一种新的编程模型,用于组织通讯线程组 。Cooperative Groups 同意开发人员 抒发线程中间的沟通粒度,协助他们更 丰硕、更有效地进行并行分解(decompositions) 。Kepler 系列以来,所有的 NVIDIA GPU 都 支撑 根本 Cooperative Groups 特点 。Pascal 和 Volta 系列还 支撑新的 Cooperative Launch API,通过该 API 可以实现 CUDA 线程块中间的同步 。另外 Volta 还添加了对新的同步模式的 支撑 。
● 最大性能和最高效率两种模式 。顾名思义,在最高性能模式下,Tesla V100 极速器将无 制约地运行,达到 300W 的 TDP(热设计功率)级别,以满足那些需要最快计算速度和最高数据吞吐量的 利用需要 。而最高效率模式则同意数据 核心治理员调整 Tesla V100 的功耗水平,以每瓦特最佳的能耗 体现输出算力 。而且,Tesla V100 还 支撑在所有 GPU 中设置上限功率,在大大减低功耗的同时,最大限度地满足机架的性能要求 。
● 针对 Volta 优化的软件 。各种新版本的深度学习框架(包括 Caffe2,MXNet,CNTK,TensorFlow 等)都 可以利用 Volta 大大缩小模型训练 工夫,同时 晋升多节点训练的性能 。各种 Volta 优化版本的 GPU 加快库(包括 cuDNN,cuBLAS 和 TensorRT 等)也都 可以在 Volta GV100 各项新 特点的 支撑下,为深度学习和 HPC 利用提供更好的性能 支撑 。此外,NVIDIA CUDA Toolkit 9.0 版也加入了新的 API 和对 Volta 新 特点的 支撑,以协助开发者更容易地针对这些新 特点编程 。
GV100 GPU 硬件架构
搭载 Volta GV100 GPU 的 NVIDIA Tesla V100 加快器是当今世界上性能最强的并行计算 解决器 。其中,GV100 GPU 存在一系列的硬件创新,为深度学习算法和框架、HPC 系统和 利用程序,均提供了 壮大的算力 支撑 。其中在 HPC 领域的性能 体现如下图所示,在各种 HPC 使命中,Tesla V100 均匀比 Tesla P100 快 1.5 倍(基于 Tesla V100 原型卡) 。
Tesla V100 占有业界率先的浮点和整型运算性能,峰值运算性能如下(基于 GPU Boost 时钟频率):
● 双精度浮点(FP64)运算性能:7.5 TFLOP/s;
● 单精度(FP32)运算性能:15 TFLOP/s;
● 混合精度矩阵乘法和累加:120 Tensor TFLOP/s 。
和之前的 Pascal GP100 一样,GV100 也由许多图形 解决集群(Graphics Processing Cluster,GPC)、纹理 解决集群(Texture Processing Cluster,TPC)、流式多 解决器(Streaming Multiprocessor,SM)以及内存操纵器构成 。一个 完全的 GV100 GPU 由 6 个 GPC、84 个 Volta SM、42 个 TPC(每个 TPC 包括 2 个 SM)和 8 个 512 位的内存操纵器(共 4096 位) 。其中,每个 SM 有 64 个 FP32 核、64 个 INT32 核、32 个 FP64 核与 8 个崭新的 Tensor Core 。同时,每个 SM 也包括了 4 个纹理 解决单元(texture units) 。
更具体地说,一个 完全版 Volta GV100 中总共包括了 5376 个 FP32 核、5376 个 INT32 核、2688 个 FP64 核、672 个 Tensor Core 以及 336 个纹理单元 。每个内存操纵器都链接一个 768 KB 的 2 级缓存,每个 HBM2 DRAM 堆栈都由一对内存操纵器操纵 。整体上,GV100 总共包括 6144KB 的二级缓存 。下图 展示了带有 84 个 SM 单元的 完全版 Volta GV100,需要 留神的是,不同的产品可能 存在不同的配置, 比方Tesla V100 就惟独 80 个 SM 。
下表 展示了 Tesla V100 与过去五年历代 Tesla 系列加快器的参数对照 。
Volta SM(流式多 解决器)
为了提供更高的性能,Volta SM 存在比旧版 SM 更低的指令和缓存延迟,而且针对深度学习 利用做了特别优化 。其主要 特点如下:
● 为深度学习矩阵计算 构建的新型混合精度 FP16/FP32 Tensor Core;
● 为更高性能、更低延迟而强化的 L1 高速数据缓存;
● 为简化解码和缩小指令延迟而改良的指令集;
● 更高的时钟频率和能效 。
下图显示了 Volta GV100 SM 单元的 根本 构造 。
Tensor Core:既是运算指令也是数据 格局
崭新的 Tensor Core 是 Volta GV100 架构中最主要的一项新 特点,在训练超大型神经网络模型时,它 可 认为系统提供强劲的运算性能 。Tesla V100 的 Tensor Core 可 认为深度学习 有关的模型训练和判断 利用提供高达 120 TFLOPS 的浮点张量计算 。具体来说,在深度学习的模型训练方面,相比于 P100 上的 FP32 操作,崭新的 Tensor Core 可以在 Tesla V100 上实现最高 12 倍速的峰值 TFLOPS 。而在深度学习的判断方面,相比于 P100 上的 FP16 操作,则 可以实现最高 6 倍速的峰值 TFLOPS 。Tesla V100 GPU 一共包括 640 个 Tensor Core,每个流式多 解决器(SM)包括 8 个 。
家喻户晓,矩阵乘法运算是神经网络训练的核心,在深度神经网络的每个衔接层中,输入矩阵都要乘以权重以 获得下一层的输入 。如下图所示,相比于上一代 Pascal 架构的 GP100,Tesla V100 中的 Tensor Core 把矩阵乘法运算的性能 晋升了至少 9 倍 。
如本节小 题目所述,Tensor Core 不只不过一个崭新的高效指令集,还是一种数据运算 格局 。
在刚公布的 Volta 架构中,每个 Tensor Core 都包括一个 4x4x4 的矩阵 解决队列,来 实现神经网络 构造中最常见的 D=AxB+C 运算 。其中 A、B、C、D 是 4 个 4×4 的矩阵, 因此被称为 4x4x4 。如下图所示,输入 A、B 是指 FP16 的矩阵,而矩阵 C 和 D 可以是 FP16,也 可以是 FP32 。
依照设计,Tensor Core 在每个时钟频率 可以执行高达 64 次 FMA 混合精度浮点操作,也便是两个 FP16 输入的乘积,再外加一个 FP32 。而由于每个 SM 单元都包括 8 个 Tensor Core, 因此总体上每个时钟 可以执行 1024 次浮点运算 。这使得在 Volta 架构中,每个 SM 单元的深度学习 利用吞吐量相比 标准 FP32 操作的 Pascal GP100 大幅 晋升了 8 倍,与Pascal P100 GPU相比,Volta V100 GPU的吞吐量总共 普及了 12 倍 。下图 展示了一个 标准的 Volta GV100 Tensor Core 流程 。
在程序执行期间,多个 Tensor Cores 通过 warp 单元协同工作 。warp 中的线程同时还提供了 可以由 Tensor Cores 解决的更大的 16x16x16 矩阵运算 。CUDA 将这些操作作为 Warp-Level 级的矩阵运算在 CUDA C++ API 中公开 。通过 CUDA C++ 编程,开发者 可以灵便 使用这些开放 API 实现基于 Tensor Cores 的乘法、加法和存储等矩阵操作 。
加强的 L1 高速数据缓存和共享内存
Volta SM 的 L1 高速数据缓存和共享内存子系统 彼此 联合,显着 普及了性能,同时也大大简化了开发者的编程步骤、以及达到或接近最优系统性能的系统调试成本 。
值得强调的是,Volta 架构将数据高速缓存和共享内存 性能组合到单个内存块中的做法,在整体上为两 品种型的内存 拜访均提供了最佳的性能 。组合后的内存容量达到了 128 KB/SM,比老版的 GP100 高速缓存大 7 倍以上,而且所有这些都 可以配置为不共享的独享 cache 块 。另外,纹理 解决单元也 可以 使用这些 cache 。例如,假如共享内存被设置为 64KB,则纹理和加载/存储操作就 可以 使用 L1 中 残余的 64 KB 容量 。
总体上,通过和共享内存 彼此组合的独创性 模式,使得 Volta GV100 L1 高速缓存 存在比过去 NVIDIA GPU 的 L1 高速缓存更低的延迟和更高的带宽 。一方面作为流数据的高吞吐量管道 施展作用,另一方面也 可 认为复用度很高的数据提供高带宽和低延迟的精准 拜访 。
下图显示了 Volta 和 Pascal 的 L1 缓存性能对照 。
计算 威力
GV100 GPU 支撑英伟达崭新的 Compute Capability 7.0 。下表显示了 NVIDIA GPU 不同架构中间的计算 威力对照 。
独立的线程调度
Volta 架构相较之前的 NVIDIA GPU 卓著减低了编程难度,消费者 可以更 专一于将各种多样的 利用产品化 。Volta GV100 是第一个 支撑独立线程调度的 GPU,也便是说,在程序中的不同线程 可以更 详尽地同步和 合作 。Volta 的一个主要设计 指标便是减低程序在 GPU 上运行所需的开发成本,以及线程中间灵便的共享机制,最终使得并行计算更为高效 。
此前的单指令多线程模式(SIMT MODELS)
在 Pascal 和之前的 GPU 中, 可以执行由 32 个线程构成的 group,在 SIMT 术语里也被称为 warps 。在 Pascal 的 warp 里,这 32 个线程 使用同一个程序计数器, 而后由一个激活掩码(active mask) 表明 warp 里的哪些线程是有效的 。这 象征着不同的执行路径里有些线程是“非激活态”的,下图给出了一个 warp 里不同分支的顺序执行过程 。在程序中,原始的掩码会先被 保留起来,直到 warps 执行 完毕,线程再度收敛,掩码会被 复原,程序再接着执行 。
从 性质上来说,Pascal 的 SIMT 模式通过削减跟踪线程状态所需的资源和 踊跃地 复原线程将并行效率最大化 。这种对整个 warps 进行线程状态跟踪的模式,其实 象征着当程序浮现并行分支时,warps 内部实际上是顺序执行的,这里已经 损失了并行的 意思,直到并行分支的 完毕 。也便是说,不同 warp 里的线程确 着实并行执行,但同一 warp 里的分 干线程却在未 复原之前顺序执行,它们中间 无奈交互信息和共享数据 。
举个例子来说,要求数据精准共享的那些算法,在不同的线程 拜访被锁和互斥机制 掩护的数据块时,由于不确定遇到的线程是来自哪个 warp,所以很容易招致死锁 。 因此,在 Pascal 和之前的 GPU 里,开发者们只能幸免细粒度同步,或者 使用那些不依赖锁,或明确 划分 warp 的算法 。
Volta 架构的单指令多线程模式
Volta 通过在所有线程间( 无论是哪个 warp 的) 施行 等同级别的并发性解决了这一问题,对每个线程,包括程序计数器和调用栈,Volta 都 保护同一个执行状态,如下图所示 。
Volta 的独立线程调配机制同意 GPU 将执行权限 斗争于任何一个线程,这样做使线程的执行效率更高,同时也让线程间的数据共享更 正当 。为了最大化并行效率,Volta 有一个调度优化器, 可以决定如何对同一个 warp 里的有效线程进行分组,并一同送到 SIMT 单元 。这不只 维持了在 NVIDIA 之前的 GPU 里较高的 SIMT 吞吐量,而且灵便性更高:现在,线程 可以在 sub-warp 级别上分支和 复原,而且,Volta 仍将那些执行 雷同代码的线程分组在一同,让他们并行运行 。
下图 展示了 Volta 多线程模式的一个样例 。这个程序里的 if/else 分支现在 可以依照时序被 间隔开来,如图12所示 。 可以看到,执行过程依旧是 SIMT 的,在任意一个时钟周期,和之前一样,同一个 warp 里的所有有效线程,CUDA 核执行的是同样的指令,这样依旧 可以 维持之前架构中的执行效率 。重点是,Volta 的这种独立调度 威力, 可以让程序员有机会用更加自然的 模式开发出复杂且 详尽的算法和数据 构造 。 固然调度器 支撑线程执行的独立性,但它依旧会优化那些非同步的代码段,在确保线程收敛的同时,最大限度地 晋升 SIMT 的高效性 。
另外,上图中还有一个 乏味的 景象:Z 在所有的线程中都不是同一时刻执行的 。这是由于 Z 可能会输出其它分支 历程需要的数据,在这种状况下,强制进行收敛并不安全 。但在之前的架构中,一般认为 A,B,X,Y 并不包括同步性操作, 因此调度器会认定在 Z 上收敛是安全的 。
在这种状况下,程序 可以调用新的 CUDA 9 中的 warp 同步函数 __syncwarp() 来强制进行线程收敛,如下图所示 。这 时候 干线程可能并不会同步执行 Z,然而通过调用 __syncwarp() 函数,同一个 warp 里的这些线程的所有执行路径将会在执行到 Z 语句之前 齐备 。 类似的,在执行 Z 之前,假如调用一下 __syncwarp() 函数,则程序将会在执行 Z 之前强制收敛 。假如开发者能提前确保这种操作的安全性,无疑这会在 定然程度上 晋升 SIMT 的执行效率 。
Starvation-Free 算法
Starvation-free 算法是独立线程调度机制的一个主要模式,具体是指:在并发计算中, 惟独系统确保所有线程 存在对竞争性资源的 适当 拜访权,就 可以 保障其正确执行 。例如,假如尝试猎取互斥锁(mutex)的线程最终 顺利 获得了该锁,就 可以在 starvation-free 算法中 使用互斥锁(或一般锁) 。在不 支撑 starvation-free 算法的系统中,可能会浮现一个或多个线程 反复猎取和 开释互斥锁的状况,这就有可能造成 其余线程始终 无奈 顺利猎取互斥锁的问题 。
下面看一个对于 Volta 独立线程调度的实例:在多线程 利用程序中将节点插入双向链表 。
__device__ void insert_after(Node *a, Node *b)
{
Node *c;
lock(a); lock(a->next);
c = a->next;
a->next = b;
b->prev = a;
b->next = c;
c->prev = b;
unlock(c); unlock(a);
}
在这个例子中,每个双向链表的元素至少含有 3 个 部分:一个后向指针,一个前向指针,以及一个 lock(惟独 owner 才有权限更新结点) 。下图 展示了在 A 和 C 中间插入 B 结点的过程 。
Volta 这种独立线程调度机制 可以确保 即便线程 T0 当前锁住了结点 A,同一个 warp 里的另一个线程 T1 依旧 可以 顺利地等到其解锁,而不影响 T0 的执行 。不过,值得 留神的丝毫是,由于同一个 warp 下的有效线程是一同执行的,所以等解锁的线程可能会让锁住的线程性能减低 。
同样需要 重视的是,如此例中这种针对每个结点上锁的用法对 GPU 的性能影响至关主要 。传统上,双向链接表的 缔造可能会用粗粒度 lock(对应前面提到的细粒度 lock),粗粒度 lock 会独占整个 构造(所有上锁),而不是对每一个结点分别予以 掩护 。由于线程间对 lock 的 抢夺, 因此这种 步骤可能会招致多线程代码的性能 降落(Volta 架构最多同意高达 163,840 个并发线程) 。这时 可以尝试在每个节点采纳细粒度 lock 的 步骤,这样除了在某些特定节点的插入操作之外,大型列表中 均匀每个节点的 lock 竞争效应就会大大减低 。
上述这种具备细粒度 lock 的双向链接表只不过个十分 容易的例子,我们想通过这个例子 转达的信息是:通过独立的线程调度机制,开发者们 可以用最自然的 模式在 NVIDIA GPU 上实现 相熟的算法和数据 构造 。
总结
NVIDIA Tesla V100 无疑是当前世界上最先进的数据 核心 GPU,专门用于 解决需要 壮大计算 威力 支撑的密集型 HPC、AI、和图形 解决 使命 。凭借最先进的 NVIDIA Volta 架构 支撑,Tesla V100 可以在单片 GPU 中提供 100 个 CPU 的运算性能,这使得数据科学家、探究人员和工程师们得以 应答曾经被认为是不可能的 挑战 。
搭载 640 个 Tensor cores,使得 Tesla V100 成为了当前世界上第一款 打破 100 TFLOPS 算力大关的深度学习 GPU 产品 。再外加新一代 NVIDIA NVLink 技术高达 300 GB/s 的衔接 威力, 事实场景中消费者 彻底 可以将多个 V100 GPU 组合起来搭建一个 壮大的深度学习运算 核心 。这样,曾经需要数周 工夫的 AI 模型现在 可以在几天之内训练 实现 。而随着训练 工夫的大幅度缩小, 将来所有的 事实问题兴许都将被 AI 解决 。
起源:英伟达开发者博客
雷锋网(公众号:雷锋网) 有关浏览:
2小时、5大AI 新品、英伟达股价暴涨17%,GTC大会上黄仁勋都讲了些啥?(内附PPT) | GTC 2017
GTC大会第二日卖点:NVIDIA将推出多消费者VR系统, 方案 造就100000名开发人员 | GTC 2017