为Hopper GPU架构优化CUDA应用程序
基于Hopper 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++ Programming Guide。
1.2. CUDA最佳实践
《CUDA C++编程指南》和《CUDA C++最佳实践指南》中描述的性能准则和最佳实践适用于所有支持CUDA的GPU架构。程序员主要需要遵循这些建议以获得最佳性能。
这些指南中的高优先级建议如下：
寻找并行化顺序代码的方法。
最小化主机与设备之间的数据传输。
调整内核启动配置以最大化设备利用率。
确保全局内存访问是合并的。
尽可能减少对全局内存的冗余访问。
避免同一warp内的线程执行过长的分支序列。
1.3. 应用兼容性
在解决本指南涵盖的具体性能调优问题之前，请参阅Hopper Compatibility Guide for CUDA Applications以确保您的应用程序以兼容NVIDIA Hopper架构的方式进行编译。
1.4. NVIDIA Hopper调优
1.4.1. 流式多处理器
NVIDIA Hopper流式多处理器(SM)相比图灵和NVIDIA安培GPU架构提供了以下改进。
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安培GPU架构相似的占用率。
1.4.1.2. 张量内存加速器
Hopper架构基于NVIDIA Ampere GPU架构引入的异步拷贝功能进行了扩展，提供了一个更先进的异步拷贝引擎：张量内存加速器(TMA)。
TMA允许应用程序在全局内存和共享内存之间双向传输1D至5D张量，同时也能在同一集群内不同SM的共享内存区域之间传输（参见线程块集群）。此外，对于从共享内存到全局内存的写入操作，它还支持为大多数常见数据类型指定按元素的归约操作，如加法/最小值/最大值，以及按位与/或操作。
这具有以下几个优势：
避免使用寄存器在不同内存空间之间移动数据。
避免使用SM指令移动数据：单个线程可以向TMA单元发出大数据传输指令。整个线程块可以在数据传输过程中继续处理其他指令，仅在真正需要消费数据时才等待数据到达。
允许用户编写warp专用代码，其中特定的warp专门负责不同内存空间之间的数据传输，而其他warp仅处理SM内的本地数据。
该功能将通过cuda::memcpy_async以及用于同步数据移动的cuda::barrier和cuda::pipeline来提供。
1.4.1.3. 线程块集群
NVIDIA Hopper架构引入了一个新的可选层次结构——线程块集群(Thread Block Clusters)，为应用程序并行化提供了更多可能性。线程块可以读取、写入以及对其集群内其他线程块的共享内存执行原子操作。这被称为分布式共享内存(Distributed Shared Memory)。如CUDA C++编程指南所示，某些应用程序无法将所需数据完全放入共享内存，必须改用全局内存。分布式共享内存可以作为这两种选择之间的中间方案。
分布式共享内存可以与L2缓存访问同时被SM使用。这有利于需要通过利用分布式共享内存和L2的组合带宽在SM之间通信数据的应用程序。
为了实现对分布式共享内存的最佳访问性能，应采用CUDA C++全局内存最佳实践指南中描述的访问模式。具体而言，应尽可能将分布式共享内存的访问合并并对齐到32字节段。如有可能，应避免使用非单位步长的访问模式，这可以通过使用本地共享内存来实现，类似于CUDA C++共享内存最佳实践指南中展示的方法。
支持的最大可移植集群大小为8；然而，NVIDIA Hopper H100 GPU通过选择启用可支持16的非可移植集群大小。启动具有非可移植集群大小的内核需要设置cudaFuncAttributeNonPortableClusterSizeAllowed函数属性。使用更大的集群大小可能会减少GPU上活动块的最大数量（请参阅占用率）。
1.4.1.4. 提升的FP32吞吐量
计算能力9.0的设备每个SM每个周期的FP32运算量是计算能力8.0设备的2倍。
1.4.1.5. 动态编程指令
NVIDIA Hopper架构新增了对动态规划算法加速的支持，例如生物信息学中用于序列比对的Smith-Waterman算法，以及图论、博弈论、机器学习和金融问题中的算法。新指令支持三操作数间的最大值/最小值计算、可生成谓词的最大值/最小值运算、结合加法与最大值/最小值的复合运算，可操作有符号/无符号32位整型、16位short2类型以及half2类型。所有支持16位short类型的DPX指令使得每个SM每周期可执行128次运算。
1.4.2. 内存系统
1.4.2.1. 高带宽内存HBM3子系统
NVIDIA H100 GPU支持HBM3和HBM2e内存，容量最高可达80GB。该GPU的HBM3内存系统支持高达3TB/s的内存带宽，相比A100-40GB的1.55TB/s提升了93%。
1.4.2.2. 提升的L2缓存容量
NVIDIA Hopper架构将L2缓存容量从A100 GPU的40 MB提升至H100 GPU的50 MB。在增加容量的同时，L2缓存到流式多处理器(SMs)的带宽也得到了提升。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的可用链路总数从A100的12个增加到18个，双向带宽达到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公司（“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公司的商标或注册商标。其他公司及产品名称可能是其各自关联公司的商标。