为NVIDIA安培GPU架构优化CUDA应用程序

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

1. NVIDIA安培GPU架构调优指南

1.1. NVIDIA安培GPU架构

NVIDIA Ampere GPU架构是NVIDIA为CUDA计算应用推出的最新架构。该架构保留并扩展了由图灵(Turing)和伏特(Volta)等前代NVIDIA GPU架构提供的相同CUDA编程模型,遵循这些架构最佳实践的应用程序通常无需任何代码修改即可在NVIDIA A100 GPU上获得加速。本指南总结了如何通过利用NVIDIA Ampere GPU架构的特性来微调应用程序以获得额外加速。1

如需了解本指南中讨论的编程功能的更多详细信息,请参阅CUDA C++ Programming Guide

1.2. CUDA最佳实践

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

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

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

  • 最小化主机与设备之间的数据传输。

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

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

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

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

1.3. 应用兼容性

在解决本指南涵盖的具体性能调优问题之前,请参阅NVIDIA Ampere GPU架构CUDA应用兼容性指南,以确保您的应用程序以兼容NVIDIA Ampere GPU架构的方式进行编译。

1.4. NVIDIA安培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安培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位有符号和无符号整数的算术运算addminmax,以及对32位无符号整数的位运算andorxor

有关新的warp范围归约操作的更多详情,请参阅CUDA C++编程指南中的Warp Reduce Functions章节。

1.4.1.5. 增强的张量核心运算

NVIDIA Ampere GPU架构包含了全新的第三代Tensor Cores,其性能比Volta和Turing SM中使用的Tensor Cores更强大。新的Tensor Cores采用了更大的基础矩阵尺寸,并新增了强大的数学运算模式,包括:

  • 支持FP64 Tensor Core,使用新的DMMA指令。

  • 支持通过HMMA指令使用Bfloat16张量核心。BFloat16格式特别适合深度学习训练场景。Bfloat16提供8位指数(与FP32相同范围)、7位尾数和1位符号位。

  • 支持通过HMMA指令使用TF32 Tensor Core。TF32是一种新的19位Tensor Core格式,可以轻松集成到程序中,相比16位HMMA格式能提供更精确的深度学习训练。TF32提供8位指数、10位尾数和1位符号位。

  • 支持通过BMMA指令实现按位AND运算,以及图灵架构中引入的按位XOR运算。

下表展示了不同GPU架构世代中Tensor Core矩阵指令大小和支持数据类型的演进历程。

指令

GPU架构

输入矩阵格式

输出累加器格式

矩阵指令大小 (MxNxK)

HMMA (16位精度)

NVIDIA Volta架构

FP16

FP16 / FP32

8x8x4

NVIDIA图灵架构

FP16

FP16 / FP32

8x8x4 / 16x8x8 / 16x8x16

NVIDIA安培架构

FP16 / BFloat16

FP16 / FP32 (BFloat16仅支持FP32作为累加器)

16x8x8 / 16x8x16

HMMA (19位精度)

NVIDIA Volta架构

N/A

N/A

N/A

NVIDIA图灵架构

N/A

N/A

N/A

NVIDIA安培架构

TF32 (19位)

FP32

16x8x4

IMMA (整数MMA)

NVIDIA Volta架构

N/A

N/A

N/A

NVIDIA图灵架构

无符号字符/有符号字符(8位精度)

int32

8x8x16

NVIDIA安培架构

无符号字符/有符号字符(8位精度)

int32

8x8x16 / 16x8x16 / 16x8x32

IMMA (整数子字节MMA)

NVIDIA Volta架构

N/A

N/A

N/A

NVIDIA Turing架构

无符号u4/有符号u4 (4位精度)

int32

8x8x32

NVIDIA安培架构

无符号u4/有符号u4(4位精度)

int32

8x8x32 / 16x8x32 / 16x8x64

BMMA (二进制MMA)

NVIDIA Volta架构

N/A

N/A

N/A

NVIDIA Turing架构

单比特

int32

8x8x128

NVIDIA安培架构

单比特

int32

8x8x128 / 16x8x128 / 16x8x256

DMMA (64位精度)

NVIDIA Volta架构

N/A

N/A

N/A

NVIDIA图灵架构

N/A

N/A

N/A

NVIDIA安培架构

FP64

FP64

8x8x4

有关新Tensor Core运算的更多详情,请参阅CUDA C++编程指南中的Warp矩阵乘法章节。

1.4.1.6. 提升FP32吞吐量

计算能力8.6的设备每个SM每个周期的FP32运算量是8.0设备的两倍。虽然为8.0编译的二进制文件可以直接在8.6上运行,但建议显式地为8.6进行编译,以获得更高的FP32吞吐量优势。

1.4.2. 内存系统

1.4.2.1. 提升内存容量与高带宽内存

NVIDIA A100 GPU将HBM2内存容量从V100 GPU的32GB提升至A100的40GB。在增加内存容量的同时,带宽也提升了72%,从Volta V100的900GB/s增加到A100的1550GB/s。

1.4.2.2. 提升的L2缓存容量与L2驻留控制

NVIDIA安培GPU架构将Tesla A100的L2缓存容量提升至40MB,比Tesla V100大7倍。在容量增加的同时,L2缓存与流式多处理器(SMs)之间的带宽也得到提升。NVIDIA安培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安培GPU架构中,与Volta等前代架构类似,L1缓存中专门用于共享内存的部分(称为carveout)可以在运行时通过cudaFuncSetAttribute()函数配合cudaFuncAttributePreferredSharedMemoryCarveout属性进行选择。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++ Programming Guide

与Volta类似,NVIDIA安培GPU架构将L1缓存和纹理缓存的功能整合为统一的L1/纹理缓存,作为内存访问的合并缓冲区,在将数据传递给线程束之前,先收集该线程束中线程所请求的数据。与Volta L1类似,其与共享内存结合的另一个优势是在延迟和带宽方面都有所改善。

2. 版本历史

版本 1.1

  • 首次公开发布

  • 新增对计算能力8.6的支持

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的设备,Volta指计算能力7.0的设备,Turing指计算能力7.5的设备,NVIDIA Ampere GPU架构指计算能力8.x的设备