针对 Blackwell GPU 架构调优 CUDA 应用程序

针对基于 Blackwell GPU 架构的 GPU 调优 CUDA 应用程序的编程指南。

1. NVIDIA Blackwell 调优指南

1.1. NVIDIA Blackwell GPU 架构

NVIDIA® Blackwell GPU 架构是 NVIDIA 最新的 CUDA® 计算应用程序架构。NVIDIA Blackwell GPU 架构保留并扩展了先前 NVIDIA GPU 架构(如 NVIDIA Ampere GPU 架构和 NVIDIA Hopper)提供的相同 CUDA 编程模型。遵循这些架构最佳实践的应用程序通常可以在 Blackwell GPU 上看到加速,而无需任何代码更改。本指南总结了应用程序可以通过利用 NVIDIA Blackwell GPU 架构的特性来获得额外加速的方法。1

有关本指南中讨论的编程特性的更多详细信息,请参阅 CUDA C++ 编程指南

1.2. CUDA 最佳实践

CUDA C++ 编程指南CUDA C++ 最佳实践指南 中描述的性能指南和最佳实践适用于所有支持 CUDA 的 GPU 架构。程序员必须主要关注遵循这些建议,以实现最佳性能。

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

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

  • 最大限度地减少主机和设备之间的数据传输。

  • 调整内核启动配置以最大限度地提高设备利用率。

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

  • 尽可能最大限度地减少对全局内存的冗余访问。

  • 避免同一 warp 内线程出现长时间的发散执行序列。

1.3. 应用程序兼容性

在解决本指南中涵盖的特定性能调优问题之前,请参阅 CUDA 应用程序的 Blackwell 兼容性指南,以确保您的应用程序以与 NVIDIA Blackwell 兼容的方式编译。

1.4. NVIDIA Blackwell 调优

1.4.1. 流式多处理器

NVIDIA Blackwell 流式多处理器 (SM) 提供了相对于 NVIDIA Hopper GPU 架构的以下改进。

1.4.1.1. 占用率

对于计算能力 10.0,每个 SM 的最大并发 warp 数量为 64,对于计算能力 12.0,则为 48。其他 影响 warp 占用率的因素 包括

  • 寄存器文件大小为每个 SM 64K 32 位寄存器。

  • 每个线程的最大寄存器数为 255。

  • 对于计算能力为 10.0 和 12.0 的设备,每个 SM 的最大线程块数为 32。

  • 对于计算能力为 10.0 的设备,每个 SM 的共享内存容量为 228 KB。对于计算能力为 12.0 的设备,每个 SM 的共享内存容量为 128KB。

  • 对于计算能力为 10.0 的设备,每个线程块的最大共享内存为 227 KB。对于计算能力为 12.0 的设备,每个线程块的最大共享内存为 99 KB。

  • 对于使用线程块集群的应用程序,始终建议使用 cudaOccupancyMaxActiveClusters 计算占用率,并相应地启动基于集群的内核。

总的来说,开发人员可以期望获得与 NVIDIA Hopper GPU 架构 GPU 类似的占用率,而无需更改其应用程序。

1.4.1.2. 线程块集群

NVIDIA Hopper 架构增加了一个新的可选层次结构,即线程块集群,从而在并行化应用程序时提供了更多可能性。Blackwell GPU 也支持线程块集群。线程块可以读取、写入和执行集群内其他线程块的共享内存中的原子操作。这被称为分布式共享内存。正如 CUDA C++ 编程指南 中所示,有些应用程序无法将所需数据放入共享内存中,而必须使用全局内存。分布式共享内存可以作为这两种选项之间的中间步骤。

分布式共享内存可以与 L2 缓存访问同时使用。这可以使需要在 SM 之间通信数据的应用程序受益,方法是利用分布式共享内存和 L2 的组合带宽。

