为Pascal架构优化CUDA应用程序

基于NVIDIA Pascal架构的GPU调优CUDA应用程序编程指南。

1. Pascal调优指南

1.1. NVIDIA Pascal计算架构

Pascal保留并扩展了与Maxwell等前代NVIDIA架构相同的CUDA编程模型。遵循这些架构最佳实践的应用程序通常无需修改代码即可在Pascal架构上获得加速。本指南总结了如何通过利用Pascal架构特性对应用程序进行微调以获得额外加速。1

Pascal架构包含两个主要变体：GP100和GP104。2 关于GP100和GP104相比早期NVIDIA架构的主要改进，在两份白皮书中有详细概述：针对GP100的NVIDIA Tesla P100: The Most Advanced Datacenter Accelerator Ever Built和针对GP104的NVIDIA GeForce GTX 1080: Gaming Perfected。

有关本指南中讨论的编程功能的更多详细信息，请参阅CUDA C++编程指南。本指南中描述的一些Pascal功能特定于GP100或GP104（如标注所示）；如果未指定，则功能适用于两种Pascal变体。

1.2. CUDA最佳实践

《CUDA C++编程指南》和《CUDA C++最佳实践指南》中描述的性能准则和最佳实践适用于所有支持CUDA的GPU架构。程序员主要需要遵循这些建议以获得最佳性能。

这些指南中的高优先级建议如下：

  • 寻找并行化顺序代码的方法，

  • 尽量减少主机和设备之间的数据传输，

  • 调整内核启动配置以最大化设备利用率，

  • 确保全局内存访问是合并的，

  • 尽可能减少对全局内存的冗余访问

  • 避免同一warp内的线程执行过长的分支序列。

1.3. 应用兼容性

在解决本指南涵盖的具体性能调优问题之前，请先查阅Pascal兼容性指南（适用于CUDA应用程序），以确保您的应用程序以兼容Pascal架构的方式进行编译。

1.4. Pascal调优

1.4.1. 流式多处理器

Pascal流式多处理器(SM)在许多方面与Maxwell类似。Pascal通过改进的16纳米FinFET制造工艺和各种架构修改，进一步提升了Maxwell架构已经非常出色的能效表现。

1.4.1.1. 指令调度

与Maxwell类似，Pascal架构在每个分区中采用2的幂次方数量的CUDA核心。这种设计简化了调度过程，因为每个流式多处理器（SM）的线程束调度器会向一组专有的CUDA核心（数量等于线程束宽度32）分发指令。每个线程束调度器仍保持双发射的灵活性（例如在同一周期内可向CUDA核心发射数学运算指令，同时向加载/存储单元发射内存操作指令），但现在单次发射已足以充分利用所有CUDA核心。

GP100和GP104架构在每个SM中集成了不同数量的CUDA核心。与Maxwell类似，每个GP104 SM配备四个线程束调度器，共管理128个单精度(FP32)核心和四个双精度(FP64)核心。GP104处理器最多可提供20个SM，而类似的GP102设计最多可提供30个SM。

相比之下，GP100提供了规模更小但数量更多的流式多处理器(SMs)。每个GP100最多可配备60个SMs。3 每个SM包含两个warp调度器，共管理64个FP32核心和32个FP64核心。由此形成的2:1 FP32与FP64核心比例与GP100新型数据路径配置完美匹配，使得Pascal架构处理FP64工作负载的效率比前代强调FP64性能的NVIDIA架构Kepler GK210更高。

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架构类似，相比FP32或FP64存储，使用FP16存储可以降低内存占用和带宽需求。Pascal还新增了对原生FP16指令的支持。通过使用配对操作在每个核心同时执行两条FP16指令，可实现峰值FP16吞吐量。要符合配对操作条件，操作数必须存储在half2向量类型中。GP100和GP104提供不同的FP16吞吐量：专为深度神经网络训练设计的GP100，其FP16吞吐量可达FP32运算的2倍；而GP104的FP16吞吐量较低，仅为FP32的1/64。不过作为补偿，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在715MHz的保守内存时钟下即可实现732GB/s的峰值带宽，其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/纹理缓存，该缓存作为内存访问的合并缓冲区，在将数据传递给线程束之前，先收集该线程束中线程所请求的数据。

默认情况下，GP100会在L1/纹理缓存中缓存全局加载。相比之下，GP104遵循Maxwell架构，仅将全局加载缓存在L2中，除非使用LDG只读数据缓存机制。与之前的架构一样，GP104允许开发者在编译时通过向-Xptxas -dlcm=ca传递标志给nvcc来选择将所有全局加载缓存在统一的L1/纹理缓存中。

