针对 Hopper GPU 架构调优 CUDA 应用程序
基于 Hopper GPU 架构的 GPU 的 CUDA 应用程序调优编程指南。
1. NVIDIA Hopper 调优指南
1.1. NVIDIA Hopper GPU 架构
NVIDIA® Hopper GPU 架构是 NVIDIA 最新的 CUDA® 计算应用程序架构。NVIDIA Hopper GPU 架构保留并扩展了先前 NVIDIA GPU 架构（如 NVIDIA Ampere GPU 架构和 NVIDIA Turing）提供的相同 CUDA 编程模型，遵循这些架构最佳实践的应用程序通常无需任何代码更改即可在 NVIDIA H100 GPU 上看到加速。本指南总结了应用程序可以通过利用 NVIDIA Hopper GPU 架构的功能进行微调以获得额外加速的方法。1
有关本指南中讨论的编程功能的更多详细信息，请参阅 CUDA C++ 编程指南。
1.2. CUDA 最佳实践
CUDA C++ 编程指南 和 CUDA C++ 最佳实践指南 中描述的性能指南和最佳实践适用于所有支持 CUDA 的 GPU 架构。程序员必须主要关注遵循这些建议以实现最佳性能。
这些指南中的高优先级建议如下
寻找并行化顺序代码的方法。
最大限度地减少主机和设备之间的数据传输。
调整内核启动配置以最大限度地提高设备利用率。
确保全局内存访问已合并。
尽可能减少对全局内存的冗余访问。
避免同一 Warp 中线程执行的发散长序列。
1.3. 应用程序兼容性
在解决本指南中涵盖的特定性能调优问题之前，请参阅 CUDA 应用程序的 Hopper 兼容性指南，以确保您的应用程序以与 NVIDIA Hopper 兼容的方式编译。
1.4. NVIDIA Hopper 调优
1.4.1. 流式多处理器
与 Turing 和 NVIDIA Ampere GPU 架构相比，NVIDIA Hopper 流式多处理器 (SM) 提供了以下改进。
1.4.1.1. 占用率
每个 SM 的最大并发 Warp 数量与 NVIDIA Ampere GPU 架构中的数量保持不变（即 64），其他 影响 Warp 占用率的因素 包括
寄存器文件大小为每个 SM 64K 32 位寄存器。
每个线程的最大寄存器数为 255。
对于计算能力为 9.0 的设备（即 H100 GPU），每个 SM 的最大线程块数为 32。
对于计算能力为 9.0 的设备（H100 GPU），每个 SM 的共享内存容量为 228 KB，与 A100 的 164 KB 容量相比增加了 39%。
对于计算能力为 9.0 的设备（H100 GPU），每个线程块的最大共享内存为 227 KB。
对于使用线程块集群的应用程序，始终建议使用
cudaOccupancyMaxActiveClusters计算占用率，并相应地启动基于集群的内核。
总的来说，开发人员可以期望在 NVIDIA Ampere GPU 架构 GPU 上获得类似的占用率，而无需更改其应用程序。
1.4.1.2. 张量内存加速器
Hopper 架构建立在 NVIDIA Ampere GPU 架构引入的异步复制之上，并提供更复杂的异步复制引擎：张量内存加速器 (TMA)。
TMA 允许应用程序在全局内存和共享内存之间（双向）以及同一集群中不同 SM 的共享内存区域之间传输 1D 和高达 5D 的张量（请参阅 线程块集群）。此外，对于从共享内存到全局内存的写入，它允许为最常见的数据类型指定逐元素缩减操作，例如加法/最小值/最大值以及按位与/或。
这有几个优点
避免使用寄存器在不同内存空间之间移动数据。
避免使用 SM 指令来移动数据：单个线程可以向 TMA 单元发出大数据移动指令。然后，整个块可以继续处理其他指令，同时数据正在传输中，并且仅在实际需要时才等待数据被消耗。
使用户能够编写 Warp 专用代码，其中特定的 Warp 专门用于不同内存空间之间的数据移动，而其他 Warp 仅处理 SM 内的本地数据。
此功能将通过 cuda::memcpy_async 以及 cuda::barrier 和 cuda::pipeline 公开，用于同步数据移动。
1.4.1.3. 线程块集群
NVIDIA Hopper 架构增加了一个新的可选层级，即线程块集群，这为并行化应用程序提供了更多可能性。线程块可以读取、写入和执行对其集群内其他线程块的共享内存的原子操作。这被称为分布式共享内存。正如 CUDA C++ 编程指南 中所示，有些应用程序无法将所需数据放入共享内存中，而必须使用全局内存。分布式共享内存可以充当这两种选择之间的中间步骤。
分布式共享内存可以与 L2 缓存访问同时被 SM 使用。这可以使需要在 SM 之间通信数据的应用程序受益，方法是利用分布式共享内存和 L2 的组合带宽。
为了实现访问分布式共享内存的最佳性能，应使用 CUDA C++ 最佳实践指南全局内存 中描述的访问模式。具体来说，如果可能，对分布式共享内存的访问应合并并与 32 字节段对齐。应尽可能避免非单位步幅的访问模式，这可以通过使用本地共享内存来实现，类似于 CUDA C++ 最佳实践指南共享内存 中所示。
支持的最大便携式集群大小为 8；但是，NVIDIA Hopper H100 GPU 允许通过选择加入非便携式集群大小 16。使用非便携式集群大小启动内核需要设置 cudaFuncAttributeNonPortableClusterSizeAllowed 函数属性。使用较大的集群大小可能会减少 GPU 上活动块的最大数量（请参阅 占用率）。
1.4.1.4. 改进的 FP32 吞吐量
计算能力为 9.0 的设备比计算能力为 8.0 的设备每个 SM 每个周期具有 2 倍的 FP32 运算。
1.4.1.5. 动态规划指令
NVIDIA Hopper 架构增加了对新指令的支持，以加速动态规划算法，例如用于生物信息学中序列比对的 Smith-Waterman 算法，以及图论、博弈论、ML 和金融问题中的算法。新指令允许计算三个操作数之间的最大值和最小值、产生谓词的最大值和最小值运算、与最大值或最小值组合的加法运算，操作于有符号和无符号 32 位 int 和 16 位 short2 类型以及 half2。所有具有 16 位 short 类型的 DPX 指令 DPX 指令每个 SM 每个周期启用 128 次运算。
1.4.2. 内存系统
1.4.2.1. 高带宽内存 HBM3 子系统
NVIDIA H100 GPU 支持 HBM3 和 HBM2e 内存，容量高达 80 GB。GPU HBM3 内存系统支持高达 3 TB/s 的内存带宽，比 A100-40GB 上的 1.55 TB/s 提高了 93%。
1.4.2.2. 增加的 L2 容量
NVIDIA Hopper 架构将 L2 缓存容量从 A100 GPU 中的 40 MB 增加到 H100 GPU 中的 50 MB。除了增加的容量外，L2 缓存到 SM 的带宽也增加了。NVIDIA Hopper 架构允许 CUDA 用户控制 L2 缓存中数据的持久性，类似于 NVIDIA Ampere GPU 架构。有关 L2 缓存中数据持久性的更多信息，请参阅 CUDA C++ 编程指南 中关于管理 L2 缓存的部分。
1.4.2.3. 内联压缩
NVIDIA Hopper 架构允许 CUDA 计算内核受益于新的内联压缩 (ILC)。此功能可以应用于单个内存分配，压缩器会自动在几种可能的压缩算法之间进行选择，如果没有合适的模式，则不进行压缩。
如果可以使用压缩，则此功能允许以明显高于全局内存带宽的带宽访问全局内存，因为只有压缩数据需要在全局内存和 SM 之间传输。
但是，此功能不允许减少内存占用空间：由于压缩是自动的，即使压缩处于活动状态，内存区域也将使用与没有压缩时相同的占用空间。这是因为底层数据可能会被用户应用程序更改，并且在应用程序的整个持续时间内可能不可压缩。
此功能通过 CUDA 驱动程序 API 提供。请参阅 CUDA C++ 编程指南中关于可压缩内存的部分
CUmemGenericAllocationHandle allocationHandle;
CUmemAllocationProp prop = {};
memset(prop, 0, sizeof(CUmemAllocationProp));
prop->type = CU_MEM_ALLOCATION_TYPE_PINNED;
prop->location.type = CU_MEM_LOCATION_TYPE_DEVICE;
prop->location.id = currentDevice;
prop->allocFlags.compressionType = CU_MEM_ALLOCATION_COMP_GENERIC;
cuMemCreate(&allocationHandle, size, &prop, 0);
可以使用以下方法检查给定设备上是否可以使用可压缩内存
cuDeviceGetAttribute(&compressionAvailable,
CU_DEVICE_ATTRIBUTE_GENERIC_COMPRESSION_SUPPORTED, currentDevice)
请注意，此示例代码不处理错误，并且编译此代码需要链接 CUDA 库 (libcuda.so)。
1.4.3. 第四代 NVLink
第四代 NVIDIA 高速 NVLink 互连在 H100 GPU 中实现，通过每个 GPU 更多链接、更快的通信带宽以及改进的错误检测和恢复功能，显着增强了多 GPU 的可扩展性、性能和可靠性。第四代 NVLink 具有相同的 50 GB/s 每链路双向数据速率。H100 中可用链接的总数增加到 18 个，而 A100 中为 12 个，从而产生 900 GB/s 的双向带宽，而 A100 为 600 GB/s。
NVLink 在现有 CUDA 模型中透明地运行。NVLink 连接的端点之间的传输会自动通过 NVLink 而不是 PCIe 路由。cudaDeviceEnablePeerAccess() API 调用仍然是启用 GPU 之间直接传输（通过 PCIe 或 NVLink）所必需的。cudaDeviceCanAccessPeer() 可用于确定任何一对 GPU 之间是否可以进行对等访问。
2. 修订历史
版本 1.0
首次公开发布
增加对计算能力 9.0 的支持
- 1
-
在本指南中，NVIDIA Volta 指的是计算能力为 7.0 的设备，NVIDIA Turing 指的是计算能力为 7.5 的设备，NVIDIA Ampere GPU 架构指的是计算能力为 8.x 的设备，NVIDIA Hopper 指的是计算能力为 9.0 的设备。
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 已被告知可能发生此类损害。尽管客户可能因任何原因而遭受任何损害，但 NVIDIA 对本文所述产品的客户承担的累计总责任应根据产品的销售条款中的规定进行限制。
3.2. OpenCL
OpenCL 是 Apple Inc. 的商标，已获得 Khronos Group Inc. 的许可使用。
3.3. 商标
NVIDIA 和 NVIDIA 徽标是 NVIDIA Corporation 在美国和其他国家/地区的商标或注册商标。其他公司和产品名称可能是与其相关的各自公司的商标。