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

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

1. NVIDIA Ampere GPU 架构调优指南

1.1. NVIDIA Ampere GPU 架构

NVIDIA Ampere GPU 架构是 NVIDIA 最新的 CUDA 计算应用程序架构。NVIDIA Ampere GPU 架构保留并扩展了先前 NVIDIA GPU 架构（如 Turing 和 Volta）提供的相同 CUDA 编程模型，并且遵循这些架构的最佳实践的应用程序通常应该在 NVIDIA A100 GPU 上看到加速，而无需任何代码更改。本指南总结了通过利用 NVIDIA Ampere GPU 架构的功能来微调应用程序以获得额外加速的方法。1

有关本指南中讨论的编程功能的更多详细信息，请参阅 CUDA C++ 编程指南。

1.2. CUDA 最佳实践

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

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

寻找并行化顺序代码的方法。
最大限度地减少主机和设备之间的数据传输。
调整内核启动配置以最大限度地提高设备利用率。
确保全局内存访问合并。
尽可能减少对全局内存的冗余访问。
避免同一 Warp 内线程中出现长时间的发散执行序列。

1.3. 应用程序兼容性

在解决本指南中介绍的特定性能调优问题之前，请参阅 NVIDIA Ampere GPU 架构 CUDA 应用程序兼容性指南，以确保您的应用程序以与 NVIDIA Ampere GPU 架构兼容的方式编译。

1.4. NVIDIA Ampere GPU 架构调优

1.4.1. 流式多处理器

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

1.4.1.1. 占用率

对于计算能力 8.0，每个 SM 的最大并发 Warp 数量与 Volta（即 64）保持不变，而对于计算能力 8.6，则为 48。其他影响 Warp 占用率的因素是

寄存器文件大小为每个 SM 64K 32 位寄存器。
每个线程的最大寄存器数为 255。
对于计算能力为 8.0 的设备（即 A100 GPU），每个 SM 的最大线程块数为 32，对于计算能力为 8.6 的 GPU，则为 16。
对于计算能力为 8.0 的设备（即 A100 GPU），每个 SM 的共享内存容量为 164 KB，与 V100 的 96 KB 容量相比增加了 71%。对于计算能力为 8.6 的 GPU，每个 SM 的共享内存容量为 100 KB。
对于计算能力为 8.0 的设备（即 A100 GPU），每个线程块的最大共享内存为 163 KB。对于计算能力为 8.6 的 GPU，每个线程块的最大共享内存为 99 KB。

总的来说，开发人员可以期望与 Volta 类似的占用率，而无需更改其应用程序。

1.4.1.2. 从全局内存到共享内存的异步数据复制

NVIDIA Ampere GPU 架构增加了用于将数据从全局内存复制到共享内存的硬件加速。这些复制指令相对于计算是异步的，并允许用户显式控制计算与从全局内存到 SM 的数据移动的重叠。这些指令还避免使用额外的寄存器进行内存复制，并且还可以绕过 L1 缓存。此新功能通过 CUDA 中的 pipeline API 公开。有关更多信息，请参阅 CUDA C++ 编程指南中关于异步复制的部分。

1.4.1.3. 用于拆分到达/等待屏障的硬件加速

NVIDIA Ampere GPU 架构为共享内存中的拆分到达/等待屏障增加了硬件加速。这些屏障可用于在 CUDA 中实现细粒度的线程控制、生产者-消费者计算流水线和发散代码模式。这些屏障也可以与异步复制一起使用。有关到达/等待屏障的更多信息，请参阅 CUDA C++ 编程指南中的到达/等待屏障部分。

1.4.1.4. Warp 级别对归约操作的支持

NVIDIA Ampere GPU 架构为 32 位有符号和无符号整数操作数增加了对 Warp 范围归约操作的本机支持。Warp 范围归约操作支持对 32 位有符号和无符号整数的算术 add、min 和 max 操作，以及对 32 位无符号整数的按位 and、or 和 xor 操作。