为了在访问分布式共享内存时获得最佳性能,应使用 CUDA C++ 最佳实践指南全局内存 中描述的访问模式。具体来说,如果可能,对分布式共享内存的访问应合并并对齐到 32 字节段。应尽可能避免使用非单位步幅的访问模式,这可以通过使用本地共享内存来实现,类似于 CUDA C++ 最佳实践指南共享内存 中所示。

支持的最大便携式集群大小为 8;但是,NVIDIA Blackwell B200 GPU 允许通过选择加入非便携式集群大小 16。使用非便携式集群大小启动内核需要设置 cudaFuncAttributeNonPortableClusterSizeAllowed 函数属性。使用较大的集群大小可能会减少 GPU 上活动块的最大数量(请参阅 占用率)。

1.4.2. 内存系统

1.4.2.1. 高带宽内存 HBM3 子系统

NVIDIA B200 GPU 支持 HBM3 和 HBM3e 内存,容量高达 180 GB。

1.4.2.2. 增加的 L2 容量

NVIDIA GB200 GPU 将 L2 缓存容量增加到 126 MB。

NVIDIA Blackwell 架构允许 CUDA 用户控制 L2 缓存中数据的持久性,类似于 NVIDIA Ampere GPU 架构。有关 L2 缓存中数据持久性的更多信息,请参阅 CUDA C++ 编程指南 中关于管理 L2 缓存的部分。

1.4.2.3. 统一共享内存/L1/纹理缓存

计算能力为 10.0 的 NVIDIA B200 GPU 与之前的 NVIDIA Hopper 架构具有相同的组合 L1 缓存、纹理缓存和共享内存的最大容量 256 KB。

在 NVIDIA Blackwell GPU 架构中,专用于共享内存的 L1 缓存部分(称为 carveout)可以在运行时选择,就像之前的架构(如 NVIDIA Ampere 架构和 NVIDIA Volta)一样,使用 cudaFuncSetAttribute() 和属性 cudaFuncAttributePreferredSharedMemoryCarveout。NVIDIA H100 GPU 和 NVIDIA B200 GPU 都支持每个 SM 0、8、16、32、64、100、132、164、196 和 228 KB 的共享内存容量。

CUDA 为每个线程块保留 1 KB 的共享内存。因此,B200 GPU 使单个线程块能够寻址高达 227 KB 的共享内存。为了保持架构兼容性,静态共享内存分配仍然限制为 48 KB,并且还需要显式选择加入才能启用高于此限制的动态分配。有关详细信息,请参阅 CUDA C++ 编程指南

与追溯到 NVIDIA Ampere 架构(计算能力 8.x)的 GPU 架构类似,NVIDIA Blackwell GPU 架构将 L1 和纹理缓存的功能组合到一个统一的 L1/纹理缓存中,该缓存充当内存访问的合并缓冲区,在将数据传递到 warp 之前,收集 warp 线程请求的数据。与之前的架构类似,与其与共享内存的结合的另一个好处是在延迟和带宽方面的改进。

2. 修订历史

版本 1.0

  • 首次公开发布

  • 增加了对计算能力 10.0 和计算能力 12.0 的支持

1

在本指南中,NVIDIA Volta 指的是计算能力为 7.0 的设备,NVIDIA Turing 指的是计算能力为 7.5 的设备,NVIDIA Ampere GPU 架构指的是计算能力为 8.x 的设备,NVIDIA Hopper 指的是计算能力为 9.0 的设备,NVIDIA Blackwell 指的是计算能力为 10.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 对本文所述产品的客户的累计总责任应根据产品的销售条款的规定进行限制。

3.2. OpenCL

OpenCL 是 Apple Inc. 的商标,已授权 Khronos Group Inc. 使用。

3.3. 商标

NVIDIA 和 NVIDIA 徽标是 NVIDIA Corporation 在美国和其他国家/地区的商标或注册商标。其他公司和产品名称可能是与其相关的各自公司的商标。