针对图灵架构调优 CUDA 应用程序
本编程指南旨在指导如何针对基于 NVIDIA 图灵架构的 GPU 调优 CUDA 应用程序。
1. 图灵调优指南
1.1. NVIDIA 图灵计算架构
图灵是 NVIDIA 最新的 CUDA 计算应用程序架构。图灵保留并扩展了先前 NVIDIA 架构（如 Pascal 和 Volta）提供的相同 CUDA 编程模型，遵循这些架构最佳实践的应用程序通常无需任何代码更改即可在图灵架构上看到加速效果。1
有关本指南中讨论的编程特性的更多详细信息，请参阅《CUDA C++ 编程指南》。
1.2. CUDA 最佳实践
《CUDA C++ 编程指南》和《CUDA C++ 最佳实践指南》中描述的性能指南和最佳实践适用于所有支持 CUDA 的 GPU 架构。程序员必须主要关注遵循这些建议以实现最佳性能。
这些指南中的高优先级建议如下：
寻找并行化顺序代码的方法，
尽量减少主机和设备之间的数据传输，
调整内核启动配置以最大化设备利用率，
确保全局内存访问已合并，
尽可能减少对全局内存的冗余访问，
避免同一 Warp 中线程出现长时间的发散执行序列。
1.3. 应用程序兼容性
在解决本指南中涵盖的特定性能调优问题之前，请参阅《CUDA 应用程序图灵兼容性指南》，以确保您的应用程序以与图灵兼容的方式编译。
1.4. 图灵调优
1.4.1. 流式多处理器
图灵流式多处理器 (SM) 基于与 Volta 相同的核心架构 (7.x)，并在 Pascal 的基础上提供了类似的改进。
1.4.1.1. 指令调度
每个图灵 SM 包括 4 个 Warp 调度器单元。每个调度器处理一组静态的 Warp，并向一组专用的算术指令单元发出指令。指令在两个周期内执行，调度器可以在每个周期发出独立的指令。核心 FMA 数学运算的依赖指令发布延迟为四个时钟周期，与 Volta 相同，而 Pascal 上为六个周期。因此，核心数学运算的执行延迟可以被每个 SM 最少 4 个 Warp 隐藏，假设每个 Warp 有 4 路指令级并行 ILP，或者在没有任何指令级并行的情况下，被每个 SM 16 个 Warp 隐藏。
与 Volta 类似，图灵 SM 提供 64 个 FP32 核心、64 个 INT32 核心和 8 个改进的混合精度 Tensor Core。图灵的双精度吞吐量低于 Volta，只有 2 个 FP64 核心。
1.4.1.2. 独立线程调度
图灵架构具有与 Volta 相同的独立线程调度特性。这使得以前不可用的 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 数量为 32（Volta 上为 64）。其他影响 Warp 占用率的因素其余方面保持相似
每个 SM 的寄存器文件大小为 64KB 32 位寄存器。
每个线程的最大寄存器数量为 255。
每个 SM 的最大线程块数量为 16。
每个 SM 的共享内存容量为 64KB。
总的来说，开发人员可以期望在图灵上获得与 Pascal 或 Volta 相似的占用率，而无需更改他们的应用程序。
1.4.1.4. 整数算术
与 Volta 类似，图灵 SM 包括专用的 FP32 和 INT32 核心。这使得可以同时执行 FP32 和 INT32 运算。应用程序可以将指针算术与浮点计算交错进行。例如，流水线循环的每次迭代都可以在更新地址并为下一次迭代加载数据的同时，以完整的 FP32 吞吐量处理当前迭代。
1.4.2. Tensor Core 运算
Volta 引入了 Tensor Core 以加速混合精度浮点数据上的矩阵乘法运算。图灵增加了对整数矩阵乘法运算的加速。Tensor Core 在 CUDA 10 C++ API 中以 Warp 级别矩阵运算的形式公开。该 API 提供了专门的矩阵加载、矩阵乘法和累加以及矩阵存储运算，其中每个 Warp 处理一个小的矩阵片段，从而允许从 CUDA-C++ 程序中高效地使用 Tensor Core。在实践中，Tensor Core 用于执行更大的 2D 或更高维度的矩阵运算，这些运算由这些较小的矩阵片段构建而成。
每个 Tensor Core 执行矩阵乘-累加运算：D = A x B + C。Tensor Core 支持半精度矩阵乘法，其中矩阵乘法输入 A 和 B 是 FP16 矩阵，而累加矩阵 C 和 D 可以是 FP16 或 FP32 矩阵。当在 FP32 中累加时，FP16 乘法运算产生一个全精度乘积，然后使用 FP32 加法进行累加。CUDA 10 支持多种片段大小，16x16x16、32x8x16 和 8x32x16，以便在 Volta 或图灵上使用 Tensor Core 和 FP16 输入。
任何为 Volta 编译的二进制文件都可以在图灵上运行，但使用 Tensor Core 的 Volta 二进制文件只能达到图灵 Tensor Core 峰值性能的一半。专门为图灵重新编译二进制文件将使其能够达到峰值性能。有关更多信息，请参阅《图灵兼容性指南》。
图灵的 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 已被告知可能发生此类损害。尽管客户可能因任何原因遭受任何损害，但 NVIDIA 对本文所述产品的客户承担的累计和累积责任应根据产品的销售条款进行限制。
3.2. OpenCL
OpenCL 是 Apple Inc. 的商标，已获得 Khronos Group Inc. 的许可使用。
3.3. 商标
NVIDIA 和 NVIDIA 徽标是 NVIDIA 公司在美国和其他国家/地区的商标或注册商标。其他公司和产品名称可能是与其相关联的各自公司的商标。