为Volta架构优化CUDA应用程序
基于NVIDIA Volta架构GPU的CUDA应用程序调优编程指南。
1. Volta调优指南
1.1. NVIDIA Volta计算架构
Volta是NVIDIA为CUDA计算应用设计的最新架构。Volta保留并扩展了由Maxwell和Pascal等前代NVIDIA架构提供的相同CUDA编程模型，遵循这些架构最佳实践的应用程序通常无需任何代码修改就能在Volta架构上获得加速。本指南总结了如何通过利用Volta架构特性对应用程序进行微调以获得额外加速的方法。1
Volta架构包含一个单一变体：GV100。关于GV100相比早期NVIDIA架构的主要改进的详细概述，请参阅题为NVIDIA Tesla V100 GPU架构：全球最先进的数据中心GPU的白皮书。
如需了解本指南中讨论的编程功能的更多详细信息，请参阅CUDA C++ Programming Guide。
1.2. CUDA最佳实践
《CUDA C++编程指南》和《CUDA C++最佳实践指南》中描述的性能准则和最佳实践适用于所有支持CUDA的GPU架构。程序员主要需要遵循这些建议以获得最佳性能。
这些指南中的高优先级建议如下：
寻找并行化顺序代码的方法，
尽量减少主机和设备之间的数据传输，
调整内核启动配置以最大化设备利用率，
确保全局内存访问是合并的，
尽可能减少对全局内存的冗余访问
避免同一warp内的线程执行过长的分支序列。
1.3. 应用兼容性
在解决本指南涵盖的具体性能调优问题之前，请参考Volta兼容性指南（适用于CUDA应用程序），以确保您的应用程序以兼容Volta架构的方式进行编译。
1.4. Volta调优
1.4.1. 流式多处理器
Volta流式多处理器(SM)相比Pascal架构提供了以下改进。
1.4.1.1. 指令调度
每个Volta SM包含4个warp调度单元。每个调度器处理一组固定的warps，并向专用的算术指令单元发出指令。指令在两个周期内执行完成，调度器每个周期可以发出独立的指令。与Pascal架构的6个周期相比，核心FMA数学运算的依赖指令发射延迟已降至4个时钟周期。因此，假设每个warp具有4路指令级并行ILP，仅需每个SM分配4个warps即可隐藏核心数学运算的执行延迟。当然，要覆盖内存事务和控制流操作更大的延迟，建议分配更多warps。
与GP100类似，GV100 SM提供64个FP32核心和32个FP64核心。GV100 SM还额外包含64个INT32核心和8个混合精度Tensor Core。GV100最多可提供84个SM。
1.4.1.2. 独立线程调度
Volta架构引入了线程束(warp)内的独立线程调度功能。这一特性使得之前无法实现的线程束内同步模式成为可能，并简化了移植CPU代码时的修改工作。然而，如果开发者基于之前硬件架构的线程束同步性2做出假设，独立线程调度可能导致实际执行代码的线程组与预期存在较大差异。
将现有代码移植到Volta架构时，需要特别注意以下三种代码模式。更多细节请参阅CUDA C++编程指南。
为避免数据损坏，使用warp内部函数（
__shfl*、__any、__all和__ballot）的应用程序应过渡到带有*_sync后缀的新安全同步版本。新的warp内部函数需要一个线程掩码，该掩码明确定义了哪些通道（warp的线程）必须参与warp内部函数。假设同一warp内的其他线程能隐式看到读写操作的应用，需要在通过全局或共享内存在线程间交换数据的步骤之间，插入新的warp级屏障同步指令
__syncwarp()。那些认为代码会锁步执行，或认为无需同步就能跨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. 张量核心运算
每个Tensor Core执行以下操作：D = AxB + C，其中A、B、C和D都是4x4矩阵。矩阵乘法输入A和B是FP16矩阵，而累加矩阵C和D可以是FP16或FP32矩阵。
当以FP32精度进行累加时，FP16乘法会生成全精度的乘积结果，随后通过FP32加法与其他中间乘积进行累加，完成4x4x4矩阵乘法运算。实际应用中，Tensor Cores会基于这些小规模运算单元来执行更大规模的二维或高维矩阵运算。
Volta张量核心在CUDA 9 C++ API中以Warp-Level Matrix Operations（线程束级矩阵运算）的形式开放。该API提供了专门的矩阵加载、矩阵乘加运算和矩阵存储操作，以便从CUDA-C++程序中高效利用张量核心。在CUDA层面，线程束级接口假设16x16大小的矩阵跨越线程束的所有32个线程。更多信息请参阅CUDA C++编程指南。
1.4.3. 内存吞吐量
1.4.3.1. 高带宽内存
GV100每个HBM2堆栈最多使用8个内存芯片，共4个堆栈，GPU显存容量最高可达32GB。相比GP100的732GB/s，更快速高效的HBM2实现方案可提供高达900GB/s的峰值内存带宽。Volta架构中新一代HBM2内存与新一代内存控制器的组合，相比Pascal GP100可提供1.5倍的实际内存带宽，在运行多数工作负载时内存带宽效率超过95%。
为了在HBM2全带宽下隐藏DRAM延迟，相比配备传统GDDR5的GPU，需要保持更多的内存访问处于执行状态。这是通过GV100中大量SM单元实现的，相比前代架构，这些SM单元通常能提升并发线程数量，从而增加飞行中的读取操作。受资源限制而只能维持低占用率的内核，可能会通过增加每个线程的并发内存访问数量而受益。
1.4.4. 协作组
Volta架构引入了独立线程调度功能，使得之前无法实现的线程束内同步模式成为可能。为了高效表达这些新模式，CUDA 9推出了协作组(Cooperative Groups)功能。这是对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访问系统内存的性能。GV100支持最多6个NVLink连接，每个连接提供高达50GB/s的双向带宽。
NVLink在现有的CUDA模型中透明运行。通过NVLink连接的端点之间的传输会自动通过NVLink路由，而不是PCIe。cudaDeviceEnablePeerAccess() API调用仍然是启用GPU之间直接传输（通过PCIe或NVLink）的必要条件。cudaDeviceCanAccessPeer()可用于确定任何一对GPU之间是否可以进行对等访问。
2. 版本历史
版本 1.0
首次公开发布
版本 1.1
新增了协作组(Cooperative Groups)章节。
更新了对CUDA C++编程指南和CUDA C++最佳实践指南的引用。
3. 通知
3.1. 注意事项
本文档仅供信息参考之用，不应视为对产品功能、状态或质量的保证。NVIDIA公司（“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对客户就本文所述产品的全部及累计责任应受产品销售条款的限制。
3.2. OpenCL
OpenCL是苹果公司的商标，经Khronos Group Inc.授权使用。
3.3. 商标
NVIDIA和NVIDIA标识是美国及其他国家NVIDIA公司的商标或注册商标。其他公司及产品名称可能是其各自关联公司的商标。