为Turing架构优化CUDA应用
基于NVIDIA Turing架构的GPU调优CUDA应用程序编程指南。
1. 图灵调优指南
1.1. NVIDIA Turing计算架构
Turing是NVIDIA为CUDA计算应用推出的最新架构。Turing保留并扩展了Pascal和Volta等前代NVIDIA架构提供的相同CUDA编程模型，遵循这些架构最佳实践的应用程序通常无需修改代码即可在Turing架构上获得加速。本指南总结了如何通过利用Turing架构特性对应用程序进行微调以获得额外加速。1
如需了解本指南中讨论的编程功能的更多详细信息，请参阅CUDA C++ Programming Guide。
1.2. CUDA最佳实践
《CUDA C++编程指南》和《CUDA C++最佳实践指南》中描述的性能准则和最佳实践适用于所有支持CUDA的GPU架构。程序员主要需要遵循这些建议以获得最佳性能。
这些指南中的高优先级建议如下：
寻找并行化顺序代码的方法，
尽量减少主机和设备之间的数据传输，
调整内核启动配置以最大化设备利用率，
确保全局内存访问是合并的，
尽可能减少对全局内存的冗余访问
避免同一warp内的线程执行过长的分支序列。
1.3. 应用兼容性
在解决本指南涵盖的具体性能调优问题之前，请参阅Turing Compatibility Guide for CUDA Applications以确保您的应用程序以兼容Turing架构的方式进行编译。
1.4. 图灵调优
1.4.1. 流式多处理器
图灵流式多处理器(SM)基于与Volta相同的主要架构(7.x)，相比Pascal架构提供了类似的改进。
1.4.1.1. 指令调度
每个Turing SM包含4个warp调度单元。每个调度器处理一组固定的warps，并向专用的算术指令单元分发指令。指令执行需要两个时钟周期，调度器每个周期都能分发独立的指令。与Pascal架构的6个周期相比，核心FMA数学运算的依赖指令分发延迟（与Volta架构相同）为4个时钟周期。因此，假设每个warp具有4路指令级并行ILP，则每个SM仅需4个warps即可隐藏核心数学运算的执行延迟；若没有任何指令级并行，则需要每个SM配备16个warps才能实现相同效果。
与Volta类似，Turing SM提供64个FP32核心、64个INT32核心和8个改进的混合精度Tensor核心。相比Volta，Turing的双精度吞吐量较低，仅配备2个FP64核心。
1.4.1.2. 独立线程调度
图灵架构采用了与Volta相同的独立线程调度技术。这项技术实现了以往无法实现的warp内部同步模式，并在移植CPU代码时简化了代码修改。然而，如果开发者基于旧硬件架构的warp同步性2做出假设，独立线程调度可能导致实际执行代码的线程组与预期出现较大差异。
将现有代码移植到Volta或Turing架构时，需要特别注意以下三种代码模式。更多详细信息请参阅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. 占用率
在Turing架构上，每个SM的最大并发warp数量为32（而Volta架构为64）。其他影响warp占用率的因素基本保持相似：
每个SM的寄存器文件大小为64k个32位寄存器。
每个线程的最大寄存器数量为255。
每个SM(流式多处理器)的最大线程块数量为16。
每个SM的共享内存容量为64KB。
总体而言，开发者无需修改应用程序即可获得与Pascal或Volta架构相似的占用率。
1.4.1.4. 整数运算
与Volta类似，Turing SM架构也配备了专用的FP32和INT32计算核心。这使得FP32与INT32运算能够同步执行。应用程序可以将指针运算与浮点计算交错进行。例如，在流水线循环的每次迭代中，可以在保持FP32全吞吐量处理当前迭代的同时，为下一次迭代更新地址并加载数据。
1.4.2. 张量核心运算
Volta架构引入了Tensor Core来加速混合精度浮点数据的矩阵乘法运算。Turing架构进一步增加了对整数矩阵乘法运算的加速支持。这些张量核心在CUDA 10 C++ API中以Warp-Level Matrix Operations的形式提供。该API专门提供了矩阵加载、矩阵乘加运算和矩阵存储操作，其中每个warp处理一个小型矩阵片段，使得开发者能够高效地从CUDA-C++程序中使用Tensor Core。实际上，Tensor Core被用于执行更大规模的二维或更高维矩阵运算，这些运算都是由这些小型矩阵片段构建而成的。
每个张量核心执行矩阵乘加运算：D = A x B + C。张量核心支持半精度矩阵乘法，其中矩阵乘法输入A和B为FP16矩阵，而累加矩阵C和D可以是FP16或FP32矩阵。当使用FP32进行累加时，FP16乘法会产生全精度乘积，随后通过FP32加法进行累加。CUDA 10支持多种片段尺寸（16x16x16、32x8x16和8x32x16），以便在Volta或Turing架构上使用FP16输入来利用张量核心。
任何为Volta编译的二进制文件都可以在Turing上运行，但使用Tensor Core的Volta二进制文件只能达到Turing Tensor Core峰值性能的一半。专门为Turing重新编译二进制文件可使其达到峰值性能。更多信息请参阅《Turing兼容性指南》。
图灵架构的Tensor Core支持整数矩阵乘法运算，可处理8位、4位和1位整数输入，并实现32位整数累加。当处理8位输入时，CUDA支持的片段尺寸为16x16x16、32x8x16和8x32x16。对于亚字节运算，4位输入可用的片段尺寸为8x8x32，1位输入则为8x8x128。
更多信息请参阅CUDA C++编程指南。
1.4.3. 内存吞吐量
2. 版本历史
版本 1.0
首次公开发布
版本 1.1
更新了对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公司的商标或注册商标。其他公司及产品名称可能是其各自关联公司的商标。