有关新的 Warp 范围归约操作的更多详细信息，请参阅 CUDA C++ 编程指南中的 Warp 归约函数。

1.4.1.5. 改进的 Tensor Core 操作

NVIDIA Ampere GPU 架构包括新的第三代 Tensor Core，它比 Volta 和 Turing SM 中使用的 Tensor Core 更强大。新的 Tensor Core 使用更大的基本矩阵大小，并增加了强大的新数学模式，包括

支持 FP64 Tensor Core，使用新的 DMMA 指令。
通过 HMMA 指令支持 Bfloat16 Tensor Core。BFloat16 格式对于 DL 训练场景尤其有效。Bfloat16 提供 8 位指数（即与 FP32 相同的范围）、7 位尾数和 1 位符号位。
通过 HMMA 指令支持 TF32 Tensor Core。TF32 是一种新的 19 位 Tensor Core 格式，可以轻松集成到程序中，以实现比 16 位 HMMA 格式更准确的 DL 训练。TF32 提供 8 位指数、10 位尾数和 1 位符号位。
支持按位 AND 以及 Turing 中引入的按位 XOR，通过 BMMA 指令。

下表介绍了跨不同 GPU 架构代的 Tensor Core 的矩阵指令大小和支持数据类型的演变。

指令	GPU 架构	输入矩阵格式	输出累加器格式	矩阵指令大小 (MxNxK)
HMMA（16 位精度）	NVIDIA Volta 架构	FP16	FP16 / FP32	8x8x4
	NVIDIA Turing 架构	FP16	FP16 / FP32	8x8x4 / 16x8x8 / 16x8x16
	NVIDIA Ampere 架构	FP16 / BFloat16	FP16 / FP32（BFloat16 仅支持 FP32 作为累加器）	16x8x8 / 16x8x16
HMMA（19 位精度）	NVIDIA Volta 架构	不适用	不适用	不适用
	NVIDIA Turing 架构	不适用	不适用	不适用
	NVIDIA Ampere 架构	TF32（19 位）	FP32	16x8x4
IMMA（整数 MMA）	NVIDIA Volta 架构	不适用	不适用	不适用
	NVIDIA Turing 架构	unsigned char/signed char（8 位精度）	int32	8x8x16
	NVIDIA Ampere 架构	unsigned char/signed char（8 位精度）	int32	8x8x16 / 16x8x16 / 16x8x32
IMMA（整数子字节 MMA）	NVIDIA Volta 架构	不适用	不适用	不适用
	NVIDIA Turing 架构	unsigned u4/signed u4（4 位精度）	int32	8x8x32
	NVIDIA Ampere 架构	unsigned u4/signed u4（4 位精度）	int32	8x8x32 / 16x8x32 / 16x8x64
BMMA（二进制 MMA）	NVIDIA Volta 架构	不适用	不适用	不适用
	NVIDIA Turing 架构	单比特	int32	8x8x128
	NVIDIA Ampere 架构	单比特	int32	8x8x128 / 16x8x128 / 16x8x256
DMMA（64 位精度）	NVIDIA Volta 架构	不适用	不适用	不适用
	NVIDIA Turing 架构	不适用	不适用	不适用
	NVIDIA Ampere 架构	FP64	FP64	8x8x4

有关新的 Tensor Core 操作的更多详细信息，请参阅 CUDA C++ 编程指南中的 Warp 矩阵乘法部分。

1.4.1.6. 改进的 FP32 吞吐量

计算能力为 8.6 的设备的每个 SM 的每个周期 FP32 操作数是计算能力为 8.0 的设备的 2 倍。虽然为 8.0 编译的二进制文件将在 8.6 上按原样运行，但建议显式为 8.6 编译以受益于增加的 FP32 吞吐量。

1.4.2. 内存系统

1.4.2.1. 增加的内存容量和高带宽内存

