为 Volta 调优 CUDA 应用程序
基于 NVIDIA Volta 架构的 GPU 的 CUDA 应用程序调优编程指南。
1. Volta 调优指南
1.1. NVIDIA Volta 计算架构
Volta 是 NVIDIA 最新的 CUDA 计算应用程序架构。Volta 保留并扩展了先前 NVIDIA 架构（如 Maxwell 和 Pascal）提供的相同 CUDA 编程模型，遵循这些架构最佳实践的应用程序通常无需任何代码更改即可在 Volta 架构上看到加速。本指南总结了应用程序可以通过利用 Volta 架构特性来获得额外加速的方法。1
Volta 架构包含一个变体：GV100。题为 NVIDIA Tesla V100 GPU 架构：世界最先进的数据中心 GPU 的白皮书详细概述了 GV100 相对于早期 NVIDIA 架构的主要改进。
有关本指南中讨论的编程特性的更多详细信息，请参阅 CUDA C++ 编程指南。
1.2. CUDA 最佳实践
CUDA C++ 编程指南 和 CUDA C++ 最佳实践指南 中描述的性能指南和最佳实践适用于所有支持 CUDA 的 GPU 架构。程序员必须主要关注遵循这些建议以获得最佳性能。
这些指南中的高优先级建议如下
寻找并行化顺序代码的方法，
最大限度地减少主机和设备之间的数据传输，
调整内核启动配置以最大限度地提高设备利用率，
确保全局内存访问合并，
尽可能最大限度地减少对全局内存的冗余访问，
避免同一 Warp 内线程执行发散的长序列。
1.3. 应用程序兼容性
在解决本指南中涵盖的特定性能调优问题之前，请参阅 CUDA 应用程序的 Volta 兼容性指南，以确保您的应用程序以与 Volta 兼容的方式编译。
1.4. Volta 调优
1.4.1. 流式多处理器
Volta 流式多处理器 (SM) 提供了相对于 Pascal 的以下改进。
1.4.1.1. 指令调度
每个 Volta SM 包含 4 个 Warp 调度器单元。每个调度器处理一组静态 Warp，并向一组专用的算术指令单元发出指令。指令在两个周期内执行，调度器可以在每个周期发出独立的指令。核心 FMA 数学运算的依赖指令发出延迟减少到四个时钟周期，而 Pascal 上为六个周期。因此，假设每个 Warp 具有 4 路指令级并行性 *ILP*，则可以通过每个 SM 仅 4 个 Warp 来隐藏核心数学运算的执行延迟。当然，建议使用更多的 Warp 来覆盖内存事务和控制流操作的更大延迟。
与 GP100 类似，GV100 SM 提供 64 个 FP32 核心和 32 个 FP64 核心。GV100 SM 还额外包含 64 个 INT32 核心和 8 个混合精度张量核心。GV100 最多提供 84 个 SM。
1.4.1.2. 独立线程调度
Volta 架构在 Warp 中的线程之间引入了独立线程调度。此功能实现了以前不可用的 Warp 内同步模式，并简化了移植 CPU 代码时的代码更改。但是，如果开发人员对先前硬件架构的 Warp 同步性2 做出假设，则独立线程调度也可能导致参与执行代码的线程集与预期大相径庭。
将现有代码移植到 Volta 时，以下三种代码模式需要特别注意。有关更多详细信息，请参阅CUDA C++ 编程指南。
为了避免数据损坏，使用 Warp 内函数（
__shfl*、__any、__all和__ballot）的应用程序应过渡到新的、安全的、同步的对应项，并带有*_sync后缀。新的 Warp 内函数接受一个线程掩码，该掩码显式定义了哪些通道（Warp 的线程）必须参与 Warp 内函数。假设读取和写入对同一 Warp 中的其他线程隐式可见的应用程序需要在通过全局或共享内存在线程之间交换数据的步骤之间插入新的
__syncwarp()Warp 范围屏障同步指令。假设代码以锁步方式执行，或者来自不同线程的读取/写入在没有同步的情况下在 Warp 中可见是无效的。以某种方式使用
__syncthreads()或 PTXbar.sync（及其派生指令）的应用程序，使得线程块中某些非退出线程无法到达屏障，必须进行修改以确保所有非退出线程都到达屏障。
compute-sanitizer 提供的 racecheck 和 synccheck 工具可以帮助定位违规行为。
1.4.1.3. 占用率
每个 SM 的最大并发 Warp 数量与 Pascal 中保持相同（即 64），并且其他 影响 Warp 占用率的因素 也保持相似
寄存器文件大小为每个 SM 64k 32 位寄存器。
每个线程的最大寄存器数为 255。
每个 SM 的最大线程块数为 32。
每个 SM 的共享内存容量为 96KB，与 GP104 相似，比 GP100 增加了 50%。
总的来说，开发人员可以期望在 Pascal 上获得类似的占用率，而无需更改其应用程序。
1.4.1.4. 整数运算
与 Pascal GPU 不同，GV100 SM 包括专用的 FP32 和 INT32 核心。这使得 FP32 和 INT32 操作能够同时执行。应用程序现在可以将指针运算与浮点计算交错进行。例如，流水线循环的每次迭代都可以更新地址并加载下一次迭代的数据，同时以完整的 FP32 吞吐量处理当前迭代。
1.4.2. 张量核心运算
每个张量核心执行以下运算：D = AxB + C，其中 A、B、C 和 D 是 4x4 矩阵。矩阵乘法输入 A 和 B 是 FP16 矩阵，而累加矩阵 C 和 D 可以是 FP16 或 FP32 矩阵。
当以 FP32 累加时，FP16 乘法会产生全精度乘积，然后使用 FP32 加法与 4x4x4 矩阵乘法的其他中间乘积累加。实际上，张量核心用于执行更大的 2D 或更高维度的矩阵运算，这些运算由这些较小的元素构建而成。
Volta 张量核心在 CUDA 9 C++ API 中作为 Warp 级矩阵运算公开。API 公开了专门的矩阵加载、矩阵乘法和累加以及矩阵存储操作，以便从 CUDA-C++ 程序中高效使用张量核心。在 CUDA 级别，Warp 级接口假定 16x16 大小的矩阵跨越 Warp 的所有 32 个线程。有关更多信息，请参阅CUDA C++ 编程指南。
1.4.3. 内存吞吐量
1.4.3.1. 高带宽内存
GV100 每个 HBM2 堆栈最多使用八个内存芯片和四个堆栈，最大 GPU 内存为 32 GB。与 GP100 的 732 GB/s 相比，更快更高效的 HBM2 实现可提供高达 900 GB/s 的峰值内存带宽。与 Pascal GP100 相比，Volta 中新一代 HBM2 内存和新一代内存控制器的结合提供了 1.5 倍的交付内存带宽，并且在运行许多工作负载时内存带宽效率超过 95%。
为了在完整的 HBM2 带宽下隐藏 DRAM 延迟，与配备传统 GDDR5 的 GPU 相比，必须保持更多的内存访问处于活动状态。这通过 GV100 中大量的 SM 来实现，这通常会提高并发线程的数量，从而提高与以前的架构相比的正在进行的读取次数。资源受限的内核（受限于低占用率）可以受益于增加每个线程的并发内存访问次数。
1.4.4. 协作组
Volta 架构引入了独立线程调度，这使得以前不可能实现的 Warp 内同步模式成为可能。为了有效地表达这些新模式，CUDA 9 引入了协作组。这是 CUDA 编程模型的扩展，用于组织通信线程组。协作组允许开发人员表达线程通信的粒度，帮助他们表达更丰富、更高效的并行分解。有关更多信息，请参阅CUDA C++ 编程指南。
1.4.5. 多进程服务
与以前的架构相比，Volta 多进程服务在性能和稳健性方面都得到了显着改进。用于先前架构 MPS 的中介软件调度器已被 GPU 内的硬件加速单元取代。MPS 客户端现在直接将任务提交到 GPU 工作队列，从而显着减少了提交延迟并提高了总吞吐量。MPS 客户端数量的限制也增加了 3 倍，达到 48 个。Volta MPS 还为每个客户端提供隔离的地址空间，3 并扩展了对 MPS 应用程序的统一内存支持。
Volta MPS 还为客户端提供控制，以将每个客户端限制为 GPU 执行资源的一部分。开发人员可以使用此功能来减少或消除队首阻塞，在这种阻塞中，来自一个 MPS 客户端的工作会淹没 GPU 执行资源并阻止其他客户端取得进展，从而提高系统的平均延迟和抖动。
1.4.6. NVLink 互连
NVLink 是 NVIDIA 的高速数据互连。NVLink 可用于显着提高 GPU 到 GPU 通信和 GPU 访问系统内存的性能。GV100 最多支持六个 NVLink 连接，每个连接的承载能力高达 50 GB/s 的双向带宽。
NVLink 在现有 CUDA 模型中透明地运行。NVLink 连接的端点之间的传输会自动通过 NVLink 而不是 PCIe 路由。cudaDeviceEnablePeerAccess() API 调用仍然是启用 GPU 之间直接传输（通过 PCIe 或 NVLink）所必需的。cudaDeviceCanAccessPeer() 可用于确定任何一对 GPU 之间是否可以进行对等访问。
2. 修订历史
版本 1.0
首次公开发布
版本 1.1
添加了协作组部分。
更新了对CUDA C++ 编程指南和CUDA C++ 最佳实践指南的引用。
3. 通知
3.1. 通知
本文档仅供参考，不应被视为对产品的特定功能、条件或质量的保证。NVIDIA Corporation（“NVIDIA”）不对本文档中包含的信息的准确性或完整性做出任何明示或暗示的陈述或保证，并且对本文档中包含的任何错误不承担任何责任。NVIDIA 对因使用此类信息或因使用此类信息而可能导致的侵犯第三方专利或其他权利的行为不承担任何责任。本文档不构成开发、发布或交付任何材料（如下定义）、代码或功能的承诺。
NVIDIA 保留随时对此文档进行更正、修改、增强、改进和任何其他更改的权利，恕不另行通知。
客户在下订单前应获取最新的相关信息，并应验证此类信息是否为最新且完整。
NVIDIA 产品根据订单确认时提供的 NVIDIA 标准销售条款和条件进行销售，除非 NVIDIA 和客户的授权代表签署的个别销售协议（“销售条款”）另有约定。NVIDIA 特此明确反对将任何客户通用条款和条件应用于购买本文档中引用的 NVIDIA 产品。本文档不直接或间接地构成任何合同义务。
NVIDIA 产品并非设计、授权或保证适用于医疗、军事、航空、航天或生命支持设备，也不适用于 NVIDIA 产品的故障或故障可能合理预期会导致人身伤害、死亡或财产或环境损害的应用。对于在上述设备或应用中包含和/或使用 NVIDIA 产品，NVIDIA 不承担任何责任，因此，此类包含和/或使用风险由客户自行承担。
NVIDIA 不保证或声明基于本文档的产品将适用于任何特定用途。NVIDIA 不一定对每个产品的所有参数进行测试。客户全权负责评估和确定本文档中包含的任何信息的适用性，确保产品适合且适用于客户计划的应用，并为该应用执行必要的测试，以避免应用或产品的默认设置。客户产品设计中的缺陷可能会影响 NVIDIA 产品的质量和可靠性，并可能导致超出本文档中包含的附加或不同条件和/或要求。NVIDIA 对可能基于或归因于以下原因的任何默认设置、损坏、成本或问题不承担任何责任：(i) 以任何违反本文档的方式使用 NVIDIA 产品，或 (ii) 客户产品设计。
本文档未授予 NVIDIA 专利权、版权或其他 NVIDIA 知识产权下的任何明示或暗示的许可。NVIDIA 发布的关于第三方产品或服务的信息不构成 NVIDIA 授予使用此类产品或服务的许可，也不构成对其的保证或认可。使用此类信息可能需要获得第三方在其专利或其他知识产权下的许可，或者获得 NVIDIA 在 NVIDIA 专利或其他知识产权下的许可。
只有在事先获得 NVIDIA 书面批准的情况下，才允许复制本文档中的信息，复制时不得进行更改，且必须完全符合所有适用的出口法律和法规，并附带所有相关的条件、限制和通知。
本文档和所有 NVIDIA 设计规范、参考板、文件、图纸、诊断程序、列表和其他文档（统称为“材料”）均按“原样”提供。NVIDIA 不对材料做出任何明示、暗示、法定或其他形式的保证，并明确声明不承担任何关于非侵权性、适销性和特定用途适用性的暗示保证。在法律未禁止的范围内，在任何情况下，NVIDIA 均不对因使用本文档而引起的任何损害（包括但不限于任何直接、间接、特殊、附带、惩罚性或后果性损害，无论如何造成且无论责任理论如何）承担责任，即使 NVIDIA 已被告知可能发生此类损害。尽管客户可能因任何原因遭受任何损害，但 NVIDIA 对本文所述产品的客户的总体和累积责任应根据产品的销售条款进行限制。
3.2. OpenCL
OpenCL 是 Apple Inc. 的商标，已获得 Khronos Group Inc. 的许可使用。
3.3. 商标
NVIDIA 和 NVIDIA 徽标是 NVIDIA Corporation 在美国和其他国家/地区的商标或注册商标。其他公司和产品名称可能是与其关联的各自公司的商标。