1
编者按：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 架构做一个深度剖析，同时帮助开发者了解它在实际开发中具体带来了哪些优势。
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 新特性的支持，以帮助开发者更方便地针对这些新特性编程。
搭载 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 具有比旧版 SM 更低的指令和缓存延迟，并且针对深度学习应用做了特殊优化。其主要特性如下：
● 为深度学习矩阵计算建立的新型混合精度 FP16/FP32 Tensor Core；
● 为更高性能、更低延迟而强化的 L1 高速数据缓存；
● 为简化解码和缩短指令延迟而改进的指令集；
● 更高的时钟频率和能效。
下图显示了 Volta GV100 SM 单元的基本结构。
全新的 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 的乘法、加法和存储等矩阵操作。
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
雷峰网版权文章，未经授权禁止转载。详情见转载须知。