NVIDIA A100 GPU 将 HBM2 内存容量从 V100 GPU 中的 32 GB 增加到 A100 GPU 中的 40 GB。除了增加的内存容量外，带宽也增加了 72%，从 Volta V100 上的 900 GB/s 增加到 A100 上的 1550 GB/s。

1.4.2.2. 增加的 L2 容量和 L2 驻留控制

NVIDIA Ampere GPU 架构将 Tesla A100 中的 L2 缓存容量增加到 40 MB，比 Tesla V100 大 7 倍。除了增加的容量外，L2 缓存到 SM 的带宽也增加了。NVIDIA Ampere GPU 架构允许 CUDA 用户控制 L2 缓存中数据的持久性。有关 L2 缓存中数据持久性的更多信息，请参阅 CUDA C++ 编程指南中关于管理 L2 缓存的部分。

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

基于计算能力 8.0 的 NVIDIA A100 GPU 将组合的 L1 缓存、纹理缓存和共享内存的最大容量增加到 192 KB，比 NVIDIA V100 GPU 中的 L1 缓存大 50%。计算能力为 8.6 的 GPU 的组合 L1 缓存容量为 128 KB。

在 NVIDIA Ampere GPU 架构中，专用于共享内存的 L1 缓存部分（称为划分区）可以在运行时选择，就像 Volta 等以前的架构一样，使用带有属性 cudaFuncAttributePreferredSharedMemoryCarveout 的 cudaFuncSetAttribute()。NVIDIA A100 GPU 支持每个 SM 0、8、16、32、64、100、132 或 164 KB 的共享内存容量。计算能力为 8.6 的 GPU 支持每个 SM 0、8、16、32、64 或 100 KB 的共享内存容量。

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

与 Volta 类似，NVIDIA Ampere GPU 架构将 L1 和纹理缓存的功能组合到统一的 L1/纹理缓存中，该缓存充当内存访问的合并缓冲区，在将数据传递到 Warp 之前收集 Warp 线程请求的数据。它与共享内存结合的另一个好处类似于 Volta L1，即在延迟和带宽方面都有所改进。

1.4.3. 第三代 NVLink

第三代 NVIDIA 高速 NVLink 互连在 A100 GPU 中实现，通过每个 GPU 更多的链路、更快的通信带宽以及改进的错误检测和恢复功能，显着增强了多 GPU 的可扩展性、性能和可靠性。第三代 NVLink 的每个链路具有相同的 50 GB/s 双向数据速率，但使用一半的信号对来实现此带宽。因此，A100 中可用的链路总数增加到 12 个，而 V100 中为 6 个，从而产生了 600 GB/s 的双向带宽，而 V100 为 300 GB/s。

NVLink 在现有的 CUDA 模型中透明地运行。NVLink 连接的端点之间的传输自动通过 NVLink 而不是 PCIe 路由。cudaDeviceEnablePeerAccess() API 调用仍然是启用 GPU 之间直接传输（通过 PCIe 或 NVLink）所必需的。cudaDeviceCanAccessPeer() 可用于确定任何一对 GPU 之间是否可以进行对等访问。

在 NVIDIA Ampere GPU 架构中，远程 NVLINK 访问会通过远程 GPU 上的 Link TLB。此 Link TLB 可以访问远程 GPU 内存的 64 GB。具有远程随机访问的应用程序可能希望将远程访问区域限制为每个对等 GPU 64 GB。

2. 修订历史

版本 1.1

初始公开发布
增加了对计算能力 8.6 的支持

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 在美国和其他国家/地区的商标或注册商标。其他公司和产品名称可能是与其相关联的各自公司的商标。

1: 在本指南中，Kepler 指的是计算能力为 3.x 的设备，Maxwell 指的是计算能力为 5.x 的设备，Pascal 指的是计算能力为 6.x 的设备，Volta 指的是计算能力为 7.0 的设备，Turing 指的是计算能力为 7.5 的设备，NVIDIA Ampere GPU 架构 指的是计算能力为 8.x 的设备