Kepler在启用全局加载的L1缓存时以128B的粒度处理负载，否则为32B。在Pascal架构上，无论全局加载是否缓存在L1中，数据访问单元都是32B。因此，不再需要关闭L1缓存来减少与非合并访问相关的全局内存事务浪费。

与Maxwell不同，Pascal将线程本地内存缓存在L1缓存中。相比Maxwell，这可以减轻寄存器溢出的成本。因此应重新评估占用率与溢出之间的平衡，以确保最佳性能。

CUDA Toolkit 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内的原子可见性。

注意

考虑到原子作用域可能被错误使用，建议应用程序使用compute-sanitizer来检测并消除错误。

在Pascal架构中实现的全系统原子操作，旨在让开发者能够尝试增强的内存模型。这些原子操作通过软件实现，需要特别注意才能获得良好性能。当原子操作的目标地址是可迁移地址且位于远程内存空间时，本地处理器会发生页错误，以便内核能将相应内存页迁移到本地内存。随后使用常规硬件指令执行原子操作。由于该页面现已驻留在本地，来自同一处理器的后续原子操作不会导致额外的页错误。然而，来自不同处理器的原子更新可能会引发频繁的页错误。

1.4.5. 共享内存

1.4.5.1. 共享内存容量

在Kepler架构中，共享内存和L1缓存共享相同的片上存储空间。相比之下，Maxwell和Pascal架构为每个SM的共享内存提供了专用空间，因为L1缓存和纹理缓存的功能已合并。与Kepler相比，这增加了每个SM可用的共享内存空间：GP100提供每个SM 64 KB共享内存，GP104则提供每个SM 96 KB。

这为应用程序开发者带来了多项优势：

  • 对共享内存容量需求较高的算法（如基数排序）除了SM数量增加带来的整体提升外，每个SM的容量还能自动获得33%至100%的提升。

  • 应用程序不再需要为获得最佳性能而选择L1/共享内存的分割偏好。

注意

线程块(Thread-blocks)的共享内存容量仍限制在48 KB。为了获得最大的灵活性，NVIDIA建议应用程序在每个线程块中最多使用32 KB共享内存。例如，这将允许至少两个线程块适配到每个GP100 SM中，或每个GP104 SM适配三个线程块。

1.4.5.2. 共享内存带宽

Kepler 提供了一种可选的8字节共享内存分块模式，这种模式有可能提高每个SM在访问8或16字节共享内存时的带宽。然而，只有当应用程序在共享内存中存储这些较大的元素时（即整数和fp32值无法受益），并且开发者通过API明确选择启用8字节分块模式时，才能从中获益。

为了简化这一点，Pascal遵循Maxwell的设计，回归到固定的四字节存储体。这使得所有使用共享内存的应用程序都能从更高的带宽中受益，而无需通过API指定任何特殊偏好。

1.4.6. GPU间通信

1.4.6.2. GPUDirect RDMA带宽

GPUDirect RDMA允许第三方设备（如网络接口卡NIC）直接访问GPU内存。这消除了不必要的复制缓冲区，降低了CPU开销，并显著减少了从/向GPU内存发送/接收MPI消息的延迟。Pascal架构在通过PCIe从源GPU内存读取数据并写入目标NIC内存时，将提供的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启发式算法表现欠佳时，可以通过在源代码中添加一组迁移提示来进行进一步调优。

在支持的操作系统平台上，任何使用默认操作系统分配器（例如malloc或new）分配的内存，都可以通过相同的指针从GPU和CPU代码中访问。实际上，所有系统虚拟内存都可以从GPU访问。在此类系统上，无需显式使用cudaMallocManaged()来分配托管内存。

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公司的商标或注册商标。其他公司及产品名称可能是其各自关联公司的商标。

1

在本指南中，Kepler指计算能力3.x的设备，Maxwell指计算能力5.x的设备，Pascal指计算能力6.x的设备。

2

GP100和GP104的具体计算能力分别为6.0和6.1。GP102架构与GP104类似。

3

Tesla P100启用了56个流式多处理器(SMs)。

4

作为例外情况，对HBM2的分散写入会受到ECC的一些开销影响，但远低于受ECC保护的GDDR5内存上类似访问模式的开销。

5

可迁移的，或称为统一内存(UM)的分配，是通过cudaMallocManaged()实现的，或者对于支持异构内存管理(HMM)的系统，使用malloc()。

6

这里的 [Op] 可以是 Add, CAS 等其中之一。