针对 Pascal 调优 CUDA 应用程序
基于 NVIDIA Pascal 架构的 GPU 的 CUDA 应用程序调优编程指南。
1. Pascal 调优指南
1.1. NVIDIA Pascal 计算架构
Pascal 保留并扩展了先前 NVIDIA 架构（如 Maxwell）提供的 CUDA 编程模型。遵循这些架构最佳实践的应用程序通常无需任何代码更改，即可在 Pascal 架构上获得性能提升。本指南概述了如何通过利用 Pascal 架构特性来微调应用程序，从而获得额外的性能提升。1
Pascal 架构包含两个主要变体：GP100 和 GP104。2 关于 GP100 和 GP104 相对于早期 NVIDIA 架构的主要改进的详细概述，请参阅两份白皮书，标题分别为 NVIDIA Tesla P100：有史以来最先进的数据中心加速器（针对 GP100）和 NVIDIA GeForce GTX 1080：游戏臻于完美（针对 GP104）。
有关本指南中讨论的编程特性的更多详细信息，请参阅 CUDA C++ 编程指南。本指南中描述的某些 Pascal 特性是 GP100 或 GP104 特有的，如已注明；如果未指定，则这些特性适用于 Pascal 的两种变体。
1.2. CUDA 最佳实践
CUDA C++ 编程指南 和 CUDA C++ 最佳实践指南 中描述的性能指南和最佳实践适用于所有支持 CUDA 的 GPU 架构。程序员应主要关注遵循这些建议以实现最佳性能。
这些指南中的高优先级建议如下：
寻找并行化顺序代码的方法，
最大限度地减少主机和设备之间的数据传输，
调整内核启动配置以最大限度地提高设备利用率，
确保全局内存访问合并，
尽可能减少对全局内存的冗余访问，
避免同一 Warp 内线程执行的长时间发散序列。
1.3. 应用程序兼容性
在解决本指南中涵盖的特定性能调优问题之前，请参阅 CUDA 应用程序 Pascal 兼容性指南，以确保您的应用程序以与 Pascal 兼容的方式编译。
1.4. Pascal 调优
1.4.1. 流式多处理器
Pascal 流式多处理器 (SM) 在许多方面与 Maxwell 的 SM 相似。Pascal 通过改进的 16 纳米 FinFET 制造工艺和各种架构修改，进一步提高了 Maxwell 架构已经提供的出色能效。
1.4.1.1. 指令调度
与 Maxwell 一样，Pascal 每个分区采用 2 的幂数的 CUDA 核心。这简化了调度，因为每个 SM 的 Warp 调度器都向一组专用的 CUDA 核心发出指令，其数量等于 Warp 宽度 (32)。每个 Warp 调度器仍然可以灵活地双发射（例如，在同一周期内向 CUDA 核心发出数学运算，同时向加载/存储单元发出内存运算），但现在单发射就足以充分利用所有 CUDA 核心。
GP100 和 GP104 设计在每个 SM 中集成了不同数量的 CUDA 核心。与 Maxwell 一样，每个 GP104 SM 提供四个 Warp 调度器，总共管理 128 个单精度 (FP32) 和四个双精度 (FP64) 核心。GP104 处理器最多提供 20 个 SM，而类似的 GP102 设计最多提供 30 个 SM。
相比之下，GP100 提供的 SM 更小但数量更多。每个 GP100 最多提供 60 个 SM。3 每个 SM 包含两个 Warp 调度器，总共管理 64 个 FP32 和 32 个 FP64 核心。FP32 与 FP64 核心的 2:1 比例与 GP100 的新数据路径配置非常吻合，使 Pascal 能够比以前强调 FP64 性能的 NVIDIA 架构 Kepler GK210 更高效地处理 FP64 工作负载。
1.4.1.2. 占用率
每个 SM 的最大并发 Warp 数与 Maxwell 中保持不变（即 64），其他 影响 Warp 占用率的因素 也保持相似。
寄存器文件大小（64k 32 位寄存器）与 Maxwell 的相同。
每个线程的最大寄存器数 255 与 Maxwell 的相同。与之前的架构一样，应通过实验来确定寄存器溢出与占用率之间的最佳平衡。
每个 SM 的最大线程块数为 32，与 Maxwell 相同。
GP100 的每个 SM 的共享内存容量为 64KB，GP104 为 96KB。相比之下，Maxwell 分别提供 96KB 和高达 112KB 的共享内存。但每个 GP100 SM 包含的 CUDA 核心更少，因此 GP100 上每个核心可用的共享内存实际上增加了。每个块的最大共享内存仍然限制为 48KB，与之前的架构相同（请参阅 共享内存容量）。
因此，开发人员可以期望与 Maxwell 相似的占用率，而无需更改其应用程序。由于相对于 Kepler 的调度改进，实现最大设备利用率所需的 Warp 占用率要求（即，可用并行性）通常会降低。
1.4.2. 新的算术原语
1.4.2.1. FP16 算术支持
Pascal 为深度学习等对低浮点精度容忍的应用提供了改进的 FP16 支持。half 类型用于表示设备上的 FP16 值。与 Maxwell 一样，FP16 存储可用于减少与 FP32 或 FP64 存储相比所需的内存占用和带宽。Pascal 还增加了对原生 FP16 指令的支持。通过使用配对操作同时执行每个核心两条 FP16 指令，可以达到峰值 FP16 吞吐量。为了符合配对操作的条件，操作数必须存储在 half2 向量类型中。GP100 和 GP104 提供不同的 FP16 吞吐量。GP100 旨在用于训练深度神经网络，其 FP16 吞吐量高达 FP32 算术的两倍。在 GP104 上，FP16 吞吐量较低，为 FP32 的 1/64。然而，为了补偿降低的 FP16 吞吐量，GP104 提供了 GP100 中不可用的额外的高吞吐量 INT8 支持。
1.4.2.2. INT8 点积
GP104 为双向和四向整数点积提供专用指令。这些指令非常适合加速深度学习推理工作负载。__dp4a 内在函数计算四个 8 位整数的点积，并将结果累加到 32 位整数中。类似地，__dp2a 在一个向量中的两个 16 位整数和另一个向量中的两个 8 位整数之间执行双元素点积，并将结果累加到 32 位整数中。两条指令都提供与 FP32 算术相等的吞吐量。
1.4.3. 内存吞吐量
1.4.3.1. 高带宽内存 2 DRAM
GP100 将高带宽内存 2 (HBM2) 用于其 DRAM。HBM2 内存与 GPU 芯片一起堆叠在单个硅封装上。与传统的 GDDR 技术相比，这允许更宽的接口和相似的功耗。GP100 连接到最多四个 HBM2 堆栈，每个堆栈使用两个 512 位内存控制器。然后内存总线的有效宽度为 4096 位，比 GM200 的 384 位有了显着提高。即使在降低的内存时钟频率下，这也允许峰值带宽大幅提升。因此，配备 GP100 的 Tesla P100 的峰值带宽为 732 GB/s，内存时钟频率适中，为 715 MHz。DRAM 访问延迟与 Maxwell 上观察到的延迟相似。
为了在全 HBM2 带宽下隐藏 DRAM 延迟，与配备传统 GDDR5 的 GPU 相比，必须保持更多的内存访问处于运行状态。值得庆幸的是，GP100 中大量的 SM 通常会增加与以前的架构相比的并发线程数（以及正在运行的读取次数）。资源受限的内核（受限于低占用率）可能会受益于增加每个线程的并发内存访问次数。
GP100 GPU 的寄存器文件、共享内存、L1 和 L2 缓存以及 DRAM 均受单错误纠正双错误检测 (SECDED) ECC 代码保护。在 Kepler GK210 上启用 ECC 支持时，可用的 DRAM 将减少 6.25%，以便存储 ECC 位。与禁用 ECC 的同一 GPU 相比，为每个内存事务获取 ECC 位也会使有效带宽降低约 20%。另一方面，HBM2 内存提供专用的 ECC 资源，从而实现零开销的 ECC 保护。4
1.4.3.2. 统一 L1/纹理缓存
与 Maxwell 一样，Pascal 将 L1 和纹理缓存的功能组合到一个统一的 L1/纹理缓存中，该缓存充当内存访问的合并缓冲区，在将数据传递到 Warp 之前，收集 Warp 的线程请求的数据。
默认情况下，GP100 在 L1/纹理缓存中缓存全局加载。相比之下，GP104 遵循 Maxwell，仅在 L2 中缓存全局加载，除非使用 LDG 只读数据缓存机制。与之前的架构一样，GP104 允许开发人员通过在编译时将 -Xptxas -dlcm=ca 标志传递给 nvcc，选择在统一的 L1/纹理缓存中缓存所有全局加载。
当启用全局加载的 L1 缓存时，Kepler 以 128B 的粒度提供加载服务，否则为 32B。在 Pascal 上，数据访问单元为 32B，无论全局加载是否在 L1 中缓存。因此，不再需要关闭 L1 缓存，以减少与非合并访问相关的浪费的全局内存事务。
与 Maxwell 不同，Pascal 在 L1 缓存中缓存线程局部内存。与 Maxwell 相比，这可以减轻寄存器溢出的成本。因此，应重新评估占用率与溢出的平衡，以确保最佳性能。
CUDA 工具包 6.0 中添加了两个新的设备属性：globalL1CacheSupported 和 localL1CacheSupported。希望为各种架构世代提供单独调整路径的开发人员可以使用这些字段来简化路径选择过程。
注意
在 GP104 中启用全局缓存可能会影响占用率。如果每个线程块的 SM 资源使用率导致启用缓存时占用率为零，则 CUDA 驱动程序将覆盖缓存选择，以允许内核启动成功。分析器会报告这种情况。
1.4.4. 原子内存操作
与 Maxwell 一样，Pascal 为 32 位整数算术提供原生共享内存原子操作，以及原生 32 位或 64 位比较和交换 (CAS)。来自 Kepler 的开发人员（在 Kepler 中，共享内存原子操作是使用锁/更新/解锁序列在软件中实现的）应该会看到巨大的性能提升，特别是对于竞争激烈的共享内存原子操作。
Pascal 还扩展了全局内存中的原子加法，使其可以在 FP64 数据上运行。因此，CUDA 中的 atomicAdd() 函数已推广为支持 32 位和 64 位整数和浮点类型。Pascal 中所有浮点原子操作的舍入模式均为就近舍入到偶数。与前几代产品一样，FP32 atomicAdd() 会将非正规化值刷新为零。
对于 GP100，原子操作可能以通过 NVLink 连接的对等 GPU 的内存为目标。NVLink 上的对等原子操作使用与以全局内存为目标的原子操作相同的 API。通过 PCIE 连接的 GPU 不支持此功能。
Pascal GPU 提供支持系统范围的原子操作，目标是可迁移分配5 如果需要系统范围的原子可见性，则以可迁移内存为目标的操作必须通过使用 atomic[Op]_system() 内在函数6 指定系统范围。在可迁移内存上使用设备范围的原子操作（例如 atomicAdd()）仍然有效，但仅在本地 GPU 内强制执行原子可见性。
注意
鉴于原子范围可能被错误使用，建议应用程序使用计算 санитайзер 来检测和消除错误。
正如为 Pascal 实现的那样，系统范围的原子操作旨在允许开发人员尝试增强的内存模型。它们是在软件中实现的，需要一定的注意才能实现良好的性能。当原子操作以远程内存空间支持的可迁移地址为目标时，本地处理器会发生页面错误，以便内核可以将相应的内存页面迁移到本地内存。然后使用通常的硬件指令来执行原子操作。由于页面现在位于本地，因此来自同一处理器的后续原子操作不会导致额外的页面错误。但是，来自不同处理器的原子更新可能会导致频繁的页面错误。
1.4.6. GPU 间通信
1.4.6.1. NVLink 互连
NVLink 是 NVIDIA 的新型高速数据互连。NVLink 可用于显着提高 GPU 到 GPU 通信和 GPU 访问系统内存的性能。GP100 最多支持四个 NVLink 连接，每个连接承载高达 40 GB/s 的双向带宽。
NVLink 在现有的 CUDA 模型中透明运行。NVLink 连接的端点之间的传输会自动通过 NVLink 而不是 PCIe 路由。cudaDeviceEnablePeerAccess() API 调用仍然是启用 GPU 之间直接传输（通过 PCIe 或 NVLink）所必需的。cudaDeviceCanAccessPeer() 可用于确定任何一对 GPU 之间是否可以进行对等访问。
1.4.6.2. GPUDirect RDMA 带宽
GPUDirect RDMA 允许第三方设备（如网络接口卡 (NIC)）直接访问 GPU 内存。这消除了不必要的复制缓冲区，降低了 CPU 开销，并显着降低了从/到 GPU 内存的 MPI 发送/接收消息的延迟。当通过 PCIe 从源 GPU 内存读取数据并写入目标 NIC 内存时，Pascal 将提供的 RDMA 带宽提高了一倍。
1.4.7. 计算抢占
计算抢占是 GP100 特有的新功能。计算抢占允许在指令级粒度中断 GPU 上运行的计算任务。执行上下文（寄存器、共享内存等）被交换到 GPU DRAM，以便可以换入另一个应用程序并运行。计算抢占为开发人员提供了两个主要优势：
长时间运行的内核不再需要分解为小的时隙，以避免在 GPU 同时用于计算和图形时，图形用户界面无响应或内核超时。
现在可以在单 GPU 系统上进行交互式内核调试。
1.4.8. 统一内存改进
Pascal 提供了新的硬件功能来扩展统一内存 (UM) 支持。扩展的 49 位虚拟寻址空间允许 Pascal GPU 通过单个虚拟地址空间寻址现代 CPU 的完整 48 位虚拟地址空间以及系统中所有 GPU 的内存，不受任何一个处理器的物理内存大小的限制。Pascal GPU 还支持内存页面错误。页面错误允许应用程序从主机和设备访问相同的托管内存分配，而无需显式同步。它还消除了 CUDA 运行时在每次内核启动之前预同步所有托管内存分配的需要。相反，当内核访问非驻留内存页面时，它会发生错误，并且可以按需将页面迁移到 GPU 内存，或者映射到 GPU 地址空间以通过 PCIe/NVLink 接口进行访问。
这些特性提高了 Pascal 上许多典型 UM 工作负载的性能。在 UM 启发式方法被证明并非最佳的情况下，可以通过添加到源代码中的一组迁移提示进行进一步的调优。
在支持的操作系统平台上，可以使用相同的指针从 GPU 和 CPU 代码访问使用默认操作系统分配器（例如，malloc 或 new）分配的任何内存。实际上，可以从 GPU 访问所有系统虚拟内存。在此类系统上，无需使用 cudaMallocManaged() 显式分配托管内存。
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 在美国和其他国家/地区的商标或注册商标。其他公司和产品名称可能是与其相关的各自公司的商标。
- 1
-
在本指南中，Kepler 指的是计算能力为 3.x 的设备，Maxwell 指的是计算能力为 5.x 的设备，Pascal 指的是计算能力为 6.x 的设备。
- 2
-
GP100 和 GP104 的具体计算能力分别为 6.0 和 6.1。GP102 架构与 GP104 相似。
- 3
-
Tesla P100 启用了 56 个 SM。
- 4
-
作为例外，分散写入 HBM2 会看到来自 ECC 的一些开销，但远低于在 ECC 保护的 GDDR5 内存上使用类似访问模式的开销。
- 5
-
可迁移或统一内存 (UM) 分配是使用
cudaMallocManaged()或对于支持异构内存管理 (HMM) 的系统使用malloc()完成的。 - 6
-
此处 [Op] 将是
Add、CAS等之一。