CUDA C++ 最佳实践指南
使用CUDA工具包从NVIDIA GPU获得最佳性能的编程指南。
1. 前言
本最佳实践指南是一本帮助开发者从NVIDIA® CUDA® GPU获取最佳性能的手册。它介绍了成熟的并行化与优化技术，并阐释了能够极大简化支持CUDA的GPU架构编程的编码范式与惯用法。
虽然本内容可作为参考手册使用，但请注意随着探讨不同编程和配置主题，某些概念会在不同上下文中重复出现。因此建议首次阅读的用户按顺序阅读本指南。这种阅读方式将极大提升您对高效编程实践的理解，并使您今后能更好地将本指南作为参考资料使用。
1.1. 本指南适合哪些读者？
本指南中的所有讨论均使用C++编程语言，因此您应能熟练阅读C++代码。
本指南参考并依赖于您应备有的其他几份文档，所有这些文档均可从CUDA网站https://docs.nvidia.com/cuda/免费获取。以下文档是特别重要的资源：
CUDA安装指南
CUDA C++ 编程指南
CUDA工具包参考手册
具体而言，本指南的优化部分假设您已成功下载并安装CUDA Toolkit（若尚未完成，请参考适用于您平台的相关CUDA安装指南），且对CUDA C++编程语言和环境有基本了解（若尚未掌握，请参阅CUDA C++编程指南）。
1.2. 评估、并行化、优化、部署
本指南介绍了评估、并行化、优化、部署(APOD)设计周期，旨在帮助应用程序开发者快速识别其代码中最能从GPU加速中获益的部分，迅速实现这一优势，并尽早开始在生产环境中利用由此带来的速度提升。
APOD是一个循环迭代的过程：只需投入少量初始时间即可实现、测试并部署初步的加速效果，随后该循环可重新启动，通过识别更多优化机会、获得额外加速收益，最终将更快的应用版本部署到生产环境。
1.2.1. 评估
对于一个现有项目，第一步是评估应用程序，以定位代码中执行时间占比最大的部分。掌握这些信息后，开发者可以评估这些瓶颈是否适合并行化，并开始研究GPU加速。
通过理解最终用户的需求和限制条件，并应用阿姆达尔定律和古斯塔夫森定律，开发者可以确定应用程序中已识别部分加速后性能提升的上限。
1.2.2. 并行化
在识别出热点并完成设定目标和预期的基础练习后，开发者需要对代码进行并行化处理。根据原始代码的情况，这可能简单到只需调用现有的GPU优化库如cuBLAS、cuFFT或Thrust，也可能简单到只需添加几条预处理器指令作为并行化编译器的提示。
另一方面，某些应用程序的设计需要进行一定程度的重构以暴露其内在并行性。由于即使是CPU架构也需要通过暴露并行性来提升或维持串行应用的性能，CUDA系列并行编程语言（CUDA C++、CUDA Fortran等）致力于使这种并行性的表达尽可能简单，同时支持在专为最大并行吞吐量设计的CUDA兼容GPU上运行。
1.2.3. 优化
在每一轮应用并行化完成后，开发者可以转向优化实现以提升性能。由于存在大量可考虑的优化方案，充分理解应用程序需求有助于使优化过程尽可能顺畅。不过，正如整个APOD流程一样，程序优化是一个迭代过程（识别优化机会、应用并测试优化、验证达到的加速效果，然后重复），这意味着开发人员无需在获得良好加速效果之前，花费大量时间记忆所有可能的优化策略。相反，可以在学习过程中逐步应用这些策略。
优化可以在多个层面进行，从重叠数据传输与计算，一直到微调浮点运算序列。可用的性能分析工具在此过程中极为宝贵，它们能帮助开发者建议下一步最佳优化方向，并提供本指南优化章节相关部分的参考。
1.2.4. 部署
在完成应用程序一个或多个组件的GPU加速后，可以将结果与原始预期进行比较。回想一下，最初的评估步骤允许开发者确定通过加速给定热点可能获得的最大潜在加速上限。
在着手解决其他热点问题以提高整体加速效果之前，开发者应考虑将部分并行化的实现方案推进到生产环境。这样做有多方面重要意义：例如，它能让用户尽早从投资中获益（即使是部分加速也极具价值），同时通过采用渐进式而非颠覆性的应用改造方案，为开发者和用户将风险降至最低。
1.3. 推荐与最佳实践
在本指南中，我们针对CUDA C++代码的设计与实现提供了具体建议。这些建议按优先级分类，优先级综合考量了建议的效果和适用范围。能为大多数CUDA应用带来显著改进的措施具有最高优先级，而仅影响特定场景的细微优化则被赋予较低优先级。
在实施较低优先级的建议之前，最好确保所有相关的高优先级建议都已被应用。这种方法往往能以投入的时间获得最佳效果，并避免过早优化的陷阱。
确定优先级的收益和范围标准会根据程序的性质而有所不同。在本指南中，它们代表了一个典型情况。您的代码可能反映了不同的优先级因素。无论存在何种可能性，在着手处理低优先级项目之前，最好先确认没有遗漏更高优先级的建议。
注意
为了简洁起见，本指南中的代码示例省略了错误检查。然而，生产代码应系统地检查每个API调用返回的错误代码，并通过调用cudaGetLastError()来检查内核启动是否失败。
1.4. 评估您的应用程序
从超级计算机到移动电话，现代处理器越来越依赖并行性来提供性能。核心计算单元（包括控制、算术、寄存器和通常的一些缓存）被复制多次，并通过网络连接到内存。因此，所有现代处理器都需要并行代码才能充分利用其计算能力。
尽管处理器正在不断发展，向程序员展示更细粒度的并行性，但许多现有应用程序仍以串行代码或粗粒度并行代码的形式演进（例如，数据被分解为并行处理的区域，子区域通过MPI共享）。为了从包括GPU在内的任何现代处理器架构中获益，首要步骤是评估应用程序以识别热点，确定它们是否可以并行化，并理解当前和未来的相关工作负载。
2. 异构计算
CUDA编程涉及在两个不同平台上并发运行代码：一个包含一个或多个CPU的主机系统，以及一个或多个支持CUDA的NVIDIA GPU设备。
虽然NVIDIA GPU通常与图形处理相关联，但它们同时也是强大的算术引擎，能够并行运行数千个轻量级线程。这种能力使其特别适合利用并行执行的计算任务。
然而，该设备基于与主机系统截然不同的设计架构，理解这些差异及其如何影响CUDA应用程序的性能表现，对于有效使用CUDA至关重要。
2.1. 主机与设备之间的差异
主要区别在于线程模型和独立的物理内存：
- Threading resources
-
主机系统上的执行流水线只能支持有限数量的并发线程。例如，配备两个32核处理器的服务器最多只能同时运行64个线程（如果CPU支持同步多线程，则可能略高于此数值）。相比之下，CUDA设备上最小的并行执行单元包含32个线程（称为线程束）。现代NVIDIA GPU每个多处理器可支持多达2048个并发活动线程（参见《CUDA C++编程指南》的功能与规格说明）。在具有80个多处理器的GPU上，这意味着能同时运行超过16万个活动线程。
- Threads
-
CPU上的线程通常是重量级实体。操作系统必须在CPU执行通道上切换线程以提供多线程能力。因此，上下文切换（当两个线程被交换时）速度慢且成本高。相比之下，GPU上的线程极其轻量级。在典型系统中，成千上万的线程排队等待工作（以每组32线程的线程束形式）。如果GPU必须等待一个线程束，它只需开始执行另一个线程束的工作。由于为所有活动线程分配了独立的寄存器，在GPU线程间切换时无需交换寄存器或其他状态。资源会一直分配给每个线程，直到其完成执行。简而言之，CPU核心设计用于最小化少量线程的延迟，而GPU设计用于处理大量并发轻量级线程以最大化吞吐量。
- RAM
-
主机系统和设备各自拥有独立的物理内存1。由于主机内存与设备内存是分离的，如CUDA设备上运行什么？所述，主机内存中的数据项偶尔需要在设备内存与主机内存之间进行传输。
这些是CPU主机与GPU设备在并行编程方面的主要硬件差异。本文档其他部分将讨论其他差异。考虑到这些差异设计的应用程序可以将主机和设备视为一个紧密协作的异构系统，其中每个处理单元都发挥其最擅长的功能：主机处理串行工作，而设备处理并行工作。
2.2. 什么能在支持CUDA的设备上运行？
在确定应用程序的哪些部分应在设备上运行时，应考虑以下问题：
该设备非常适合可同时在众多数据元素上并行运行的计算任务。这通常涉及对大型数据集(如矩阵)进行算术运算，其中可同时对成千上万(甚至数百万)个元素执行相同的操作。这是CUDA获得良好性能的必要条件：软件必须使用大量(通常数千或数万)的并发线程。支持并行运行大量线程的特性源自CUDA采用的上述轻量级线程模型。
-
要使用CUDA，必须将数据值从主机传输到设备。这些传输在性能方面代价高昂，应尽量减少。（参见主机与设备之间的数据传输。）这种成本会产生几个影响：
-
操作的复杂度应能证明数据在设备与主机之间传输的开销是合理的。仅少量线程短暂使用数据而进行传输的代码，其性能提升将微乎其微甚至没有。理想情况是大量线程执行大量计算工作。
例如，将两个矩阵传输到设备执行矩阵加法，再将结果传回主机，这种操作几乎不会带来性能提升。问题关键在于每次数据传输对应的计算操作数量。假设矩阵尺寸为NxN，上述过程中共执行N2次加法运算，同时传输了3N2个元素，计算操作与数据传输量之比为1:3即O(1)。当该比值更高时更容易获得性能优势。以相同矩阵的乘法运算为例，其需要N3次乘加运算，此时计算操作与数据传输量之比达到O(N)，矩阵越大性能收益越显著。此外，运算类型也是关键因素，例如三角函数与加法运算具有不同的复杂度特征。在决定操作应在主机还是设备上执行时，必须充分考虑数据传输带来的额外开销。
数据应尽可能长时间保留在设备上。由于需要尽量减少传输，对于同一数据运行多个内核的程序，应倾向于在内核调用之间将数据保留在设备上，而不是将中间结果传输到主机然后再发送回设备进行后续计算。因此，在前面的示例中，如果要相加的两个矩阵已经是之前某些计算的结果而存在于设备上，或者如果加法结果将用于某些后续计算，则矩阵加法应在设备本地执行。即使一系列计算中的某个步骤在主机上执行可能更快，也应采用这种方法。即使是一个相对较慢的内核，如果它能避免主机和设备内存之间的一次或多次传输，也可能是值得的。主机与设备之间的数据传输提供了更多细节，包括主机与设备之间以及设备内部带宽的测量数据。
-
为了获得最佳性能，设备上运行的相邻线程在内存访问时应保持一定的连贯性。某些内存访问模式能让硬件将多个数据项的读写操作合并为单次操作。如果数据无法按照支持合并访问的方式排布，或者缺乏足够的局部性来有效利用L1或纹理缓存，那么在GPU计算中往往难以获得显著的加速效果。值得注意的是，完全随机的内存访问模式是个例外。通常应当避免这种模式，因为任何架构处理这类随机访问模式的效率都远低于峰值性能。但与基于缓存的架构（如CPU）相比，具备延迟隐藏特性的GPU架构在处理完全随机内存访问模式时往往表现更优。
- 1
-
在集成GPU的系统级芯片(SoC)上，例如NVIDIA® Tegra®，主机内存和设备内存在物理上是相同的，但在逻辑上仍存在主机内存与设备内存的区别。详情请参阅Tegra平台CUDA应用指南。
3. 应用性能分析
3.1. 性能分析
许多代码通过相对少量的代码完成了工作的主要部分。开发者可以使用性能分析工具识别这些热点区域，并开始列出并行化的候选目标。
3.1.1. 创建配置文件
分析代码性能有许多可能的方法，但所有情况下的目标都是相同的：确定应用程序在执行过程中花费最多时间的函数。
注意
高优先级： 为了最大化开发者的生产力，需要对应用程序进行分析以确定热点和瓶颈。
任何性能分析活动中最重要的考量是确保工作负载真实可靠——也就是说，从测试中获得的信息以及基于该信息做出的决策都要与实际数据相关。使用不切实际的工作负载可能导致次优结果和精力浪费，既会让开发者针对不切实际的问题规模进行优化，也会使开发者将注意力集中在错误的函数上。
有多种工具可用于生成性能分析报告。以下示例基于gprof，这是GNU Binutils集合中适用于Linux平台的开源性能分析工具。
$ gcc -O2 -g -pg myprog.c
$ gprof ./a.out > profile.txt
Each sample counts as 0.01 seconds.
% cumulative self self total
time seconds seconds calls ms/call ms/call name
33.34 0.02 0.02 7208 0.00 0.00 genTimeStep
16.67 0.03 0.01 240 0.04 0.12 calcStats
16.67 0.04 0.01 8 1.25 1.25 calcSummaryData
16.67 0.05 0.01 7 1.43 1.43 write
16.67 0.06 0.01 mcount
0.00 0.06 0.00 236 0.00 0.00 tzset
0.00 0.06 0.00 192 0.00 0.00 tolower
0.00 0.06 0.00 47 0.00 0.00 strlen
0.00 0.06 0.00 45 0.00 0.00 strchr
0.00 0.06 0.00 1 0.00 50.00 main
0.00 0.06 0.00 1 0.00 0.00 memcpy
0.00 0.06 0.00 1 0.00 10.11 print
0.00 0.06 0.00 1 0.00 0.00 profil
0.00 0.06 0.00 1 0.00 50.00 report
3.1.2. 识别热点区域
在上面的示例中，我们可以清楚地看到函数genTimeStep()占用了应用程序总运行时间的三分之一。这应该是我们进行并行化的首要候选函数。Understanding Scaling讨论了这种并行化可能带来的潜在收益。
值得注意的是，上述示例中的其他几个函数也占据了整体运行时间的很大一部分，例如calcStats()和calcSummaryData()。并行化这些函数也应能提高我们的加速潜力。不过，由于APOD是一个循环过程，我们可能会选择在后续的APOD阶段并行化这些函数，从而将我们在任何特定阶段的工作范围限制在较小的增量变更集内。
3.1.3. 理解扩展性
应用程序通过运行在CUDA上能获得的性能提升程度，完全取决于其可并行化的程度。无法充分并行化的代码应该在主机上运行，除非这样做会导致主机与设备之间过多的数据传输。
注意
高优先级：为了从CUDA获得最大收益，首先应专注于寻找并行化顺序代码的方法。
通过理解应用程序如何扩展，可以设定预期并规划增量并行化策略。强扩展与阿姆达尔定律描述了强扩展，这使我们能够为固定问题规模下的加速设定上限。弱扩展与古斯塔夫森定律则描述了弱扩展，即通过增大问题规模来获得加速。在许多应用中，强扩展和弱扩展的结合是可取的。
3.1.3.1. 强扩展性与阿姆达尔定律
强扩展性衡量的是在固定总体问题规模的情况下，随着系统增加更多处理器，求解时间如何减少。展现线性强扩展性的应用程序，其加速比等于所用处理器的数量。
强扩展性通常等同于阿姆达尔定律，该定律规定了通过并行化串行程序的部分所能预期的最大加速比。本质上，它表明程序的最大加速比S为：
\(S = \frac{1}{(1 - P) + \frac{P}{N}}\)
这里P表示可并行化代码部分占总串行执行时间的比例，N表示运行代码并行部分的处理器数量。
当N越大（即处理器数量越多），P/N的比例就越小。可以简单地将N视为一个非常大的数字，这样方程就转化为\(S = 1/(1 - P)\)。现在，如果一个顺序程序的运行时间有3/4被并行化，那么相对于串行代码的最大加速比为1 / (1 - 3/4) = 4。
实际上，大多数应用程序并不能呈现完美的线性强扩展性，即便它们确实展现出某种程度的强扩展。在多数情况下，关键在于可并行部分P越大，潜在的加速效果就越显著。反之，如果P数值较小（意味着应用程序不具备实质性的可并行性），增加处理器数量N对性能提升几乎没有帮助。因此，为了在固定问题规模下获得最大加速比，值得投入精力增大P值，即最大化可并行化的代码量。
3.1.3.2. 弱扩展性与古斯塔夫森定律
弱扩展性衡量的是在保持每个处理器的问题规模固定的情况下，随着系统增加更多处理器，求解时间如何变化；也就是说，随着处理器数量的增加，整体问题规模也会相应增大。
弱扩展性通常等同于古斯塔夫森定律(Gustafson's Law)，该定律指出在实际应用中，问题规模会随着处理器数量的增加而扩大。因此，程序的最大加速比S为：
\(S = N + (1 - P)(1 - N)\)
这里P表示可并行化代码部分占总串行执行时间的比例，N表示运行该并行代码部分的处理器数量。
理解Gustafson定律的另一种视角是：当我们扩展系统规模时，保持恒定的不是问题规模，而是执行时间。需要注意的是，Gustafson定律假设串行与并行执行的比例保持恒定，这反映了处理更大问题时所需的额外设置和管理成本。
3.1.3.3. 应用强扩展与弱扩展
理解哪种扩展类型最适合某个应用是评估加速效果的重要部分。对于某些应用，问题规模将保持不变，因此仅适用强扩展。例如模拟两个分子如何相互作用的情况，其中分子大小是固定的。
对于其他应用场景，问题规模会扩大以充分利用可用的处理器资源。例如，将流体或结构建模为网格或格点，以及某些蒙特卡洛模拟，在这些情况下，增大问题规模可以提高计算精度。
在理解了应用性能分析后，开发者应当了解如果计算性能发生变化时问题规模会如何改变，然后运用阿姆达尔定律或古斯塔夫森定律来确定加速比的上限。
4. 并行化您的应用程序
在识别出热点并完成设定目标和预期的基础练习后，开发者需要对代码进行并行化处理。根据原始代码的情况，这可能简单到只需调用现有的GPU优化库如cuBLAS、cuFFT或Thrust，也可能简单到只需添加几条预处理器指令作为并行化编译器的提示。
另一方面，某些应用程序的设计需要进行一定程度的重构才能展现其内在的并行性。正如CPU架构需要通过暴露这种并行性来提升或维持串行应用程序的性能一样，CUDA系列并行编程语言（CUDA C++、CUDA Fortran等）致力于让这种并行性的表达尽可能简单，同时支持在专为最大并行吞吐量设计的CUDA兼容GPU上运行。
5. 入门指南
将顺序代码并行化有几种关键策略。虽然如何将这些策略应用于特定应用程序的细节是一个复杂且与问题相关的话题，但这里列出的一般主题适用于我们是将代码并行化以在多核CPU上运行，还是在CUDA GPU上使用。
5.1. 并行库
并行化应用程序最直接的方法是利用现有的库，这些库能代表我们充分利用并行架构。CUDA工具包包含许多为NVIDIA CUDA GPU精心优化的此类库，例如cuBLAS、cuFFT等。
关键在于，当库与应用程序的需求高度匹配时，它们才最具价值。例如，已经使用其他BLAS库的应用程序通常可以轻松切换到cuBLAS，而几乎不涉及线性代数的应用程序则很少会用到cuBLAS。其他CUDA工具包库也是如此：cuFFT的接口与FFTW类似，等等。
同样值得注意的是Thrust库，这是一个类似于C++标准模板库的并行C++模板库。Thrust提供了丰富的数据并行原语集合，如扫描(scan)、排序(sort)和归约(reduce)，这些原语可以组合在一起，用简洁易读的源代码实现复杂算法。通过使用这些高级抽象来描述计算，您赋予Thrust自动选择最高效实现的自由。因此，Thrust既可用于CUDA应用程序的快速原型开发（此时程序员的生产力最为关键），也可用于生产环境（此时健壮性和绝对性能至关重要）。
5.2. 并行化编译器
另一种常见的序列代码并行化方法是利用并行化编译器。这通常意味着采用基于指令的方法，程序员使用pragma或其他类似符号向编译器提示可以找到并行性的位置，而无需修改或调整底层代码本身。通过向编译器暴露并行性，指令允许编译器执行将计算映射到并行架构上的详细工作。
OpenACC标准提供了一组编译器指令，用于指定标准C、C++和Fortran中应卸载到附加加速器（如CUDA GPU）的循环和代码区域。加速器设备的管理细节由支持OpenACC的编译器和运行时隐式处理。
详情请参阅 http://www.openacc.org/。
5.3. 编码以暴露并行性
对于需要超越现有并行库或并行化编译器所能提供的额外功能或性能的应用，与现有串行代码无缝集成的并行编程语言（如CUDA C++）至关重要。
一旦我们在应用程序的性能分析中定位到热点区域，并确定自定义代码是最佳解决方案，就可以使用CUDA C++将该代码段的并行性暴露为CUDA内核。随后我们可以将这个内核部署到GPU上并获取结果，而无需对应用程序的其他部分进行大规模重写。
当应用程序的大部分总运行时间集中在代码中几个相对独立的部分时，这种方法最为直接。而对于具有非常扁平性能分析的应用（即时间消耗相对均匀分布在代码库的广泛部分中），并行化则更为困难。对于后一类应用，可能需要进行一定程度的代码重构以暴露应用中固有的并行性，但请记住，这种重构工作往往会使所有未来架构（CPU和GPU）都受益，因此如果确实必要，这些努力是非常值得的。
6. 获取正确答案
获取正确答案显然是所有计算的主要目标。在并行系统中，可能会遇到传统串行编程中不常见的困难。这些问题包括线程问题、由于浮点值计算方式导致的意外结果，以及CPU和GPU处理器运行方式差异带来的挑战。本章将探讨可能影响返回数据正确性的问题，并提供相应的解决方案。
6.1. 验证
6.1.1. 参考比较
验证对现有程序修改正确性的一个关键方面是建立一种机制，将先前已知良好的代表性输入参考输出与新结果进行比较。每次更改后，确保结果符合适用于特定算法的任何标准。有些算法期望位级完全相同的结果，但这并不总是可能的，尤其是在涉及浮点运算时；关于数值精度，请参阅数值准确性与精度。对于其他算法，如果实现结果与参考值在某个小误差范围内匹配，则可以认为其正确。
需要注意的是，用于验证数值结果的过程可以轻松扩展到性能验证领域。我们希望确保每个修改不仅正确无误，而且能提升性能（以及提升幅度）。将这些检查作为APOD循环流程的有机组成部分频繁执行，将帮助我们尽可能快速地达成预期目标。
6.1.2. 单元测试
与上述参考比较方法相辅相成的是，将代码本身结构化，使其在单元级别易于验证。例如，我们可以将CUDA内核编写为一系列短小的__device__函数集合，而不是一个庞大的__global__函数；每个设备函数在整合前都可以独立测试。
例如，许多内核除了实际计算外，还包含复杂的内存寻址逻辑。如果我们在引入主要计算之前先单独验证寻址逻辑，这将简化后续的调试工作。（请注意，CUDA编译器会将任何未对全局内存写入产生贡献的设备代码视为可消除的死代码，因此我们必须至少通过寻址逻辑向全局内存写入一些内容才能成功应用此策略。）
更进一步来说，如果将大多数函数定义为__host__ __device__而不仅仅是__device__函数，那么这些函数就可以同时在CPU和GPU上进行测试，从而增强我们对函数正确性的信心，并确保结果不会出现任何意外差异。如果确实存在差异，这些差异就能在早期被发现，并且可以在简单函数的上下文中被理解。
作为一项有益的副作用，如果我们希望在应用程序中同时包含CPU和GPU执行路径，这种策略将为我们提供减少代码重复的手段：如果CUDA内核的大部分工作都在__host__ __device__函数中完成，我们就可以轻松地从主机代码和设备代码中调用这些函数而无需重复。
6.2. 调试
CUDA-GDB是GNU调试器的一个移植版本，可在Linux和Mac系统上运行；详见：https://developer.nvidia.com/cuda-gdb。
NVIDIA Nsight Visual Studio Edition可作为Microsoft Visual Studio的免费插件使用；详见：https://developer.nvidia.com/nsight-visual-studio-edition。
多个第三方调试器也支持CUDA调试；详情请参阅：https://developer.nvidia.com/debugging-solutions
6.3. 数值精度与准确度
错误或意外结果主要源于浮点数精度问题，这是由于浮点数值的计算和存储方式导致的。以下部分将解释主要关注点。关于浮点运算的其他特性，请参阅《CUDA C++编程指南》的功能与技术规范，以及NVIDIA开发者官网上提供的关于浮点精度与性能的白皮书和配套网络研讨会https://developer.nvidia.com/content/precision-performance-floating-point-and-ieee-754-compliance-nvidia-gpus。
6.3.1. 单精度与双精度对比
支持CUDA Compute Capability 1.3及以上版本的设备原生提供双精度浮点值（即64位宽度）支持。由于双精度运算具有更高精度和舍入误差特性，其计算结果往往与单精度算术运算结果存在差异。因此必须确保比较相同精度的数值，并将结果允许误差范围纳入考量，而非期待完全精确的匹配。
6.3.2. 浮点运算不具备结合律
每个浮点算术运算都涉及一定程度的舍入误差。因此，算术运算的执行顺序至关重要。如果A、B和C是浮点值，(A+B)+C不能保证等于A+(B+C)，就像符号数学中那样。当您并行化计算时，可能会改变运算顺序，因此并行结果可能与顺序结果不匹配。这个限制并非CUDA特有，而是浮点值并行计算固有的特性。
6.3.3. IEEE 754 合规性
所有CUDA计算设备都遵循IEEE 754二进制浮点数表示标准，但存在一些细微差异。这些差异在《CUDA C++编程指南》的功能与技术规范章节中有详细说明，可能导致计算结果与主机系统上计算的IEEE 754标准值存在偏差。
其中一个关键区别是融合乘加(FMA)指令，它将乘加运算合并为单条指令执行。其结果通常与分别执行两个运算得到的结果略有不同。
6.3.4. x86 80位浮点运算
x86处理器在执行浮点计算时可以使用80位双扩展精度数学运算。这些计算结果常常与CUDA设备上执行的纯64位运算结果不同。为了使数值更接近匹配，可以将x86主机处理器设置为使用常规双精度或单精度（分别为64位和32位）。这可以通过FLDCWx86汇编指令或等效的操作系统API来实现。
7. 优化CUDA应用
在每一轮应用并行化完成后，开发者可以转向优化实现以提升性能。由于存在大量可考虑的优化方案，充分理解应用程序需求有助于使优化过程尽可能顺畅。不过，正如整个APOD流程一样，程序优化是一个迭代过程（识别优化机会、应用并测试优化、验证达到的加速效果，然后重复），这意味着开发人员无需在获得良好加速效果之前，花费大量时间记忆所有可能的优化策略。相反，可以在学习过程中逐步应用这些策略。
优化可以在多个层面进行，从重叠数据传输与计算，一直到微调浮点运算序列。可用的性能分析工具在此过程中极为宝贵，它们能帮助开发者建议下一步最佳优化方向，并提供本指南优化章节相关部分的参考。
8. 性能指标
在尝试优化CUDA代码时，了解如何准确测量性能以及理解带宽在性能评估中的作用至关重要。本章将讨论如何使用CPU计时器和CUDA事件正确测量性能，然后探讨带宽如何影响性能指标，以及如何缓解其带来的一些挑战。
8.1. 计时
可以使用CPU或GPU计时器对CUDA调用和内核执行进行计时。本节探讨这两种方法的功能、优势和潜在问题。
8.1.1. 使用CPU计时器
任何CPU计时器都可以用来测量CUDA调用或内核执行的耗时。各种CPU计时方法的细节不在本文档讨论范围内，但开发者始终需要注意其计时调用所提供的精度。
使用CPU计时器时，必须牢记许多CUDA API函数是异步的；也就是说，它们在完成工作之前就将控制权返回给调用CPU线程。所有内核启动都是异步的，名称带有Async后缀的内存复制函数也是如此。因此，要准确测量特定调用或一系列CUDA调用的耗时，必须在启动和停止CPU计时器之前立即调用cudaDeviceSynchronize()来同步CPU线程与GPU。cudaDeviceSynchronize()会阻塞调用CPU线程，直到该线程之前发出的所有CUDA调用都完成为止。
虽然也可以将CPU线程与GPU上的特定流或事件同步，但这些同步函数不适合用于测量默认流之外其他流中的代码执行时间。cudaStreamSynchronize()会阻塞CPU线程，直到先前提交到给定流的所有CUDA调用都完成。cudaEventSynchronize()则会阻塞直到GPU记录了特定流中的给定事件。由于驱动程序可能会交错执行来自其他非默认流的CUDA调用，因此计时结果可能包含其他流中的调用。
由于默认流（stream 0）对设备上的工作表现出串行化行为（默认流中的操作只有在任何流中所有先前调用完成后才能开始；且任何流中的后续操作必须等待该操作完成后才能开始），这些函数可以可靠地用于默认流中的计时。
请注意，本节提到的CPU到GPU同步点意味着GPU处理管道的停滞，因此应谨慎使用以最小化其对性能的影响。
8.1.2. 使用CUDA GPU计时器
CUDA事件API提供了创建和销毁事件、记录事件（包括时间戳）以及将时间戳差异转换为毫秒级浮点数值的调用。如何使用CUDA事件进行代码计时展示了其用法。
如何使用CUDA事件进行代码计时
cudaEvent_t start, stop;
float time;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord( start, 0 );
kernel<<<grid,threads>>> ( d_odata, d_idata, size_x, size_y,
NUM_REPS);
cudaEventRecord( stop, 0 );
cudaEventSynchronize( stop );
cudaEventElapsedTime( &time, start, stop );
cudaEventDestroy( start );
cudaEventDestroy( stop );
这里使用cudaEventRecord()将start和stop事件放入默认流（流0）中。当设备在流中到达该事件时，会为事件记录时间戳。cudaEventElapsedTime()函数返回记录start和stop事件之间经过的时间。该值以毫秒表示，分辨率约为半微秒。与本列表中的其他调用一样，它们的具体操作、参数和返回值在CUDA Toolkit参考手册中有详细说明。请注意，计时是基于GPU时钟测量的，因此计时分辨率与操作系统无关。
8.2. 带宽
带宽——数据传输的速率——是影响性能最重要的限制因素之一。几乎所有代码修改都应考虑其对带宽的影响。正如本指南内存优化章节所述，数据存储的内存选择、数据布局方式、访问顺序以及其他因素都会显著影响带宽。
为了准确测量性能，计算理论带宽和实际带宽非常有用。当后者远低于前者时，设计或实现细节可能会降低带宽，此时提高实际带宽应成为后续优化工作的首要目标。
注意
高优先级：在衡量性能和优化效果时，请将计算的有效带宽作为指标。
8.2.1. 理论带宽计算
理论带宽可以通过产品手册中的硬件规格计算得出。例如，NVIDIA Tesla V100采用HBM2（双倍数据速率）内存，其内存时钟频率为877 MHz，并具有4096位宽的内存接口。
使用这些数据项，NVIDIA Tesla V100的理论峰值内存带宽为898 GB/s：
\(\left. \left( 0.877 \times 10^{9} \right. \times (4096/8) \times 2 \right) \div 10^{9} = 898\text{GB/s}\)
在此计算中，内存时钟频率被转换为赫兹(Hz)，乘以接口位宽（除以8以将比特转换为字节），再乘以2以考虑双倍数据速率。最后，该乘积除以109将结果转换为GB/秒。
注意
某些计算在最终结果中使用10243而非109作为除数。在这种情况下，带宽将为836.4 GiB/s。为确保理论带宽与实际带宽的比较有效性，计算时必须使用相同的除数。
注意
在启用ECC的GDDR内存GPU上，可用DRAM容量会减少6.25%以存储ECC校验位。与禁用ECC的同款GPU相比，每次内存事务获取ECC校验位还会使有效带宽降低约20%，不过ECC对带宽的实际影响可能更高，具体取决于内存访问模式。而HBM2内存则提供专用ECC资源，可实现无开销的ECC保护。2
8.2.2. 有效带宽计算
有效带宽通过计时特定程序活动并了解程序如何访问数据来计算。为此，请使用以下公式：
\(\text{Effective\ bandwidth} = \left( {\left( B_{r} + B_{w} \right) \div 10^{9}} \right) \div \text{time}\)
这里，有效带宽的单位是GB/s，Br表示每个内核读取的字节数，Bw表示每个内核写入的字节数，时间以秒为单位给出。
例如，要计算2048 x 2048矩阵复制的有效带宽，可以使用以下公式：
\(\text{Effective\ bandwidth} = \left( {\left( 2048^{2} \times 4 \times 2 \right) \div 10^{9}} \right) \div \text{time}\)
元素数量乘以每个元素的大小（浮点数为4字节），再乘以2（因为涉及读取和写入操作），然后除以109（或1,0243）得到传输的内存GB数。最后将该数值除以时间（秒）即可得到GB/s的传输速率。
8.2.3. Visual Profiler报告的吞吐量
对于计算能力为2.0或更高的设备，可以使用Visual Profiler收集多种不同的内存吞吐量测量值。以下吞吐量指标可以在Details或Detail Graphs视图中显示：
请求的全局加载吞吐量
请求的全局存储吞吐量
全局加载吞吐量
全局存储吞吐量
DRAM读取吞吐量
DRAM写入吞吐量
请求的全局加载吞吐量和请求的全局存储吞吐量值表示内核请求的全局内存吞吐量，因此对应于有效带宽计算下所示计算获得的有效带宽。
由于最小内存事务大小通常大于大多数字的大小，内核实际所需的内存吞吐量可能包含内核未使用的数据传输。对于全局内存访问，这一实际吞吐量由全局加载吞吐量(Global Load Throughput)和全局存储吞吐量(Global Store Throughput)值报告。
需要注意的是，这两个数值都具有参考价值。实际内存吞吐量显示了代码接近硬件极限的程度，而将有效（或请求）带宽与实际带宽进行对比，可以很好地估算出由于内存访问未充分合并而浪费了多少带宽（参见全局内存的合并访问）。对于全局内存访问，这种请求内存带宽与实际内存带宽的对比会通过"全局内存加载效率"和"全局内存存储效率"指标来反映。
- 2
-
作为例外情况，对HBM2的分散写入会受到ECC的一些开销影响，但远低于受ECC保护的GDDR5内存上类似访问模式的开销。
9. 内存优化
内存优化是性能提升最重要的领域。其目标是通过最大化带宽来充分利用硬件性能。要获得最佳带宽，应尽可能使用更多高速内存，减少低速访问内存的使用。本章将讨论主机和设备上的各类内存，以及如何最佳配置数据项以高效利用内存。
9.1. 主机与设备之间的数据传输
设备内存与GPU之间的峰值理论带宽（例如NVIDIA Tesla V100可达898 GB/s）远高于主机内存与设备内存之间的峰值理论带宽（PCIe x16 Gen3为16 GB/s）。因此，为了获得最佳整体应用性能，尽量减少主机与设备之间的数据传输至关重要，即使这意味着在GPU上运行的内核相比在主机CPU上运行没有任何加速效果。
注意
高优先级：尽量减少主机和设备之间的数据传输，即使这意味着在设备上运行某些内核相比在主机CPU上运行不会带来性能提升。
中间数据结构应在设备内存中创建，由设备进行操作，并在无需主机映射或复制到主机内存的情况下销毁。
此外，由于每次数据传输都会产生额外开销，将多个小批量传输合并为一次大批量传输的性能会显著优于单独执行每次传输，即使这样做需要将非连续的内存区域打包到连续缓冲区中并在传输后解包。
最后，如CUDA C++编程指南和本文档固定内存部分所述，使用页锁定（或称固定）内存可实现主机与设备间更高的带宽。
9.1.1. 固定内存
页锁定或固定内存传输可实现主机与设备之间的最高带宽。例如，在PCIe x16 Gen3卡上，固定内存可达到约12 GB/s的传输速率。
固定内存是通过运行时API中的cudaHostAlloc()函数分配的。CUDA示例bandwidthTest展示了如何使用这些函数以及如何测量内存传输性能。
对于已经预分配的系统内存区域，可以使用cudaHostRegister()动态固定内存，无需分配单独的缓冲区并将数据复制到其中。
固定内存不应过度使用。过度使用会降低整体系统性能，因为固定内存是一种稀缺资源，但多少算过多很难提前预知。此外，与大多数常规系统内存分配相比，系统内存的固定操作是一项重量级操作，因此与所有优化一样，应测试应用程序及其运行系统以获得最佳性能参数。
9.1.2. 计算与异步及重叠传输
主机与设备之间使用cudaMemcpy()进行的数据传输是阻塞式传输，也就是说，只有在数据传输完成后控制权才会返回给主机线程。而cudaMemcpyAsync()函数是cudaMemcpy()的非阻塞版本，它会立即将控制权返回给主机线程。与cudaMemcpy()不同，异步传输版本要求使用固定主机内存（参见Pinned Memory），并且包含一个额外的参数——流ID。流本质上是在设备上按顺序执行的一系列操作。不同流中的操作可以交错执行，在某些情况下甚至可以重叠——这一特性可用于隐藏主机与设备之间的数据传输。
异步传输通过两种不同的方式实现数据传输与计算的重叠。在所有支持CUDA的设备上，主机计算可以与异步数据传输及设备计算重叠执行。例如，异步传输与计算的重叠展示了在数据传输到设备并执行使用该设备的内核时，如何同时执行例程cpuFunction()中的主机计算。
重叠计算与数据传输
cudaMemcpyAsync(a_d, a_h, size, cudaMemcpyHostToDevice, 0);
kernel<<<grid, block>>>(a_d);
cpuFunction();
cudaMemcpyAsync()函数的最后一个参数是流ID，在本例中使用的是默认流（流0）。内核同样使用默认流，它将在内存拷贝完成后才开始执行；因此不需要显式同步。由于内存拷贝和内核都会立即将控制权返回给主机，主机函数cpuFunction()会与它们的执行重叠。
在异步计算与重叠传输中，内存拷贝与内核执行是按顺序进行的。在支持并发拷贝与计算的设备上，可以实现设备内核执行与主机-设备间数据传输的重叠操作。设备是否具备此能力可通过cudaDeviceProp结构体的asyncEngineCount字段判断（或查看CUDA示例程序deviceQuery的输出结果）。对于支持此功能的设备，要实现重叠操作仍需使用固定主机内存，此外数据传输与内核必须使用不同的非默认流（流ID非零的流）。之所以要求非默认流，是因为使用默认流的内存拷贝、内存设置函数及内核调用，必须等待设备上所有先前调用（任何流中的）完成后才会开始，且这些操作完成前设备上（任何流中的）其他操作都无法启动。
异步和重叠传输与计算 展示了基本技术。
并发复制与执行
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
cudaMemcpyAsync(a_d, a_h, size, cudaMemcpyHostToDevice, stream1);
kernel<<<grid, block, 0, stream2>>>(otherData_d);
在这段代码中，创建了两个流，并按照cudaMemcpyAsync调用和内核执行配置的最后一个参数所指定的那样，用于数据传输和内核执行。
异步计算与重叠传输展示了如何将内核执行与异步数据传输重叠进行。当数据依赖关系允许将数据分块并以多阶段传输时，可以启动多个内核在每个数据块到达时进行处理，此时可采用此技术。顺序复制执行和分阶段并发复制执行对此进行了演示。两者产生的结果相同。第一段展示了参考顺序实现，该实现传输并操作一个包含N个浮点数的数组（假设N能被nThreads整除）。
顺序复制并执行
cudaMemcpy(a_d, a_h, N*sizeof(float), dir);
kernel<<<N/nThreads, nThreads>>>(a_d);
分阶段并发拷贝与执行展示了如何将数据传输和内核执行分解为nStreams阶段。这种方法允许数据传输和执行部分重叠。
分阶段并发复制与执行
size=N*sizeof(float)/nStreams;
for (i=0; i<nStreams; i++) {
offset = i*N/nStreams;
cudaMemcpyAsync(a_d+offset, a_h+offset, size, dir, stream[i]);
kernel<<<N/(nThreads*nStreams), nThreads, 0,
stream[i]>>>(a_d+offset);
}
(在分阶段并发拷贝与执行中，假设N能被nThreads*nStreams整除。)由于流内的执行是顺序进行的，在各自流中的数据拷贝完成之前，所有内核都不会启动。当前GPU可以同时处理异步数据传输和执行内核运算。具有单个拷贝引擎的GPU可以执行一项异步数据传输并运行内核，而具有两个拷贝引擎的GPU可以同时执行从主机到设备的一项异步数据传输、从设备到主机的一项异步数据传输以及内核运算。GPU上的拷贝引擎数量由cudaDeviceProp结构体的asyncEngineCount字段给出，该字段也会列在CUDA示例程序deviceQuery的输出中。（需要说明的是，阻塞传输无法与异步传输重叠执行，因为阻塞传输发生在默认流中，所以它要等到之前所有CUDA调用完成后才会开始，并且在它完成之前不会允许任何其他CUDA调用开始。）描述这两段代码执行时间线的示意图见图1，其中下半部分分阶段并发拷贝与执行的nStreams值为4。
图1 拷贝与内核执行的时间线对比
- Top
-
顺序
- Bottom
-
并发
在这个示例中，假设数据传输时间和内核执行时间相当。在这种情况下，当执行时间(tE)超过传输时间(tT)时，分阶段版本的总体时间粗略估计为tE + tT/nStreams，而顺序版本则为tE + tT。如果传输时间超过执行时间，则总体时间的粗略估计为tT + tE/nStreams。
9.1.3. 零拷贝
零拷贝是CUDA Toolkit 2.2版本中引入的一项功能。它允许GPU线程直接访问主机内存。为此，需要使用映射的固定（不可分页）内存。在集成GPU（即CUDA设备属性结构中integrated字段设置为1的GPU）上，映射固定内存总能带来性能提升，因为集成GPU和CPU内存物理上是同一块内存，从而避免了多余的拷贝。在独立GPU上，映射固定内存仅在特定情况下才有优势。由于数据不会缓存在GPU上，映射固定内存应仅读取或写入一次，并且读写该内存的全局加载和存储操作应保持合并访问。零拷贝可以替代流的使用，因为内核发起的数据传输会自动与内核执行重叠，无需额外设置和确定最优流数量的开销。
注意
低优先级：在集成GPU上对CUDA Toolkit 2.2及更高版本使用零拷贝操作。
主机代码在零拷贝主机代码中展示了如何典型地设置零拷贝。
零拷贝主机代码
float *a_h, *a_map;
...
cudaGetDeviceProperties(&prop, 0);
if (!prop.canMapHostMemory)
exit(0);
cudaSetDeviceFlags(cudaDeviceMapHost);
cudaHostAlloc(&a_h, nBytes, cudaHostAllocMapped);
cudaHostGetDevicePointer(&a_map, a_h, 0);
kernel<<<gridSize, blockSize>>>(a_map);
在这段代码中，通过cudaGetDeviceProperties()返回的结构体中的canMapHostMemory字段用于检查设备是否支持将主机内存映射到设备地址空间。通过调用cudaSetDeviceFlags()并传入cudaDeviceMapHost参数来启用页锁定内存映射功能。需要注意的是，cudaSetDeviceFlags()必须在设置设备或进行需要状态的CUDA调用之前调用（本质上就是在创建上下文之前）。页锁定的映射主机内存通过cudaHostAlloc()分配，而映射到设备地址空间的指针则通过cudaHostGetDevicePointer()函数获取。在Zero-copy host code的代码中，kernel()可以使用指针a_map引用映射的固定主机内存，其使用方式与a_map指向设备内存中的位置时完全相同。
注意
映射的固定主机内存允许您在避免使用CUDA流的同时，重叠CPU-GPU内存传输与计算。但由于对此类内存区域的任何重复访问都会导致重复的CPU-GPU传输，建议在设备内存中创建第二个区域，用于手动缓存先前读取的主机内存数据。
9.1.4. 统一虚拟寻址
支持计算能力2.0及更高版本的设备在64位Linux和Windows系统上支持一种称为统一虚拟寻址(UVA)的特殊寻址模式。通过UVA，主机内存和所有已安装支持设备的内存共享一个虚拟地址空间。
在使用UVA之前，应用程序必须通过额外的元数据位（或在程序中硬编码信息）来跟踪每个指针是指向设备内存（以及对应哪个设备）还是主机内存。而采用UVA后，只需通过cudaPointerGetAttributes()检查指针值，就能确定该指针指向的物理内存空间。
在UVA架构下，通过cudaHostAlloc()分配的固定主机内存将拥有相同的主机和设备指针，因此无需对此类分配调用cudaHostGetDevicePointer()。然而，通过cudaHostRegister()事后固定的主机内存分配仍将保持与主机指针不同的设备指针，因此在这种情况下仍需使用cudaHostGetDevicePointer()。
UVA也是实现点对点(P2P)数据传输的必要前提条件，在支持的配置中，对于支持的GPU可以直接通过PCIe总线或NVLink传输数据，绕过主机内存。
有关UVA和P2P的进一步解释和软件要求，请参阅CUDA C++编程指南。
9.2. 设备内存空间
CUDA设备使用多种内存空间，这些空间具有不同的特性，反映了它们在CUDA应用中的不同用途。这些内存空间包括全局内存、本地内存、共享内存、纹理内存和寄存器，如图2所示。
图2 CUDA设备上的内存空间
在这些不同的内存空间中，全局内存容量最大；具体各计算能力级别下每个内存空间的可用容量，请参阅《CUDA C++编程指南》的特性与技术规范章节。全局内存、本地内存和纹理内存的访问延迟最高，其次是常量内存、共享内存以及寄存器文件。
各种内存类型的主要特性如表1所示。
内存 |
片上/片外位置 |
是否缓存 |
访问方式 |
作用域 |
生命周期 |
|---|---|---|---|---|---|
寄存器 |
开启 |
不适用 |
读写 |
1线程 |
线程 |
本地 |
关闭 |
是†† |
读写 |
1线程 |
线程 |
共享 |
开启 |
不适用 |
读写 |
块内所有线程 |
块 |
全局 |
关闭 |
† |
读写 |
所有线程+主机 |
主机分配 |
常量 |
关闭 |
是 |
R |
所有线程+主机 |
主机分配 |
纹理 |
关闭 |
是 |
R |
所有线程+主机 |
主机分配 |
† 在计算能力6.0和7.x的设备上默认缓存在L1和L2中；在较低计算能力的设备上默认仅缓存在L2中，不过某些设备也允许通过编译标志选择性地缓存到L1中。 |
|||||
†† 默认情况下，除计算能力5.x的设备外，本地变量会缓存在L1和L2中；计算能力5.x的设备仅将本地变量缓存在L2中。 |
在纹理访问的情况下，如果纹理引用绑定到全局内存中的线性数组，则设备代码可以写入底层数组。绑定到CUDA数组的纹理引用可以通过表面写入操作进行写入（方法是将表面绑定到相同的底层CUDA数组存储）。应避免在同一内核启动期间从纹理读取数据的同时向其底层全局内存数组写入数据，因为纹理缓存是只读的，并且在修改关联的全局内存时不会失效。
9.2.1. 全局内存的合并访问
在为支持CUDA的GPU架构编程时，一个非常重要的性能考量因素是全局内存访问的合并。设备会将一个线程束中所有线程的全局内存加载和存储操作合并为尽可能少的事务。
注意
高优先级：尽可能确保全局内存访问是合并的。
合并访问的要求取决于设备的计算能力，具体文档可参考CUDA C++编程指南。
对于计算能力6.0或更高的设备，其要求可以非常简单地概括：一个warp中线程的并发访问将合并成若干次事务，这些事务的数量等于服务该warp所有线程所需的32字节事务的数量。
对于计算能力5.2的特定设备，可以选择启用对全局内存访问的L1缓存。如果在这些设备上启用了L1缓存，所需的事务数量等于所需的128字节对齐段的数量。
注意
在计算能力6.0或更高的设备上，L1缓存是默认启用的，但无论全局加载是否缓存在L1中，数据访问单元都是32字节。
在使用GDDR内存的设备上，当ECC功能开启时，以合并方式访问内存变得更为重要。分散的访问会增加ECC内存传输开销，尤其是在向全局内存写入数据时。
以下简单示例说明了合并概念。除非另有说明，这些示例假设计算能力为6.0或更高，并且访问的是4字节字。
9.2.1.1. 简单访问模式
第一种也是最简单的合并情况可以由任何计算能力6.0或更高的CUDA设备实现：第k个线程访问32字节对齐数组中的第k个字。并非所有线程都需要参与。
例如，如果一个线程束(warp)的线程访问相邻的4字节字(如相邻的float值)，四个合并的32字节事务将处理该内存访问。这种模式如图3
图3 合并访问
这种访问模式导致了四次32字节的传输，如红色矩形所示。
如果四个32字节段中的任何一个仅请求了部分字（例如，如果多个线程访问了同一个字，或者某些线程未参与访问），系统仍会获取整个段。此外，即使线程束中的访问在四个段内或跨段进行了置换，对于计算能力6.0或更高的设备，仍然只会执行四次32字节的事务。
9.2.1.2. 顺序但未对齐的访问模式
如果warp中的连续线程访问的内存是连续的但未与32字节段对齐，则会请求五个32字节段，如图4所示。
图4 分布在五个32字节段内的未对齐连续地址
通过CUDA Runtime API（例如cudaMalloc()）分配的内存保证至少对齐256字节。因此，选择合理的线程块大小（例如当前GPU中warp大小的倍数，即32的倍数）有助于warp进行正确对齐的内存访问。（例如，考虑如果线程块大小不是warp大小的倍数时，第二个、第三个及后续线程块访问的内存地址会发生什么情况。）
9.2.1.3. 未对齐访问的影响
通过一个简单的复制内核（如展示未对齐访问的复制内核中的示例）来探究未对齐访问的影响，既简单又富有启发性。
一个演示未对齐访问的复制内核
__global__ void offsetCopy(float *odata, float* idata, int offset)
{
int xid = blockIdx.x * blockDim.x + threadIdx.x + offset;
odata[xid] = idata[xid];
}
在展示未对齐访问的复制内核中，数据从输入数组idata复制到输出数组，两者都位于全局内存中。该内核在主机代码的循环中执行，循环会将参数offset从0变化到32（例如，图4对应这些未对齐情况）。在NVIDIA Tesla V100（计算能力7.0）上使用不同偏移量进行复制的有效带宽如图5所示。
图5 offsetCopy内核性能
对于NVIDIA Tesla V100显卡，当全局内存访问没有偏移量或偏移量为8个字的倍数时，会产生四次32字节的事务传输。此时实现的带宽约为790 GB/s。其他情况下，每个线程束会加载五个32字节段，我们预计其内存吞吐量约为无偏移情况下的4/5th。
在这个特定示例中，由于相邻线程束复用了其邻居获取的缓存行，最终实现的偏移内存吞吐量约为9/10th。因此虽然影响仍然明显，但并没有我们预期的那么大。如果相邻线程束没有表现出如此高程度的过量获取缓存行复用，影响本应更为显著。
9.2.1.4. 跨步访问
如上所述，在顺序访问未对齐的情况下，缓存有助于减轻性能影响。但对于非单位步长的访问模式（这在处理多维数据或矩阵时经常出现），情况可能有所不同。因此，确保所获取的每个缓存行中的数据尽可能被实际使用，是优化这些设备上内存访问性能的重要环节。
为了演示跨步访问对有效带宽的影响，请查看用于演示非单位跨步数据复制的内核中的strideCopy()内核，该内核以跨步元素为间隔，将数据从idata复制到odata。
一个演示非单位步长数据复制的内核
__global__ void strideCopy(float *odata, float* idata, int stride)
{
int xid = (blockIdx.x*blockDim.x + threadIdx.x)*stride;
odata[xid] = idata[xid];
}
图6展示了这种情况; 在这个例子中，warp内的线程以步长2访问内存中的字。这个操作导致Tesla V100(计算能力7.0)上每个warp加载8个L2缓存段。
图6 以跨度为2访问内存的相邻线程
步长为2会导致50%的加载/存储效率，因为事务中有一半元素未被使用，造成了带宽浪费。随着步长增加，有效带宽持续下降，直到一个线程束中的32个线程需要加载32个32字节段为止，如图7所示。
图7 strideCopy内核性能
如图7所示，应尽可能避免非单位步长的全局内存访问。实现这一点的一种方法是利用共享内存，这将在下一节中讨论。
9.2.2. 二级缓存
从CUDA 11.0开始，计算能力8.0及以上的设备能够影响数据在L2缓存中的持久性。由于L2缓存位于芯片上，它可能为全局内存提供更高的带宽和更低的访问延迟。
更多详情请参阅CUDA C++编程指南中的L2访问管理章节。
9.2.2.1. L2缓存访问窗口
当CUDA内核重复访问全局内存中的某个数据区域时，这类数据访问可被视为持久性访问。反之，若数据仅被访问一次，则这类访问可视为流式访问。可以将L2缓存的一部分预留出来，专门用于全局内存中数据区域的持久性访问。如果预留部分未被持久性访问占用，那么流式或常规数据访问也可以使用这部分缓存。
持久访问的L2缓存预留大小可以在限制范围内进行调整：
cudaGetDeviceProperties(&prop, device_id);
cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, prop.persistingL2CacheMaxSize); /* Set aside max possible size of L2 cache for persisting accesses */
用户数据到L2预留部分的映射可以通过在CUDA流或CUDA图内核节点上使用访问策略窗口来控制。下面的示例展示了如何在CUDA流上使用访问策略窗口。
cudaStreamAttrValue stream_attribute; // Stream level attributes data structure
stream_attribute.accessPolicyWindow.base_ptr = reinterpret_cast<void*>(ptr); // Global Memory data pointer
stream_attribute.accessPolicyWindow.num_bytes = num_bytes; // Number of bytes for persisting accesses.
// (Must be less than cudaDeviceProp::accessPolicyMaxWindowSize)
stream_attribute.accessPolicyWindow.hitRatio = 1.0; // Hint for L2 cache hit ratio for persisting accesses in the num_bytes region
stream_attribute.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting; // Type of access property on cache hit
stream_attribute.accessPolicyWindow.missProp = cudaAccessPropertyStreaming; // Type of access property on cache miss.
//Set the attributes to a CUDA stream of type cudaStream_t
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &stream_attribute);
访问策略窗口需要为hitRatio和num_bytes提供值。根据num_bytes参数的值和L2缓存的大小，可能需要调整hitRatio的值以避免L2缓存行的抖动。
9.2.2.2. 调优访问窗口命中率
hitRatio参数可用于指定接收hitProp属性的访问比例。例如，若hitRatio值为0.6，则表示全局内存区域[ptr..ptr+num_bytes)中60%的内存访问具有持久化属性，40%具有流式属性。为理解hitRatio和num_bytes的影响，我们使用了滑动窗口微基准测试。
该微基准测试使用GPU全局内存中的1024 MB区域。首先，如上文所述，我们通过cudaDeviceSetLimit()预留30 MB的L2缓存用于持久化访问。如下图所示，我们指定对内存区域前freqSize * sizeof(int)字节的访问为持久化访问。这部分数据将因此使用预留的L2缓存部分。在实验中，我们将持久化数据区域的大小从10 MB调整到60 MB，以模拟数据适应或超出30 MB可用L2预留缓存的各种场景。请注意，NVIDIA Tesla A100 GPU的总L2缓存容量为40 MB。对内存区域剩余数据（即流式数据）的访问被视为常规或流式访问，因此将使用剩余的10 MB非预留L2缓存部分（除非预留的L2缓存部分有未使用空间）。
图8 滑动窗口实验中持久化数据访问映射到预留L2缓存的示意图
考虑以下内核代码和访问窗口参数，作为滑动窗口实验的实现。
__global__ void kernel(int *data_persistent, int *data_streaming, int dataSize, int freqSize) {
int tid = blockIdx.x * blockDim.x + threadIdx.x;
/*Each CUDA thread accesses one element in the persistent data section
and one element in the streaming data section.
Because the size of the persistent memory region (freqSize * sizeof(int) bytes) is much
smaller than the size of the streaming memory region (dataSize * sizeof(int) bytes), data
in the persistent region is accessed more frequently*/
data_persistent[tid % freqSize] = 2 * data_persistent[tid % freqSize];
data_streaming[tid % dataSize] = 2 * data_streaming[tid % dataSize];
}
stream_attribute.accessPolicyWindow.base_ptr = reinterpret_cast<void*>(data_persistent);
stream_attribute.accessPolicyWindow.num_bytes = freqSize * sizeof(int); //Number of bytes for persisting accesses in range 10-60 MB
stream_attribute.accessPolicyWindow.hitRatio = 1.0; //Hint for cache hit ratio. Fixed value 1.0
上述内核的性能如下图所示。当持久数据区域完全适配L2缓存预留的30MB空间时，性能提升可达50%。然而，一旦该持久数据区域的大小超过L2缓存预留部分容量，由于L2缓存行的频繁换出，性能会下降约10%。
图9 固定命中率为1.0时滑动窗口基准测试的性能表现
为了优化性能，当持久化数据的大小超过预留L2缓存部分时，我们按如下方式调整访问窗口中的num_bytes和hitRatio参数。
stream_attribute.accessPolicyWindow.base_ptr = reinterpret_cast<void*>(data_persistent);
stream_attribute.accessPolicyWindow.num_bytes = 20*1024*1024; //20 MB
stream_attribute.accessPolicyWindow.hitRatio = (20*1024*1024)/((float)freqSize*sizeof(int)); //Such that up to 20MB of data is resident.
我们将访问窗口中的num_bytes固定为20MB，并调整hitRatio，使得总持久化数据中随机20MB驻留在L2预留缓存部分。这部分持久化数据的剩余部分将通过流式属性进行访问。这有助于减少缓存抖动。结果如下图所示，我们可以看到无论持久化数据是否适合L2预留缓存，性能表现都很好。
图10 经过命中率调优后的滑动窗口基准测试性能
9.2.3. 共享内存
由于位于芯片上，共享内存相比本地和全局内存具有更高的带宽和更低的延迟——前提是线程之间不存在存储体冲突，具体细节将在下一节中详述。
9.2.3.2. 矩阵乘法中的共享内存(C=AB)
共享内存支持块内线程之间的协作。当块中的多个线程需要使用全局内存中的相同数据时，共享内存可以确保只需从全局内存中访问一次数据。共享内存还能通过以合并模式从全局内存加载/存储数据，然后在共享内存中重新排序，从而避免非合并的内存访问。除了存储体冲突外，warp在共享内存中进行非连续或未对齐访问不会产生性能损失。
共享内存的使用通过一个简单的矩阵乘法示例C = AB来说明，其中矩阵A的维度为Mxw，矩阵B的维度为wxN，矩阵C的维度为MxN。为了使内核保持简单，M和N都是32的倍数，因为当前设备的线程束大小(w)为32。
该问题的一个自然分解方式是使用wxw线程的块和瓦片大小。因此，在wxw瓦片的概念下，A是一个列矩阵，B是一个行矩阵，而C是它们的外积；参见图11。启动一个N/w乘M/w块的网格，其中每个线程块根据A的单个瓦片和B的单个瓦片计算C中不同瓦片的元素。
图11 块列矩阵乘以块行矩阵。块列矩阵(A)乘以块行矩阵(B)得到结果乘积矩阵(C)。
为此，simpleMultiply内核（未优化的矩阵乘法）计算矩阵C的一个分块的输出元素。
未优化的矩阵乘法
__global__ void simpleMultiply(float *a, float* b, float *c,
int N)
{
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
float sum = 0.0f;
for (int i = 0; i < TILE_DIM; i++) {
sum += a[row*TILE_DIM+i] * b[i*N+col];
}
c[row*N+col] = sum;
}
在未优化的矩阵乘法中，a、b和c分别是指向矩阵A、B和C的全局内存指针；blockDim.x、blockDim.y和TILE_DIM都等于w。wxw线程块中的每个线程计算C矩阵图块中的一个元素。row和col表示特定线程正在计算的C矩阵元素的行列位置。for循环遍历i将A的一行与B的一列相乘，结果写入C。
该内核在NVIDIA Tesla V100上的有效带宽为119.9 GB/s。为了分析性能，需要考虑在for循环中线程束如何访问全局内存。每个线程束计算C的一个分块行，这依赖于A的单个行和B的整个分块，如图12所示。
图12 计算图块的一行。在C中使用A的一行和B的整个图块来计算图块的一行。
在for循环的每次迭代i中，一个warp内的线程会读取B tile的一行数据，这对所有计算能力来说都是顺序且合并的访问方式。
然而，对于每次迭代i，一个warp中的所有线程都会从全局内存中读取矩阵A的相同值，因为索引row*TILE_DIM+i在warp内是恒定的。尽管在计算能力2.0或更高的设备上，这种访问只需要1次事务，但该事务中存在带宽浪费，因为在32字节缓存段中，8个字里只有1个4字节字被使用。我们可以在循环的后续迭代中重用这个缓存行，最终会利用全部8个字；然而，当多个warp同时在同一个多处理器上执行时（通常情况就是如此），缓存行很容易在迭代i和i+1之间从缓存中被逐出。
通过将A矩阵的一个分块读取到共享内存中，可以提升任何计算能力设备上的性能，具体方法如使用共享内存提升矩阵乘法中全局内存加载效率所示。
__global__ void coalescedMultiply(float *a, float* b, float *c,
int N)
{
__shared__ float aTile[TILE_DIM][TILE_DIM];
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
float sum = 0.0f;
aTile[threadIdx.y][threadIdx.x] = a[row*TILE_DIM+threadIdx.x];
__syncwarp();
for (int i = 0; i < TILE_DIM; i++) {
sum += aTile[threadIdx.y][i]* b[i*N+col];
}
c[row*N+col] = sum;
}
在使用共享内存提高矩阵乘法中全局内存加载效率一文中，A矩阵分块中的每个元素仅以完全合并的方式（无带宽浪费）从全局内存读取一次到共享内存。在每次for循环迭代中，共享内存中的值会广播给整个线程束中的所有线程。由于只有将数据写入共享内存的线程束内线程才会读取这些数据，因此在将A矩阵分块读入共享内存后，使用__syncwarp()同步屏障调用就足够了，而无需调用__syncthreads()。该内核在NVIDIA Tesla V100上实现了144.4 GB/s的有效带宽。这说明了当硬件L1缓存驱逐策略与应用程序需求不匹配，或者L1缓存未用于全局内存读取时，将共享内存作为用户管理缓存的使用方法。
可以进一步改进使用共享内存提升矩阵乘法中全局内存加载效率处理矩阵B的方式。在计算矩阵C每个分块的行时，会重复读取B分块的全部数据。通过将B分块一次性读入共享内存(通过将额外数据读入共享内存进行优化)，可以消除这种重复读取。
__global__ void sharedABMultiply(float *a, float* b, float *c,
int N)
{
__shared__ float aTile[TILE_DIM][TILE_DIM],
bTile[TILE_DIM][TILE_DIM];
int row = blockIdx.y * blockDim.y + threadIdx.y;
int col = blockIdx.x * blockDim.x + threadIdx.x;
float sum = 0.0f;
aTile[threadIdx.y][threadIdx.x] = a[row*TILE_DIM+threadIdx.x];
bTile[threadIdx.y][threadIdx.x] = b[threadIdx.y*N+col];
__syncthreads();
for (int i = 0; i < TILE_DIM; i++) {
sum += aTile[threadIdx.y][i]* bTile[i][threadIdx.x];
}
c[row*N+col] = sum;
}
请注意，在通过将额外数据读入共享内存进行优化时，读取B数据块后需要调用__syncthreads()，因为一个线程束会读取由其他线程束写入共享内存的数据。该例程在NVIDIA Tesla V100上的有效带宽为195.5 GB/s。需要注意的是，性能提升并非来自任何一种情况下的合并访问改进，而是避免了从全局内存的冗余数据传输。
各种优化结果总结在表2中。
优化 |
NVIDIA Tesla V100 |
|---|---|
未优化 |
119.9 GB/s |
使用共享内存合并存储A的图块 |
144.4 GB/s |
使用共享内存消除对B图块的多余读取 |
195.5 GB/s |
注意
中等优先级： 使用共享内存以避免从全局内存的冗余传输。
9.2.3.4. 从全局内存到共享内存的异步拷贝
CUDA 11.0 引入了一项异步拷贝功能，可在设备代码中用于显式管理从全局内存到共享内存的异步数据拷贝。该功能使CUDA内核能够将全局内存到共享内存的数据拷贝与计算操作重叠执行，同时避免了传统架构中全局内存读取和共享内存写入之间必须经过寄存器文件访问的中间步骤。
更多详情请参阅CUDA C++编程指南中的memcpy_async章节。
为了理解从全局内存到共享内存的同步拷贝与异步拷贝之间的性能差异，请参考以下用于演示同步和异步方法的CUDA微基准测试内核。异步拷贝在NVIDIA A100 GPU上可获得硬件加速。
template <typename T>
__global__ void pipeline_kernel_sync(T *global, uint64_t *clock, size_t copy_count) {
extern __shared__ char s[];
T *shared = reinterpret_cast<T *>(s);
uint64_t clock_start = clock64();
for (size_t i = 0; i < copy_count; ++i) {
shared[blockDim.x * i + threadIdx.x] = global[blockDim.x * i + threadIdx.x];
}
uint64_t clock_end = clock64();
atomicAdd(reinterpret_cast<unsigned long long *>(clock),
clock_end - clock_start);
}
template <typename T>
__global__ void pipeline_kernel_async(T *global, uint64_t *clock, size_t copy_count) {
extern __shared__ char s[];
T *shared = reinterpret_cast<T *>(s);
uint64_t clock_start = clock64();
//pipeline pipe;
for (size_t i = 0; i < copy_count; ++i) {
__pipeline_memcpy_async(&shared[blockDim.x * i + threadIdx.x],
&global[blockDim.x * i + threadIdx.x], sizeof(T));
}
__pipeline_commit();
__pipeline_wait_prior(0);
uint64_t clock_end = clock64();
atomicAdd(reinterpret_cast<unsigned long long *>(clock),
clock_end - clock_start);
}
内核的同步版本会先将元素从全局内存加载到中间寄存器，然后再将中间寄存器的值存储到共享内存。而在异步版本的内核中，一旦调用__pipeline_memcpy_async()函数，就会立即发出从全局内存加载并直接存储到共享内存的指令。__pipeline_wait_prior(0)会等待管道对象中的所有指令执行完毕。使用异步拷贝不会占用任何中间寄存器，这有助于降低寄存器压力并提高内核占用率。通过异步拷贝指令从全局内存复制到共享内存的数据可以被缓存在L1缓存中，也可以选择绕过L1缓存。如果单个CUDA线程正在拷贝16字节的元素，则可以绕过L1缓存。这一差异在图13中进行了说明。
图13 全局内存与共享内存同步/异步拷贝性能对比
我们评估了两个内核在使用每个线程4B、8B和16B大小元素（即使用int、int2和int4作为模板参数）时的性能表现。我们在内核中调整了copy_count参数，使每个线程块复制的数据量从512字节到48MB不等。内核性能表现如图14所示。
图14 全局内存与共享内存间同步复制与异步复制的性能对比
从性能图表可以看出，本次实验有以下观察结果。
当
copy_count参数对于所有三种元素大小都是4的倍数时，同步复制可获得最佳性能。编译器可以优化4组加载和存储指令。从锯齿曲线中可以明显看出这一点。在几乎所有情况下，异步复制都能实现更好的性能。
异步复制不需要
copy_count参数是4的倍数，通过编译器优化实现最佳性能。总体而言，当使用8或16字节大小的元素进行异步复制时，可获得最佳性能。
9.2.4. 本地内存
本地内存之所以称为"本地"，是因为其作用域仅限于线程内部，而非其物理位置。实际上，本地内存位于芯片外。因此，访问本地内存与访问全局内存的开销同样高昂。换言之，名称中的"本地"一词并不意味着更快的访问速度。
Local memory is used only to hold automatic variables. This is done by the nvcc compiler when it determines that there is insufficient register space to hold the variable. Automatic variables that are likely to be placed in local memory are large structures or arrays that would consume too much register space and arrays that the compiler determines may be indexed dynamically.
检查PTX汇编代码（通过使用-ptx或-keep命令行选项编译nvcc获得）可以揭示变量是否在初始编译阶段被放置在本地内存中。如果是，它将使用.local助记符声明，并通过ld.local和st.local助记符访问。如果没有，后续编译阶段仍可能做出不同决定，如果发现该变量为目标架构占用了过多寄存器空间。无法针对特定变量检查这一点，但编译器在使用--ptxas-options=-v选项运行时，会报告每个内核的本地内存总使用量（lmem）。
9.2.5. 纹理内存
只读纹理内存空间是缓存的。因此，纹理获取仅在缓存未命中时消耗一次设备内存读取；否则，它只需从纹理缓存中读取一次。纹理缓存针对二维空间局部性进行了优化，因此同一线程束中读取地址相近的线程将获得最佳性能。纹理内存还设计用于具有恒定延迟的流式获取；也就是说，缓存命中会降低DRAM带宽需求，但不会减少获取延迟。
在某些寻址场景中，通过纹理获取读取设备内存可能比从全局内存或常量内存读取更具优势。
9.2.5.1. 额外纹理功能
如果使用tex1D()、tex2D()或tex3D()而非tex1Dfetch()来获取纹理，硬件会提供其他可能对某些应用（如图像处理）有用的功能，如表4所示。
功能 |
用途 |
注意事项 |
|---|---|---|
过滤 |
纹理单元之间的快速低精度插值 |
仅在纹理引用返回浮点数据时有效 |
归一化的纹理坐标 |
分辨率无关编码 |
无 |
寻址模式 |
自动处理边界情况1 |
仅能与归一化纹理坐标一起使用 |
1 表4底部行中边界情况的自动处理指的是当纹理坐标超出有效寻址范围时如何解析。有两种选项：clamp（钳制）和wrap（环绕）。如果x是坐标，N是一维纹理的纹素数，那么在clamp模式下，当x < 0时会被替换为0，当1 <x时会被替换为1-1/N。在wrap模式下，x会被替换为frac(x)，其中frac(x) = x - floor(x)。floor函数返回小于等于x的最大整数。因此，在N = 1的clamp模式下，1.3会被钳制到1.0；而在wrap模式下，它会被转换为0.3 |
在内核调用期间，纹理缓存不会与全局内存写入保持一致性，因此从同一内核调用中通过全局存储写入的地址进行纹理提取将返回未定义数据。也就是说，如果某个内存位置是通过先前内核调用或内存复制更新的，线程可以通过纹理安全地读取该位置；但如果该位置是通过同一内核调用中的同一线程或其他线程先前更新的，则无法安全读取。
9.2.6. 常量内存
设备上共有64 KB的常量内存。常量内存空间具有缓存机制。因此，从常量内存读取数据时，只有在缓存未命中的情况下才会产生一次设备内存读取开销；否则仅需从常量缓存读取一次。同一个线程束内线程对不同地址的访问是串行化的，因此开销会随着线程束内所有线程读取的唯一地址数量线性增长。由此可见，当同一线程束内的线程仅访问少量不同位置时，常量缓存性能最佳。如果线程束内所有线程访问同一位置，常量内存的访问速度可达到寄存器访问级别。
9.2.7. 寄存器
通常情况下，访问寄存器每条指令不会消耗额外的时钟周期，但由于寄存器读写依赖和寄存器存储体冲突可能会导致延迟。
编译器和硬件线程调度器会尽可能优化地调度指令，以避免寄存器内存库冲突。应用程序无法直接控制这些库冲突。特别需要注意的是，没有与寄存器相关的原因需要将数据打包成向量数据类型，例如float4或int4类型。
9.2.7.1. 寄存器压力
寄存器压力发生在没有足够寄存器可用于给定任务时。尽管每个多处理器包含数千个32位寄存器（参见CUDA C++编程指南的功能和技术规格），但这些寄存器是在并发线程之间分配的。为防止编译器分配过多寄存器，可使用-maxrregcount=N编译器命令行选项或启动边界内核定义限定符（参见CUDA C++编程指南的执行配置）来控制每个线程分配的最大寄存器数量。
9.3. 分配
通过cudaMalloc()和cudaFree()进行设备内存分配和释放是昂贵的操作。建议使用流序池分配器cudaMallocAsync()和cudaFreeAsync()来管理设备内存。
9.4. NUMA最佳实践
部分最新的Linux发行版默认启用了自动NUMA平衡（或称"AutoNUMA"）。在某些情况下，自动NUMA平衡执行的操作可能会降低运行在NVIDIA GPU上的应用程序性能。为获得最佳性能，用户应手动调整其应用程序的NUMA特性。
最佳的NUMA调优取决于每个应用程序和节点的特性以及所需的硬件亲和性，但通常建议在NVIDIA GPU上进行计算的应用程序选择禁用自动NUMA平衡的策略。例如，在IBM Newell POWER9节点上（其中CPU对应NUMA节点0和8），使用：
numactl --membind=0,8
将内存分配绑定到CPU。
10. 执行配置优化
实现良好性能的关键之一，是尽可能保持设备上的多处理器处于繁忙状态。如果工作在多处理器之间分配不均，设备将无法发挥最佳性能。因此，设计应用程序时，必须合理规划线程和块的使用方式，以最大化硬件利用率，并避免阻碍工作自由分配的做法。其中关键概念是占用率（occupancy），这将在后续章节详细说明。
在某些情况下，通过设计应用程序使多个独立的内核能够同时执行，也可以提高硬件利用率。多个内核同时执行被称为并发内核执行。下文将描述并发内核执行。
另一个重要概念是管理为特定任务分配的系统资源。本章最后几节将讨论如何管理这些资源利用率。
10.1. 占用率
在CUDA中，线程指令是按顺序执行的，因此，当一个线程束暂停或停滞时，执行其他线程束是隐藏延迟并保持硬件忙碌的唯一方法。因此，与多处理器上活动线程束数量相关的某些指标对于确定硬件保持忙碌的有效性非常重要。这个指标就是占用率。
占用率是指每个多处理器上活跃warp数量与最大可能活跃warp数量的比值。（要确定后者数值，请参阅deviceQuery CUDA示例或参考计算能力。）另一种理解占用率的方式是：硬件处理warp的能力中被实际使用的百分比。
更高的占用率并不总是等同于更高的性能——存在一个临界点，超过该点后增加占用率将不再提升性能。然而，低占用率总会影响隐藏内存延迟的能力，从而导致性能下降。
CUDA内核所需的每线程资源可能会以不希望的方式限制最大块大小。为了保持对未来硬件和工具包的向前兼容性，并确保至少一个线程块可以在SM上运行，开发者应包含单参数__launch_bounds__(maxThreadsPerBlock)，该参数指定内核将启动的最大块大小。若不这样做，可能会导致"请求启动的资源过多"错误。在某些情况下，提供双参数版本的__launch_bounds__(maxThreadsPerBlock,minBlocksPerMultiprocessor)可以提高性能。minBlocksPerMultiprocessor的正确值应通过详细的内核分析来确定。
10.1.1. 计算占用率
决定占用率的几个因素之一是寄存器可用性。寄存器存储使线程能够将局部变量保存在附近以实现低延迟访问。然而，寄存器组（称为寄存器文件）是所有驻留在多处理器上的线程必须共享的有限资源。寄存器会一次性分配给整个线程块。因此，如果每个线程块使用大量寄存器，那么多处理器上可驻留的线程块数量就会减少，从而降低多处理器的占用率。每个线程的最大寄存器数量可以在编译时通过文件级选项-maxrregcount或内核级限定符__launch_bounds__手动设置（参见寄存器压力）。
在计算占用率时，每个线程使用的寄存器数量是关键因素之一。例如，在CUDA计算能力7.0的设备上，每个多处理器拥有65,536个32位寄存器，最多可同时驻留2048个线程（64个线程束×每束32线程）。这意味着在此类设备上，要实现多处理器100%的占用率，每个线程最多只能使用32个寄存器。不过，这种判断寄存器数量如何影响占用率的方法未考虑寄存器分配的粒度。例如，在计算能力7.0的设备上，使用每线程37个寄存器的128线程块内核可实现75%的占用率（每个多处理器运行12个128线程块），而使用相同37个寄存器的320线程块内核只能实现63%的占用率，因为每个多处理器仅能容纳4个320线程块。此外，寄存器分配会按每线程束向上取整至最接近的256个寄存器。
The number of registers available, the maximum number of simultaneous threads resident on each multiprocessor, and the register allocation granularity vary over different
compute capabilities. Because of these nuances in register allocation and the fact that a multiprocessor’s shared memory is also partitioned between resident thread blocks,
the exact relationship between register usage and occupancy can be difficult to determine. The --ptxas options=v option of nvcc details the number of registers
used per thread for each kernel. See Hardware Multithreading of the CUDA C++ Programming Guide for the register allocation formulas for devices of various compute
capabilities and Features and Technical Specifications of the CUDA C++ Programming Guide for the total number of registers available on those devices. Alternatively,
NVIDIA provides an occupancy calculator as part of Nsight Compute; refer to https://docs.nvidia.com/nsight-compute/NsightCompute/index.html#occupancy-calculator.
图15 使用CUDA占用率计算器预测GPU多处理器占用情况
应用程序还可以使用CUDA Runtime中的Occupancy API，例如cudaOccupancyMaxActiveBlocksPerMultiprocessor，根据运行时参数动态选择启动配置。
10.2. 隐藏寄存器依赖关系
注意
中等优先级：为了隐藏寄存器依赖导致的延迟，需在每个多处理器上维持足够数量的活跃线程（即保持足够的占用率）。
寄存器依赖发生在一条指令需要使用前一条指令写入寄存器中的结果时。在计算能力7.0的设备上，大多数算术指令的延迟通常为4个时钟周期。因此线程在使用算术结果前必须等待大约4个周期。不过，通过执行其他线程束中的线程可以完全隐藏这种延迟。详情请参阅Registers。
10.3. 线程与块启发式策略
注意
中等优先级：每个块的线程数应为32的倍数，因为这样可以提供最佳计算效率并促进合并操作。
每个网格的块维度与大小以及每个块的线程维度与大小都是重要因素。这些参数的多维特性使得多维问题更容易映射到CUDA，且对性能没有影响。因此，本节讨论大小而非维度。
延迟隐藏和占用率取决于每个多处理器上的活动线程束数量，这由执行参数以及资源（寄存器和共享内存）限制隐式决定。选择执行参数需要在延迟隐藏（占用率）和资源利用率之间取得平衡。
选择执行配置参数时应综合考虑，但每个参数都有其适用的启发式方法。在确定第一个执行配置参数——每个网格的块数（即网格大小）时，首要考虑因素是保持整个GPU处于忙碌状态。网格中的块数应多于多处理器数量，确保每个多处理器至少有一个块可执行。此外，每个多处理器应有多个活跃块，这样那些无需等待__syncthreads()的块就能持续占用硬件资源。该建议需结合资源可用性进行调整，因此应结合第二个执行参数——每个块的线程数（即块大小）以及共享内存使用情况来综合确定。为适配未来设备，每次内核启动的块数应达到数千量级。
在选择块大小时，需要记住多个并发块可以驻留在多处理器上，因此占用率不仅仅由块大小决定。特别要注意的是，较大的块大小并不意味着更高的占用率。
如Occupancy中所述，更高的占用率并不总是等同于更好的性能。例如，将占用率从66%提高到100%通常不会带来同比例的性能提升。较低占用率的内核每个线程可用的寄存器比高占用率内核更多，这可能减少寄存器溢出到本地内存的情况；特别是在具有高度显式指令级并行（ILP）的情况下，某些情况下甚至可以通过低占用率完全掩盖延迟。
选择块大小时涉及许多因素，不可避免地需要进行一些实验。不过，应遵循一些经验法则：
每个块的线程数应为warp大小的倍数，以避免在未充分利用的warps上浪费计算资源，并促进合并操作。
每个块至少应使用64个线程，且仅当每个多处理器有多个并发块时才适用。
每个块128到256个线程是尝试不同块大小的良好初始范围。
如果延迟影响性能，使用多个较小的线程块，而不是每个多处理器使用一个大的线程块。这对于频繁调用
__syncthreads()的内核特别有益。
请注意，当线程块分配的多处理器寄存器数量超过可用数量时，内核启动将失败，这种情况同样会发生在请求过多共享内存或线程数时。
10.5. 并发内核执行
如异步和重叠计算传输中所述，CUDA流可用于将内核执行与数据传输重叠。在支持并发内核执行的设备上，流还可用于同时执行多个内核，以更充分地利用设备的多处理器。设备是否具有此功能由cudaDeviceProp结构的concurrentKernels字段指示（或在deviceQuery CUDA示例的输出中列出）。并发执行需要使用非默认流（流0以外的流），因为使用默认流的内核调用仅在设备上所有先前调用（在任何流中）完成后才开始，并且在它们完成之前设备上的任何操作（在任何流中）都不会开始。
以下示例展示了基本技术。由于kernel1和kernel2在不同的非默认流中执行，性能足够的设备可以同时执行这些内核。
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
kernel1<<<grid, block, 0, stream1>>>(data_1);
kernel2<<<grid, block, 0, stream2>>>(data_2);
10.6. 多上下文环境
CUDA工作在一个特定GPU的进程空间内进行，这个空间被称为上下文。上下文封装了该GPU的内核启动和内存分配，以及页表等支持结构。在CUDA驱动API中上下文是显式存在的，而在CUDA运行时API中则完全隐式，后者会自动创建和管理上下文。
通过CUDA Driver API，一个CUDA应用程序进程可以为给定的GPU创建多个上下文。如果多个CUDA应用程序进程同时访问同一个GPU，这通常意味着存在多个上下文，因为上下文与特定的主机进程绑定，除非正在使用多进程服务。
虽然可以在给定的GPU上同时分配多个上下文（及其关联资源，如全局内存分配），但在任何时刻该GPU上只能有一个上下文执行任务；共享同一GPU的上下文会进行时间片轮转。创建额外的上下文会产生每个上下文数据的内存开销以及上下文切换的时间开销。此外，当多个上下文的工作本可以并发执行时，上下文切换的需求会降低利用率（另请参阅并发内核执行）。
因此，在同一个CUDA应用程序中最好避免每个GPU上存在多个上下文。为了帮助实现这一点，CUDA驱动API提供了访问和管理每个GPU上特殊上下文（称为主上下文）的方法。这些上下文与CUDA运行时在线程没有当前上下文时隐式使用的上下文相同。
// When initializing the program/library
CUcontext ctx;
cuDevicePrimaryCtxRetain(&ctx, dev);
// When the program/library launches work
cuCtxPushCurrent(ctx);
kernel<<<...>>>(...);
cuCtxPopCurrent(&ctx);
// When the program/library is finished with the context
cuDevicePrimaryCtxRelease(dev);
注意
NVIDIA-SMI可用于将GPU配置为独占进程模式，该模式将每个GPU的上下文数量限制为一个。此上下文可在创建进程内根据需要关联任意数量的线程，如果设备上已存在使用CUDA驱动API创建的非主上下文，则cuDevicePrimaryCtxRetain将执行失败。
11. 指令优化
了解指令如何执行通常有助于进行底层优化，这在频繁运行的代码（即程序中的所谓热点）中尤为有用。最佳实践建议，在所有高层优化完成后才进行此类优化。
11.1. 算术指令
单精度浮点数能提供最佳性能，强烈推荐使用。具体算术运算的吞吐量详见CUDA C++编程指南。
11.1.1. 除法模运算
注意
低优先级： 使用位移操作来避免昂贵的除法和取模计算。
整数除法和模运算特别耗费资源，应尽可能避免或用位运算替代：如果\(n\)是2的幂次方，( \(i/n\) )等价于( \(i \gg {log2}(n)\) )，而( \(i\% n\) )等价于( \(i\&\left( {n - 1} \right)\) )。
如果n是字面量，编译器将执行这些转换。（更多信息，请参阅CUDA C++编程指南中的性能指南部分）。
11.1.2. 循环计数器：有符号与无符号对比
注意
中低优先级： 使用有符号整数而非无符号整数作为循环计数器。
在C语言标准中，无符号整数溢出的语义有明确定义，而有符号整数溢出会导致未定义结果。因此，编译器对有符号算术运算可以进行比无符号算术更激进的优化。这一点在循环计数器上尤为值得注意：由于循环计数器的值通常总是正数，人们可能会倾向于将计数器声明为无符号类型。但为了获得稍好的性能，应该将其声明为有符号类型。
例如，考虑以下代码：
for (i = 0; i < n; i++) {
out[i] = in[offset + stride*i];
}
在这里，子表达式 stride*i 可能会溢出32位整数，因此如果 i 被声明为无符号类型，溢出语义会阻止编译器应用某些原本可能适用的优化（例如强度折减）。而如果 i 被声明为有符号类型（其溢出语义是未定义的），编译器就有更多余地来使用这些优化。
11.1.3. 倒数平方根
倒数平方根运算应始终显式调用单精度函数rsqrtf()或双精度函数rsqrt()。编译器仅在符合IEEE-754语义规范时，才会将1.0f/sqrtf(x)优化为rsqrtf()。
11.1.4. 其他算术指令
注意
低优先级： 避免自动将双精度浮点数转换为单精度浮点数。
编译器有时需要插入转换指令，这会引入额外的执行周期。这种情况适用于：
对
char或short类型进行操作的函数，其操作数通常需要转换为int类型用作单精度浮点计算输入的双精度浮点常量（定义时未添加任何类型后缀）
后一种情况可以通过使用单精度浮点常量来避免，这些常量用f后缀定义，例如3.141592653589793f、1.0f、0.5f。
对于单精度代码，强烈建议使用float类型和单精度数学函数。
还需注意的是，CUDA数学库的互补误差函数erfcf()在保持完整单精度准确性的同时速度极快。
11.1.5. 小分数参数的指数运算
对于某些分数指数，相比使用pow()函数，通过平方根、立方根及其逆运算可以显著加快幂运算速度。对于那些指数无法精确表示为浮点数的幂运算（例如1/3），这种方法还能提供更准确的结果，因为使用pow()会放大初始的表示误差。
下表中的公式适用于 x >= 0, x != -0，即 signbit(x) == 0。
计算 |
公式 |
|---|---|
x1/9 |
|
x-1/9 |
|
x1/6 |
|
x-1/6 |
|
x1/4 |
|
x-1/4 |
|
x1/3 |
|
x-1/3 |
|
x1/2 |
|
x-1/2 |
|
x2/3 |
|
x-2/3 |
|
x3/4 |
|
x-3/4 |
|
x7/6 |
|
x-7/6 |
|
x5/4 |
|
x-5/4 |
|
x4/3 |
|
x-4/3 |
|
x3/2 |
|
x-3/2 |
|
11.1.6. 数学库
注意
中等优先级：当速度比精度更重要时，使用快速数学库。
支持两种运行时数学运算。可以通过名称区分它们：有些函数名称带有下划线前缀，而另一些则没有（例如__functionName()与functionName()）。遵循__functionName()命名约定的函数直接映射到硬件级别。它们速度更快但精度稍低（例如__sinf(x)和__expf(x)）。遵循functionName()命名约定的函数速度较慢但精度更高（例如sinf(x)和expf(x)）。__sinf(x)、__cosf(x)和__expf(x)的吞吐量远高于sinf(x)、cosf(x)和expf(x)。如果需要减小参数x的幅值，后者的开销会更大（大约慢一个数量级）。此外，在这种情况下，参数归约代码会使用本地内存，由于本地内存的高延迟，可能会进一步影响性能。更多细节请参阅CUDA C++编程指南。
还需注意的是，当需要计算同一参数的正弦和余弦时，应使用sincos指令系列来优化性能：
__sincosf()用于单精度快速数学运算（参见下一段）sincosf()用于常规单精度计算sincos()双精度版本
-use_fast_math 编译器选项会将每个 functionName() 调用强制转换为等效的 __functionName() 调用。该选项还会禁用单精度非规格化数支持，并普遍降低单精度除法的精度。这是一种激进的优化手段，可能会同时降低数值精度并改变特殊情况的处理方式。更稳健的做法是仅在性能提升显著且可接受行为改变的情况下，有选择性地引入快速内置函数调用。请注意此开关仅对单精度浮点数有效。
注意
中等优先级：在可能的情况下，优先使用更快、更专用的数学函数，而不是较慢、更通用的函数。
对于小的整数幂（如x2或x3），显式乘法几乎肯定比使用通用求幂例程（如pow()）更快。虽然编译器优化改进不断缩小这一差距，但显式乘法（或使用等效的专用内联函数或宏）仍可能具有显著优势。当需要计算同一底数的多个幂时（例如在相近位置同时计算x2和x5），这种优势会更加明显，因为这有助于编译器进行公共子表达式消除（CSE）优化。
对于以2或10为底的指数运算，请使用函数exp2()或expf2()以及exp10()或expf10()，而不是pow()或powf()函数。由于通用指数运算中存在大量特殊情况，并且难以在基数和指数的整个范围内实现良好的精度，pow()和powf()在寄存器压力和指令数量方面都是开销较大的函数。另一方面，函数exp2()、exp2f()、exp10()和exp10f()在性能方面与exp()和expf()相似，并且可能比等效的pow()/powf()快十倍。
对于指数为1/3的幂运算，请使用cbrt()或cbrtf()函数，而不是通用的幂函数pow()或powf()，因为前者比后者快得多。同样，对于指数为-1/3的幂运算，请使用rcbrt()或rcbrtf()。
将sin(π*替换为sinpi(，cos(π*替换为cospi(，以及sincos(π*替换为sincospi(。这种替换在精度和性能方面都具有优势。举个具体例子，若要以度而非弧度计算正弦函数，请使用sinpi(x/180.0)。同理，当函数参数形式为π*时，单精度函数sinpif()、cospif()和sincospif()应替代对sinf()、cosf()和sincosf()的调用。（sinpi()相比sin()的性能优势源于简化的参数归约过程；精度优势则是因为sinpi()仅隐式乘以π，实际上使用了无限精度的数学π而非单精度或双精度的近似值。）
11.2. 内存指令
注意
高优先级： 尽量减少全局内存的使用。尽可能优先使用共享内存访问。
内存指令包括任何从共享内存、局部内存或全局内存读取或写入的指令。当访问未缓存的局部或全局内存时，会有数百个时钟周期的内存延迟。
例如，以下示例代码中的赋值运算符具有高吞吐量，但关键是从全局内存读取数据存在数百个时钟周期的延迟：
__shared__ float shared[32];
__device__ float device[32];
shared[threadIdx.x] = device[threadIdx.x];
如果在线程调度器等待全局内存访问完成的同时，有足够多的独立算术指令可以发出，那么大部分全局内存延迟可以被隐藏。然而，最好尽可能避免访问全局内存。
12. 控制流程
12.1. 分支与发散
注意
高优先级：避免在同一warp内出现不同的执行路径。
流程控制指令（if、switch、do、for、while）会通过导致同一warp内的线程发生分支（即执行不同路径）而显著影响指令吞吐量。当这种情况发生时，不同的执行路径必须被分别执行，这会增加该warp执行指令的总数。
为了在线程ID影响控制流的情况下获得最佳性能，应编写控制条件以最小化发散warp的数量。
这是可行的，因为正如《CUDA C++编程指南》中的SIMT架构所述，跨块的warp分布是确定性的。一个简单的例子是当控制条件仅依赖于(threadIdx / WSIZE)，其中WSIZE表示warp大小。
在这种情况下，由于控制条件与warp完美对齐，因此没有warp发生分歧。
对于仅包含少量指令的分支，warp发散通常只会导致轻微的性能损失。例如，编译器可能使用谓词化(predication)来避免实际的分支跳转。相反，所有指令都会被调度执行，但通过每个线程的条件码或谓词来控制哪些线程执行这些指令。谓词为假的线程不会写入结果，也不会计算地址或读取操作数。
从Volta架构开始，独立线程调度允许warp在数据依赖的条件块之外保持发散状态。可以使用显式的__syncwarp()来确保warp已重新收敛以执行后续指令。
12.2. 分支预测
注意
低优先级： 让编译器更容易使用分支预测来代替循环或控制语句。
有时，编译器可能会通过使用分支预测而非循环展开或优化掉if或switch语句。在这些情况下，不会出现线程束发散的情况。程序员也可以通过以下方式控制循环展开
#pragma unroll
有关此编译指令的更多信息，请参阅CUDA C++编程指南。
在使用分支预测时，依赖于控制条件的指令都不会被跳过。相反，每条此类指令都与一个每线程条件码（或称为谓词）相关联，该谓词会根据控制条件设置为真或假。尽管这些指令都被调度执行，但只有谓词为真的指令才会实际执行。谓词为假的指令不会写入结果，也不会计算地址或读取操作数。
编译器仅在分支条件控制的指令数量小于或等于特定阈值时，才会将分支指令替换为谓词化指令。
13. 部署CUDA应用程序
在完成应用程序一个或多个组件的GPU加速后，可以将结果与原始预期进行比较。回想一下，最初的评估步骤允许开发者确定通过加速给定热点可能获得的最大潜在加速上限。
在着手解决其他热点问题以提高整体加速效果之前，开发者应考虑将部分并行化的实现方案推进到生产环境。这样做有多方面重要意义：例如，它能让用户尽早从投资中获益（即使是部分加速也极具价值），同时通过采用渐进式而非颠覆性的应用改造方案，为开发者和用户将风险降至最低。
14. 理解编程环境
随着每一代NVIDIA处理器的推出，GPU都会新增一些CUDA可以利用的功能特性。因此，理解该架构的特性非常重要。
程序员需要注意两个版本号。第一个是计算能力，第二个是CUDA运行时和CUDA驱动API的版本号。
14.1. CUDA计算能力
计算能力描述了硬件的特性，反映了设备支持的指令集以及其他规格，例如每个块的最大线程数和每个多处理器的寄存器数量。更高版本的计算能力是较低（即较早）版本的超集，因此它们是向后兼容的。
可以通过编程方式查询设备中GPU的计算能力，如deviceQuery CUDA示例所示。该程序的输出如图16所示。这些信息是通过调用cudaGetDeviceProperties()并访问其返回结构体中的信息获得的。
图16 deviceQuery报告的CUDA配置数据示例
计算能力的主版本号和次版本号显示在图16的第七行。该系统的设备0具有计算能力7.0。
有关各种GPU计算能力的更多详细信息，请参阅《CUDA C++编程指南》中的支持CUDA的GPU和计算能力部分。开发者尤其需要注意设备上的多处理器数量、寄存器数量、可用内存大小以及设备的任何特殊功能。
14.2. 附加硬件数据
某些硬件特性并未通过计算能力进行描述。例如，在大多数（但并非全部）GPU上，无论计算能力如何，都可以实现内核执行与主机和设备之间异步数据传输的重叠操作。在这种情况下，可以调用cudaGetDeviceProperties()来确定设备是否支持特定功能。例如，设备属性结构中的asyncEngineCount字段表示是否支持内核执行与数据传输的重叠（如果支持，还能显示可能的并发传输数量）；同样地，canMapHostMemory字段表示是否支持零拷贝数据传输。
14.3. 选择哪个计算能力目标
要针对特定版本的NVIDIA硬件和CUDA软件，请使用nvcc的-arch、-code和-gencode选项。例如，使用warp shuffle操作的代码必须使用-arch=sm_30（或更高计算能力）进行编译。
有关同时为多代支持CUDA的设备构建代码时所用标志的进一步讨论，请参阅构建以实现最大兼容性。
14.4. CUDA运行时
CUDA软件环境的主机运行时组件仅能被主机函数使用。它提供以下功能处理：
设备管理
上下文管理
内存管理
代码模块管理
执行控制
纹理引用管理
与OpenGL和Direct3D的互操作性
与较低级别的CUDA Driver API相比，CUDA Runtime通过提供隐式初始化、上下文管理和设备代码模块管理，极大地简化了设备管理。由nvcc生成的C++主机代码使用了CUDA Runtime，因此链接到此代码的应用程序将依赖于CUDA Runtime；同样，任何使用cuBLAS、cuFFT和其他CUDA Toolkit库的代码也将依赖于CUDA Runtime，这些库在内部使用了它。
构成CUDA运行时API的函数在《CUDA工具包参考手册》中有详细说明。
CUDA运行时负责在内核启动前处理内核加载、设置内核参数和启动配置。隐式的驱动程序版本检查、代码初始化、CUDA上下文管理、CUDA模块管理（将cubin映射到函数）、内核配置以及参数传递都由CUDA运行时执行。
它包含两个主要部分：
一个C风格函数接口 (
cuda_runtime_api.h)。基于C风格函数构建的C++风格便捷封装器(
cuda_runtime.h)。
有关运行时API的更多信息，请参阅CUDA C++编程指南中的CUDA运行时章节。
15. CUDA兼容性开发者指南
CUDA Toolkit 每月发布一次，提供新功能、性能改进和关键错误修复。CUDA兼容性允许用户更新最新的CUDA Toolkit软件（包括编译器、库和工具），而无需更新整个驱动程序堆栈。
CUDA软件环境由三部分组成：
CUDA工具包（包含库、CUDA运行时和开发工具）——供开发者构建CUDA应用程序的软件开发套件。
CUDA驱动 - 用于运行CUDA应用程序的用户态驱动组件（例如Linux系统中的libcuda.so）。
NVIDIA GPU设备驱动程序 - 用于NVIDIA GPU的内核模式驱动组件。
在Linux系统上，CUDA驱动程序和内核模式组件被打包在NVIDIA显示驱动程序包中一起提供。如图1所示。
图17 CUDA组件
CUDA编译器(nvcc)提供了一种处理CUDA和非CUDA代码的方法(通过拆分和引导编译)，与CUDA运行时一起构成了CUDA编译器工具链的一部分。CUDA运行时API为开发者提供了高级C++接口，用于简化设备管理、内核执行等操作，而CUDA驱动API(CUDA Driver API)则为应用程序提供了面向NVIDIA硬件的底层编程接口。
基于这些技术构建的是CUDA库，其中一些包含在CUDA工具包中，而其他如cuDNN可能独立于CUDA工具包发布。
15.1. CUDA工具包版本控制
从CUDA 11开始，工具包版本采用行业标准的语义化版本控制方案：.X.Y.Z，其中：
.X代表主版本号 - API已变更且二进制兼容性被破坏。
.Y代表次要版本 - 引入新API、弃用旧API，可能会破坏源代码兼容性但保持二进制兼容性。
.Z 代表发布/补丁版本 - 新的更新和补丁将会增加这个数字。
工具包中的每个组件都建议采用语义化版本控制。从CUDA 11.3开始，NVRTC也采用了语义化版本控制。我们将在文档后续部分标注其中一些组件的版本信息。工具包中各组件的版本信息可查阅此表格。
因此，CUDA平台的兼容性旨在解决以下几种场景：
对于企业或数据中心中运行GPU的生产系统来说，升级NVIDIA驱动程序可能较为复杂，需要提前规划。延迟部署新版NVIDIA驱动程序意味着这些系统的用户可能无法使用CUDA版本中的新功能。而新CUDA版本无需更新驱动程序，则意味着可以更快地向用户提供新版软件。
许多基于CUDA构建的软件库和应用程序（例如数学库或深度学习框架）并不直接依赖于CUDA运行时、编译器或驱动程序。在这种情况下，用户或开发者仍然可以受益，无需升级整个CUDA工具包或驱动程序即可使用这些库或框架。
升级依赖项容易出错且耗时，在某些极端情况下甚至可能改变程序的语义。持续使用最新的CUDA Toolkit重新编译意味着强制应用程序产品的终端客户进行升级。虽然包管理器简化了这个过程，但仍可能出现意外问题，如果发现错误，就需要重复上述升级流程。
CUDA支持多种兼容性选项：
CUDA 10首次引入的CUDA前向兼容升级功能，旨在让用户能够获取新的CUDA特性，并在安装旧版NVIDIA数据中心驱动程序的系统上运行基于新CUDA版本构建的应用程序。
-
首次在CUDA 11.1中引入的CUDA增强兼容性提供了两大优势：
通过在CUDA工具包各组件中采用语义化版本控制，应用程序可以针对某个CUDA次要版本（例如11.1）进行构建，并兼容该主版本系列中所有未来的次要版本（即11.x）。
CUDA运行时放宽了最低驱动程序版本检查，因此在升级到新的次要版本时不再需要升级驱动程序。
CUDA驱动程序确保为已编译的CUDA应用程序保持向后二进制兼容性。使用旧至3.2版本的CUDA工具包编译的应用程序可在较新的驱动程序上运行。
15.2. 源代码兼容性
我们将源兼容性定义为库提供的一组保证，即针对特定版本库（使用SDK）构建的正确应用程序在安装新版本SDK后，仍能继续构建和运行而不出错。
CUDA驱动程序和CUDA运行时在不同SDK版本之间不保证源代码兼容性。API可能会被弃用或移除。因此，在旧版本工具包上成功编译的应用程序可能需要修改才能针对新版本工具包进行编译。
开发者会通过弃用通知和文档机制获知当前或即将进行的任何变更。这并不意味着使用旧版工具包编译的应用程序二进制文件将不再受支持。应用程序二进制文件依赖于CUDA Driver API接口，尽管CUDA Driver API本身可能也会随工具包版本而变化，但CUDA保证CUDA Driver API接口的二进制兼容性。
15.3. 二进制兼容性
我们将二进制兼容性定义为库提供的一组保证，即针对该库开发的应用程序在动态链接到不同版本的库时仍能继续工作。
CUDA驱动API采用版本化的C风格ABI接口，这确保了针对旧版驱动（如CUDA 3.2）开发的应用程序仍能在现代驱动（例如随CUDA 11.0发布的版本）上正常运行。这意味着虽然应用程序源代码可能需要重新针对新版CUDA工具包进行编译才能使用新功能，但系统安装的新版驱动组件始终会向下兼容现有应用程序及其功能。
因此，CUDA Driver API 是二进制兼容的（操作系统加载器可以选用新版本而应用程序仍能正常工作），但不是源代码兼容的（针对新版SDK重新构建应用程序可能需要修改源代码）。
图18 CUDA工具包与最低驱动程序版本要求
在深入探讨这个话题之前，开发者理解最低驱动程序版本的概念及其可能产生的影响非常重要。
每个版本的CUDA工具包（及运行时）都需要对应最低版本的NVIDIA驱动程序。基于某个CUDA工具包版本编译的应用程序，仅能在满足该工具包版本最低驱动要求的系统上运行。在CUDA 11.0之前，工具包要求的最低驱动版本与其随附的驱动程序版本相同。
因此，当应用程序使用CUDA 11.0构建时，它只能在安装了R450或更高版本驱动的系统上运行。如果此类应用程序在安装了R418驱动的系统上运行，CUDA初始化将返回错误，如下例所示。
在本示例中，deviceQuery示例程序使用CUDA 11.1编译，并在装有R418驱动的系统上运行。此场景中，由于不满足最低驱动程序要求，CUDA初始化返回错误。
ubuntu@:~/samples/1_Utilities/deviceQuery
$ make
/usr/local/cuda-11.1/bin/nvcc -ccbin g++ -I../../common/inc -m64 -gencode arch=compute_35,code=sm_35 -gencode arch=compute_37,code=sm_37 -gencode arch=compute_50,code=sm_50 -gencode arch=compute_52,code=sm_52 -gencode arch=compute_60,code=sm_60 -gencode arch=compute_61,code=sm_61 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_75,code=sm_75 -gencode arch=compute_80,code=sm_80 -gencode arch=compute_86,code=sm_86 -gencode arch=compute_86,code=compute_86 -o deviceQuery.o -c deviceQuery.cpp
/usr/local/cuda-11.1/bin/nvcc -ccbin g++ -m64 -gencode arch=compute_35,code=sm_35 -gencode arch=compute_37,code=sm_37 -gencode arch=compute_50,code=sm_50 -gencode arch=compute_52,code=sm_52 -gencode arch=compute_60,code=sm_60 -gencode arch=compute_61,code=sm_61 -gencode arch=compute_70,code=sm_70 -gencode arch=compute_75,code=sm_75 -gencode arch=compute_80,code=sm_80 -gencode arch=compute_86,code=sm_86 -gencode arch=compute_86,code=compute_86 -o deviceQuery deviceQuery.o
$ nvidia-smi
+-----------------------------------------------------------------------------+
| NVIDIA-SMI 418.165.02 Driver Version: 418.165.02 CUDA Version: 10.1 |
|-------------------------------+----------------------+----------------------+
| GPU Name Persistence-M| Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap| Memory-Usage | GPU-Util Compute M. |
|===============================+======================+======================|
| 0 Tesla T4 On | 00000000:00:1E.0 Off | 0 |
| N/A 42C P0 28W / 70W | 0MiB / 15079MiB | 0% Default |
+-------------------------------+----------------------+----------------------+
+-----------------------------------------------------------------------------+
| Processes: GPU Memory |
| GPU PID Type Process name Usage |
|=============================================================================|
| No running processes found |
+-----------------------------------------------------------------------------+
$ samples/bin/x86_64/linux/release/deviceQuery
samples/bin/x86_64/linux/release/deviceQuery Starting...
CUDA Device Query (Runtime API) version (CUDART static linking)
cudaGetDeviceCount returned 3
-> initialization error
Result = FAIL
有关最低驱动程序版本及工具包随附驱动程序版本的详细信息，请参阅CUDA Toolkit Release Notes。
15.3.1. CUDA二进制(cubin)兼容性
一个略微相关但重要的话题是CUDA中跨GPU架构的应用程序二进制兼容性。
CUDA C++为熟悉C++编程语言的用户提供了一条简单路径，使其能够轻松编写在设备上执行的程序。内核(kernels)可以使用CUDA指令集架构PTX编写，该架构在PTX参考手册中有详细说明。然而，使用高级编程语言如C++通常更为高效。无论采用哪种方式，内核都必须通过nvcc编译成二进制代码(称为cubins)才能在设备上执行。
cubin文件是与架构相关的。cubin的二进制兼容性保证从一个计算能力的小版本升级到下一个版本时有效，但不保证从一个小版本回退到前一个版本或跨越主要计算能力版本时有效。换句话说，为计算能力X.y生成的cubin对象只能在计算能力为X.z（其中z≥y）的设备上执行。
要在特定计算能力的设备上执行代码，应用程序必须加载与该计算能力兼容的二进制或PTX代码。为了实现可移植性，即能够在具有更高计算能力的未来GPU架构上执行代码（目前还无法生成对应的二进制代码），应用程序必须加载PTX代码，这些代码将由NVIDIA驱动程序为这些未来设备进行即时编译。
有关cubin、PTX和应用兼容性的更多信息，请参阅CUDA C++编程指南。
15.4. CUDA跨小版本兼容性
通过采用语义化版本控制，从CUDA 11开始，CUDA工具包中的组件将在工具包的次要版本之间保持二进制兼容性。为了在次要版本之间保持二进制兼容性，CUDA运行时不再要求每个次要版本都提升最低驱动程序版本要求——这一调整仅在主版本发布时才会发生。
新工具链需要新最低驱动程序的主要原因之一是为了处理PTX代码的即时编译和二进制代码的即时链接。
在本节中，我们将回顾利用CUDA平台的兼容性功能时可能需要新用户工作流程的使用模式。
15.4.1. CUDA小版本中的现有CUDA应用程序
$ nvidia-smi
+-----------------------------------------------------------------------------+
| NVIDIA-SMI 450.80.02 Driver Version: 450.80.02 CUDA Version: 11.0 |
|-------------------------------+----------------------+----------------------+
| GPU Name Persistence-M| Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap| Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|===============================+======================+======================|
| 0 Tesla T4 On | 00000000:00:1E.0 Off | 0 |
| N/A 39C P8 9W / 70W | 0MiB / 15109MiB | 0% Default |
| | | N/A |
+-------------------------------+----------------------+----------------------+
+-----------------------------------------------------------------------------+
| Processes: |
| GPU GI CI PID Type Process name GPU Memory |
| ID ID Usage |
|=============================================================================|
| No running processes found |
+-----------------------------------------------------------------------------+
当我们的CUDA 11.1应用程序（即静态链接了cudart 11.1）在系统上运行时，我们发现即使驱动程序报告的是11.0版本，它也能成功运行——也就是说，不需要在系统上更新驱动程序或其他工具包组件。
$ samples/bin/x86_64/linux/release/deviceQuery
samples/bin/x86_64/linux/release/deviceQuery Starting...
CUDA Device Query (Runtime API) version (CUDART static linking)
Detected 1 CUDA Capable device(s)
Device 0: "Tesla T4"
CUDA Driver Version / Runtime Version 11.0 / 11.1
CUDA Capability Major/Minor version number: 7.5
...<snip>...
deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 11.0, CUDA Runtime Version = 11.1, NumDevs = 1
Result = PASS
通过使用新版本的CUDA，用户可以受益于新的CUDA编程模型API、编译器优化和数学库功能。
以下部分讨论了一些注意事项和考量因素。
15.4.1.1. 处理新的CUDA功能和驱动程序API
CUDA API的一个子集不需要新驱动程序，它们都可以在没有任何驱动程序依赖的情况下使用。例如，cuMemMap API或CUDA 11.0之前引入的任何API（如cudaDeviceSynchronize）都不需要升级驱动程序。要使用次要版本中引入的其他CUDA API（需要新驱动程序），用户必须实现回退方案或优雅地处理失败情况。这种情况与当前开发者使用宏根据CUDA版本编译功能的现状并无不同。用户应参考CUDA头文件和文档以了解版本中引入的新CUDA API。
当使用工具包次要版本中提供的功能时，如果应用程序运行在较旧的CUDA驱动程序上，该功能可能在运行时不可用。希望利用此类功能的用户应在代码中通过动态检查来查询其可用性：
static bool hostRegisterFeatureSupported = false;
static bool hostRegisterIsDeviceAddress = false;
static error_t cuFooFunction(int *ptr)
{
int *dptr = null;
if (hostRegisterFeatureSupported) {
cudaHostRegister(ptr, size, flags);
if (hostRegisterIsDeviceAddress) {
qptr = ptr;
}
else {
cudaHostGetDevicePointer(&qptr, ptr, 0);
}
}
else {
// cudaMalloc();
// cudaMemcpy();
}
gemm<<<1,1>>>(dptr);
cudaDeviceSynchronize();
}
int main()
{
// rest of code here
cudaDeviceGetAttribute(
&hostRegisterFeatureSupported,
cudaDevAttrHostRegisterSupported,
0);
cudaDeviceGetAttribute(
&hostRegisterIsDeviceAddress,
cudaDevAttrCanUseHostPointerForRegisteredMem,
0);
cuFooFunction(/* malloced pointer */);
}
或者，如果应用程序的界面在没有新CUDA驱动的情况下完全无法运行，那么最好立即返回错误：
#define MIN_VERSION 11010
cudaError_t foo()
{
int version = 0;
cudaGetDriverVersion(&version);
if (version < MIN_VERSION) {
return CUDA_ERROR_INSUFFICIENT_DRIVER;
}
// proceed as normal
}
新增了一个错误代码，用于指示当前运行的驱动程序缺少所需功能：cudaErrorCallRequiresNewerDriver。
15.4.1.2. 使用PTX
PTX定义了一个用于通用并行线程执行的虚拟机和指令集架构。PTX程序在加载时通过CUDA驱动程序中的JIT编译器转换为目标硬件指令集。由于PTX由CUDA驱动程序编译，新工具链生成的PTX可能与旧版CUDA驱动程序不兼容。当PTX用于未来设备兼容性时（最常见情况）这不是问题，但在用于运行时编译时可能会导致问题。
对于继续使用PTX的代码，为了支持在较旧驱动程序上编译，您的代码必须首先通过静态ptxjitcompiler库或NVRTC转换为设备代码，并选择为特定架构（如sm_80）而非虚拟架构（如compute_80）生成代码。针对此工作流程，CUDA Toolkit附带了一个新的nvptxcompiler_static库。
我们可以在以下示例中看到这种用法：
char* compilePTXToNVElf()
{
nvPTXCompilerHandle compiler = NULL;
nvPTXCompileResult status;
size_t elfSize, infoSize, errorSize;
char *elf, *infoLog, *errorLog;
int minorVer, majorVer;
const char* compile_options[] = { "--gpu-name=sm_80",
"--device-debug"
};
nvPTXCompilerGetVersion(&majorVer, &minorVer);
nvPTXCompilerCreate(&compiler, (size_t)strlen(ptxCode), ptxCode);
status = nvPTXCompilerCompile(compiler, 2, compile_options);
if (status != NVPTXCOMPILE_SUCCESS) {
nvPTXCompilerGetErrorLogSize(compiler, (void*)&errorSize);
if (errorSize != 0) {
errorLog = (char*)malloc(errorSize+1);
nvPTXCompilerGetErrorLog(compiler, (void*)errorLog);
printf("Error log: %s\n", errorLog);
free(errorLog);
}
exit(1);
}
nvPTXCompilerGetCompiledProgramSize(compiler, &elfSize));
elf = (char*)malloc(elfSize);
nvPTXCompilerGetCompiledProgram(compiler, (void*)elf);
nvPTXCompilerGetInfoLogSize(compiler, (void*)&infoSize);
if (infoSize != 0) {
infoLog = (char*)malloc(infoSize+1);
nvPTXCompilerGetInfoLog(compiler, (void*)infoLog);
printf("Info log: %s\n", infoLog);
free(infoLog);
}
nvPTXCompilerDestroy(&compiler);
return elf;
}
15.4.1.3. 动态代码生成
NVRTC是一个用于CUDA C++的运行时编译库。它接收字符串形式的CUDA C++源代码，并创建可用于获取PTX的句柄。由NVRTC生成的PTX字符串可以通过cuModuleLoadData和cuModuleLoadDataEx加载。
目前尚不支持处理可重定位对象，因此CUDA驱动程序中的cuLink*系列API无法与增强兼容性功能配合使用。当前这些API需要使用与CUDA运行时版本匹配的升级版驱动程序。
如PTX部分所述，PTX到设备代码的编译与CUDA驱动程序一起存在，因此生成的PTX可能比部署系统上驱动程序支持的版本更新。在使用NVRTC时，建议首先按照PTX用户工作流程中概述的步骤将生成的PTX代码转换为最终的设备代码。这可以确保您的代码兼容。或者，从CUDA 11.1开始，NVRTC可以直接生成cubins。使用新API的应用程序可以直接使用驱动程序API cuModuleLoadData 和 cuModuleLoadDataEx加载最终的设备代码。
NVRTC过去仅通过-arch选项支持虚拟架构，因为它只生成PTX。现在它也将支持实际架构以生成SASS。该接口经过增强，如果指定了实际架构，可以检索PTX或cubin。
以下示例展示了如何调整现有示例以使用新功能，在本例中由USE_CUBIN宏进行保护：
#include <nvrtc.h>
#include <cuda.h>
#include <iostream>
void NVRTC_SAFE_CALL(nvrtcResult result) {
if (result != NVRTC_SUCCESS) {
std::cerr << "\nnvrtc error: " << nvrtcGetErrorString(result) << '\n';
std::exit(1);
}
}
void CUDA_SAFE_CALL(CUresult result) {
if (result != CUDA_SUCCESS) {
const char *msg;
cuGetErrorName(result, &msg);
std::cerr << "\ncuda error: " << msg << '\n';
std::exit(1);
}
}
const char *hello = " \n\
extern \"C\" __global__ void hello() { \n\
printf(\"hello world\\n\"); \n\
} \n";
int main()
{
nvrtcProgram prog;
NVRTC_SAFE_CALL(nvrtcCreateProgram(&prog, hello, "hello.cu", 0, NULL, NULL));
#ifdef USE_CUBIN
const char *opts[] = {"-arch=sm_70"};
#else
const char *opts[] = {"-arch=compute_70"};
#endif
nvrtcResult compileResult = nvrtcCompileProgram(prog, 1, opts);
size_t logSize;
NVRTC_SAFE_CALL(nvrtcGetProgramLogSize(prog, &logSize));
char *log = new char[logSize];
NVRTC_SAFE_CALL(nvrtcGetProgramLog(prog, log));
std::cout << log << '\n';
delete[] log;
if (compileResult != NVRTC_SUCCESS)
exit(1);
size_t codeSize;
#ifdef USE_CUBIN
NVRTC_SAFE_CALL(nvrtcGetCUBINSize(prog, &codeSize));
char *code = new char[codeSize];
NVRTC_SAFE_CALL(nvrtcGetCUBIN(prog, code));
#else
NVRTC_SAFE_CALL(nvrtcGetPTXSize(prog, &codeSize));
char *code = new char[codeSize];
NVRTC_SAFE_CALL(nvrtcGetPTX(prog, code));
#endif
NVRTC_SAFE_CALL(nvrtcDestroyProgram(&prog));
CUdevice cuDevice;
CUcontext context;
CUmodule module;
CUfunction kernel;
CUDA_SAFE_CALL(cuInit(0));
CUDA_SAFE_CALL(cuDeviceGet(&cuDevice, 0));
CUDA_SAFE_CALL(cuCtxCreate(&context, 0, cuDevice));
CUDA_SAFE_CALL(cuModuleLoadDataEx(&module, code, 0, 0, 0));
CUDA_SAFE_CALL(cuModuleGetFunction(&kernel, module, "hello"));
CUDA_SAFE_CALL(cuLaunchKernel(kernel, 1, 1, 1, 1, 1, 1, 0, NULL, NULL, 0));
CUDA_SAFE_CALL(cuCtxSynchronize());
CUDA_SAFE_CALL(cuModuleUnload(module));
CUDA_SAFE_CALL(cuCtxDestroy(context));
delete[] code;
}
15.4.1.4. 构建小版本兼容库的建议
我们建议静态链接CUDA运行时以最小化依赖项。请确保您的库不会在既定的ABI合约之外泄露依赖项、破坏或命名空间等。
为您的库的soname遵循语义化版本控制。拥有语义化版本化的ABI意味着需要维护和版本化接口。当更改影响此ABI契约时，库应遵循语义规则并增加版本号。缺少依赖项也会导致二进制兼容性中断，因此您应为依赖这些接口的功能提供回退或保护措施。当发生ABI破坏性变更（如API弃用和修改）时，应增加主版本号。新API可以在次版本中添加。
有条件地使用功能以保持与旧版驱动程序的兼容性。如果未使用新功能（或在使用时提供备用方案作为条件），您将能够保持兼容。
不要暴露可能变化的ABI结构。嵌入结构大小的指针是更好的解决方案。
当链接工具包中的动态库时，所用库的版本必须等于或高于应用程序链接过程中涉及的任何一个组件所需的版本。例如，如果您链接的是CUDA 11.1动态运行时库，并使用了11.1版本的功能，同时还使用了一个单独的共享库（该库链接的是需要11.2版本功能的CUDA 11.2动态运行时库），那么最终链接步骤必须包含CUDA 11.2或更高版本的动态运行时库。
15.4.1.5. 在应用程序中利用次要版本兼容性的建议
某些功能可能不可用，因此您应在适用时进行查询。这对于构建与GPU架构、平台和编译器无关的应用程序很常见。不过我们现在还要在其中加入"底层驱动程序"这个因素。
与前一节关于库构建建议类似，如果使用CUDA运行时，我们建议在构建应用程序时静态链接到CUDA运行时。当直接使用驱动程序API时，我们建议使用新的驱动程序入口点访问API（cuGetProcAddress），其文档在此处：CUDA Driver API :: CUDA Toolkit Documentation。
使用共享或静态库时，请遵循该库的发布说明以确定其是否支持次要版本兼容性。
16. 部署准备
16.1. 检测CUDA可用性
在部署CUDA应用程序时，通常需要确保即使目标机器没有支持CUDA的GPU和/或安装足够版本的NVIDIA驱动程序，应用程序仍能正常运行。（针对已知配置的单一机器进行开发的开发者可以选择跳过本节。）
检测支持CUDA的GPU
当应用程序需要部署到配置未知的任意目标机器时，应用程序应明确检测是否存在支持CUDA的GPU，以便在无可用设备时采取适当措施。cudaGetDeviceCount()函数可用于查询可用设备数量。与所有CUDA运行时API函数一样，如果没有支持CUDA的GPU，该函数会优雅地失败并向应用程序返回cudaErrorNoDevice；如果未安装适当版本的NVIDIA驱动，则返回cudaErrorInsufficientDriver。如果cudaGetDeviceCount()报告错误，应用程序应回退到备用代码路径。
一个配备多块GPU的系统可能包含不同硬件版本和性能的GPU。当在同一应用程序中使用多块GPU时，建议使用相同类型的GPU，而不是混合不同代次的硬件。cudaChooseDevice()函数可用于选择最符合所需功能集的设备。
检测硬件和软件配置
当应用程序依赖特定硬件或软件功能的可用性来实现某些功能时，可以查询CUDA API以获取有关可用设备配置和已安装软件版本的详细信息。
cudaGetDeviceProperties()函数可报告可用设备的各种特性，包括设备的CUDA计算能力（另请参阅CUDA C++编程指南的计算能力章节）。有关如何查询可用CUDA软件API版本的详细信息，请参阅版本管理。
16.2. 错误处理
所有CUDA运行时API调用都会返回一个类型为cudaError_t的错误代码；如果没有发生错误，返回值将等于cudaSuccess。（例外情况是内核启动，它返回void；以及cudaGetErrorString()，它返回一个描述传入的cudaError_t代码的字符串。）CUDA工具包库（cuBLAS、cuFFT等）同样会返回它们自己的一组错误代码。
由于某些CUDA API调用和所有内核启动相对于主机代码是异步的，错误也可能异步报告给主机；这通常发生在主机和设备下一次相互同步时，例如在调用cudaMemcpy()或cudaDeviceSynchronize()期间。
始终检查所有CUDA API函数的错误返回值，即使是预期不会失败的函数，因为这能让应用程序在错误发生时尽快检测并恢复。对于使用<<<...>>>语法启动的内核（该语法不返回任何错误代码），应在内核启动后立即检查cudaGetLastError()的返回码。不检查CUDA API错误的应用程序有时会运行完成，但未注意到GPU计算的数据不完整、无效或未初始化。
注意
CUDA工具包示例提供了多个辅助函数，用于检查各种CUDA API的错误；这些辅助函数位于CUDA工具包中的samples/common/inc/helper_cuda.h文件内。
16.3. 构建最大兼容性
每一代支持CUDA的设备都有一个关联的计算能力版本，用于指示该设备支持的功能集（参见CUDA计算能力）。在构建文件时，可以向nvcc编译器指定一个或多个计算能力版本；针对应用程序目标GPU的原生计算能力进行编译非常重要，这能确保应用程序内核获得最佳性能，并能够利用特定GPU世代所提供的功能。
当一个应用程序被构建为同时支持多种计算能力（通过使用多个-gencode标志实例传递给nvcc时），针对指定计算能力的二进制文件会被合并到可执行文件中，CUDA驱动程序在运行时根据当前设备的计算能力自动选择最合适的二进制版本。如果没有合适的原生二进制文件（cubin）可用，但存在中间表示PTX代码（针对抽象虚拟指令集设计，用于实现前向兼容性），那么内核将通过即时编译（JIT）方式（参见编译器JIT缓存管理工具）从PTX代码动态编译为设备对应的原生cubin。如果连PTX代码也不可用，则内核启动将失败。
Windows
nvcc.exe -ccbin "C:\vs2008\VC\bin"
-Xcompiler "/EHsc /W3 /nologo /O2 /Zi /MT"
-gencode=arch=compute_30,code=sm_30
-gencode=arch=compute_35,code=sm_35
-gencode=arch=compute_50,code=sm_50
-gencode=arch=compute_60,code=sm_60
-gencode=arch=compute_70,code=sm_70
-gencode=arch=compute_75,code=sm_75
-gencode=arch=compute_75,code=compute_75
--compile -o "Release\mykernel.cu.obj" "mykernel.cu"
Mac/Linux
/usr/local/cuda/bin/nvcc
-gencode=arch=compute_30,code=sm_30
-gencode=arch=compute_35,code=sm_35
-gencode=arch=compute_50,code=sm_50
-gencode=arch=compute_60,code=sm_60
-gencode=arch=compute_70,code=sm_70
-gencode=arch=compute_75,code=sm_75
-gencode=arch=compute_75,code=compute_75
-O2 -o mykernel.o -c mykernel.cu
或者，可以使用nvcc命令行选项-arch=sm_XX作为上述更详细的-gencode=命令行选项的简写等价形式：
-gencode=arch=compute_XX,code=sm_XX
-gencode=arch=compute_XX,code=compute_XX
然而，虽然-arch=sm_XX命令行选项确实会默认包含PTX后端目标（因为它隐含了code=compute_XX目标），但它一次只能指定单个目标cubin架构，且无法在同一条nvcc命令行中使用多个-arch=选项，这就是为什么上述示例显式使用-gencode=的原因。
16.4. 分发CUDA运行时和库
CUDA应用程序是基于CUDA运行时库构建的，该库负责设备、内存和内核管理。与CUDA驱动程序不同，CUDA运行时库不保证跨版本的向前或向后二进制兼容性。因此在使用动态链接时，最好将CUDA运行时库与应用程序一起重新分发，或者选择静态链接CUDA运行时库。这样可以确保即使最终用户没有安装与应用程序构建时相同的CUDA工具包，可执行文件仍能正常运行。
注意
当静态链接到CUDA运行时库时，同一应用程序进程中可以同时和平共存多个版本的运行时库。例如，如果应用程序使用一个版本的CUDA运行时，而该应用程序的插件静态链接到另一个不同版本，只要安装的NVIDIA驱动程序能满足两者的需求，这种情况是完全可行的。
静态链接的CUDA运行时
最简单的选择是静态链接到CUDA运行时库。如果使用nvcc在CUDA 5.5及更高版本中进行链接，这是默认选项。静态链接会使可执行文件稍大一些，但它能确保应用程序二进制文件中包含正确版本的运行时库函数，而无需单独重新分发CUDA运行时库。
动态链接的CUDA运行时
如果由于某些原因静态链接CUDA运行时不可行，也可以使用动态链接版本的CUDA运行时库。（这是CUDA 5.0及更早版本中默认且唯一提供的选项。）
当使用CUDA 5.5或更高版本的nvcc链接应用程序时，若需采用动态链接CUDA运行时，请在链接命令行中添加--cudart=shared标志；默认情况下将使用静态链接的CUDA运行时库。
当应用程序动态链接到CUDA Runtime后，此版本的运行时库应随应用程序一起打包。它可以被复制到应用程序可执行文件所在的目录，或安装路径的子目录中。
其他CUDA库
尽管CUDA运行时提供了静态链接的选项，但CUDA工具包中包含的一些库仅以动态链接形式提供。与动态链接版本的CUDA运行时库一样，在分发应用程序时，这些库应该与应用程序可执行文件捆绑在一起。
16.4.1. CUDA工具库再分发许可
CUDA工具包的最终用户许可协议(EULA)允许在特定条款和条件下重新分发许多CUDA库。这使得依赖这些库的应用程序能够重新分发它们构建和测试时使用的库的精确版本，从而避免最终用户可能遇到问题——他们的机器上可能安装了不同版本的CUDA工具包(或者根本没有安装)。详情请参阅EULA。
注意
这不适用于NVIDIA驱动程序；终端用户仍需下载并安装适合其GPU和操作系统的NVIDIA驱动程序。
16.4.1.1. 需要重新分发的文件
在重新分发一个或多个CUDA库的动态链接版本时，准确识别需要重新分发的文件至关重要。以下示例使用CUDA Toolkit 5.5中的cuBLAS库进行说明：
Linux
在Linux的共享库中，有一个名为SONAME的字符串字段，用于表示库的二进制兼容级别。应用程序构建时所依赖库的SONAME必须与随应用程序分发的库文件名相匹配。
例如，在标准的CUDA工具包安装中，文件libcublas.so和libcublas.so.5.5都是指向特定cuBLAS构建版本的符号链接，其命名格式为libcublas.so.5.5.x，其中x表示构建编号（例如libcublas.so.5.5.17）。然而，该库的SONAME被指定为"libcublas.so.5.5"：
$ objdump -p /usr/local/cuda/lib64/libcublas.so | grep SONAME
SONAME libcublas.so.5.5
因此，即使应用程序在链接时使用了-lcublas（未指定版本号），链接时找到的SONAME表明动态加载器在加载应用程序时将寻找名为"libcublas.so.5.5"的文件，因此该文件名（或其符号链接）必须与应用程序一起重新分发。
ldd工具可用于识别应用程序在运行时预期查找的库文件的确切名称，以及在当前库搜索路径下动态加载器加载应用程序时会选择的该库副本的路径（如果有的话）：
$ ldd a.out | grep libcublas
libcublas.so.5.5 => /usr/local/cuda/lib64/libcublas.so.5.5
Mac
在Mac OS X的共享库中，有一个名为install name的字段，用于指定库的预期安装路径和文件名；CUDA库也使用此文件名来指示二进制兼容性。该字段的值会被传递到基于该库构建的应用程序中，并在运行时用于定位正确版本的库。
例如，如果cuBLAS库的安装名称指定为@rpath/libcublas.5.5.dylib，则表示该库版本为5.5，且随应用程序重新分发的此库副本必须命名为libcublas.5.5.dylib，尽管在链接时仅使用了-lcublas（未指定版本号）。此外，该文件应安装到应用程序的@rpath中；详见Where to Install Redistributed CUDA Libraries。
要查看库的安装名称，请使用 otool -L 命令：
$ otool -L a.out
a.out:
@rpath/libcublas.5.5.dylib (...)
Windows
Windows上CUDA库的二进制兼容版本在文件名中有所体现。
例如，一个链接到cuBLAS 5.5的64位应用程序在运行时将寻找cublas64_55.dll，因此这是应该与该应用程序一起分发的文件，尽管应用程序实际链接的是cublas.lib文件。对于32位应用程序，对应的文件将是cublas32_55.dll。
要验证应用程序在运行时期望找到的确切DLL文件名，请使用Visual Studio命令提示符中的dumpbin工具：
$ dumpbin /IMPORTS a.exe
Microsoft (R) COFF/PE Dumper Version 10.00.40219.01
Copyright (C) Microsoft Corporation. All rights reserved.
Dump of file a.exe
File Type: EXECUTABLE IMAGE
Section contains the following imports:
...
cublas64_55.dll
...
16.4.1.2. 重新分发的CUDA库安装位置
一旦确定了用于重新分发的正确库文件，必须配置它们以便安装到应用程序能够找到的位置。
在Windows系统上，如果将CUDA运行时或其他动态链接的CUDA工具库文件与可执行文件放在同一目录下，Windows会自动定位到它们。而在Linux和Mac系统上，应使用-rpath链接器选项来指示可执行文件先在本地路径中搜索这些库文件，然后再搜索系统路径：
Linux/Mac
nvcc -I $(CUDA_HOME)/include
-Xlinker "-rpath '$ORIGIN'" --cudart=shared
-o myprogram myprogram.cu
Windows
nvcc.exe -ccbin "C:\vs2008\VC\bin"
-Xcompiler "/EHsc /W3 /nologo /O2 /Zi /MT" --cudart=shared
-o "Release\myprogram.exe" "myprogram.cu"
注意
可能需要调整-ccbin的值以反映您Visual Studio的安装位置。
要指定库文件分发的备用路径，请使用类似以下的链接器选项：
Linux/Mac
nvcc -I $(CUDA_HOME)/include
-Xlinker "-rpath '$ORIGIN/lib'" --cudart=shared
-o myprogram myprogram.cu
Windows
nvcc.exe -ccbin "C:\vs2008\VC\bin"
-Xcompiler "/EHsc /W3 /nologo /O2 /Zi /MT /DELAY" --cudart=shared
-o "Release\myprogram.exe" "myprogram.cu"
对于Linux和Mac系统，-rpath选项的用法与之前相同。在Windows系统上，则使用/DELAY选项；这要求应用程序在首次调用任何CUDA API函数之前，必须先调用SetDllDirectory()来指定包含CUDA DLL文件的目录。
注意
对于Windows 8系统，应使用SetDefaultDLLDirectories()和AddDllDirectory()而非SetDllDirectory()。有关这些例程的更多信息，请参阅MSDN文档。
17. 部署基础设施工具
17.1. Nvidia-SMI
NVIDIA系统管理界面(nvidia-smi)是一款命令行工具，用于帮助管理和监控NVIDIA GPU设备。该工具允许管理员查询GPU设备状态，并在具备适当权限的情况下修改GPU设备状态。nvidia-smi主要面向Tesla和部分Quadro GPU，同时也为其他NVIDIA GPU提供有限支持。nvidia-smi随NVIDIA GPU显示驱动程序一起发布，支持Linux系统以及64位Windows Server 2008 R2和Windows 7系统。nvidia-smi可将查询信息以XML格式或人类可读的纯文本格式输出到标准输出或文件中。详情请参阅nvidia-smi文档。请注意，新版本的nvidia-smi不保证与旧版本向后兼容。
17.1.1. 可查询状态
- ECC error counts
-
可纠正的单比特错误和可检测的双比特错误都会被报告。错误计数会提供当前启动周期和GPU整个生命周期的数据。
- GPU utilization
-
当前报告了GPU计算资源和内存接口的利用率。
- Active compute process
-
报告了GPU上正在运行的活跃进程列表，包括对应的进程名称/ID以及分配的GPU内存。
- Clocks and performance state
-
报告了多个重要时钟域的最大和当前时钟频率，以及当前GPU性能状态(pstate)。
- Temperature and fan speed
-
报告当前GPU核心温度，以及带有主动散热功能产品的风扇转速。
- Power management
-
当前板卡功耗和功率限制会针对支持这些测量功能的产品进行报告。
- Identification
-
报告各种动态和静态信息，包括主板序列号、PCI设备ID、VBIOS/Inforom版本号和产品名称。
17.1.2. 可变状态
- ECC mode
-
启用和禁用ECC报告。
- ECC reset
-
清除单比特和双比特ECC错误计数。
- Compute mode
-
指示计算进程是否可以在GPU上运行，以及它们是独占运行还是与其他计算进程并发运行。
- Persistence mode
-
指示当没有应用程序连接到GPU时，NVIDIA驱动程序是否保持加载状态。在大多数情况下，最好启用此选项。
- GPU reset
-
通过辅助总线重置重新初始化GPU硬件和软件状态。
17.2. NVML
NVIDIA管理库(NVML)是一个基于C语言的接口，可直接访问通过nvidia-smi暴露的查询和命令，旨在作为构建第三方系统管理应用程序的平台。NVML API随CUDA工具包(自8.0版本起)一起发布，也可在NVIDIA开发者网站上作为GPU部署工具包的一部分独立获取，包含单个头文件、PDF文档、存根库和示例应用程序；详见https://developer.nvidia.com/gpu-deployment-kit。每个新版本的NVML都保持向后兼容性。
为NVML API额外提供了一套Perl和Python绑定。这些绑定暴露了与基于C的接口相同的功能，并提供了向后兼容性。Perl绑定通过CPAN提供，Python绑定则通过PyPI提供。
所有这些产品（nvidia-smi、NVML以及NVML语言绑定）都会随着每个新CUDA版本发布而更新，并提供大致相同的功能。
更多信息请参阅https://developer.nvidia.com/nvidia-management-library-nvml。
17.3. 集群管理工具
管理您的GPU集群将有助于实现GPU的最大利用率，并帮助您和用户获得最佳性能。许多行业最受欢迎的集群管理工具通过NVML支持CUDA GPU。有关这些工具的列表，请参阅https://developer.nvidia.com/cluster-management。
17.4. 编译器JIT缓存管理工具
应用程序在运行时加载的任何PTX设备代码都会由设备驱动程序进一步编译为二进制代码。这被称为即时编译(JIT)。即时编译会增加应用程序的加载时间，但能让应用程序受益于最新的编译器改进。对于在应用程序编译时尚未存在的设备来说，这也是应用程序能够运行的唯一方式。
当使用PTX设备代码的即时编译时，NVIDIA驱动程序会将生成的二进制代码缓存到磁盘上。此行为的某些方面（如缓存位置和最大缓存大小）可以通过环境变量进行控制；详见《CUDA C++编程指南》中的即时编译章节。
17.5. CUDA_VISIBLE_DEVICES
可以通过设置CUDA_VISIBLE_DEVICES环境变量，在CUDA应用程序启动前重新排列已安装的CUDA设备集合，这些设备将对应用程序可见并被枚举。
要使应用程序可见的设备，应以系统范围内可枚举设备列表的逗号分隔形式包含在内。例如，若仅使用系统设备列表中的设备0和2，应在启动应用程序前设置CUDA_VISIBLE_DEVICES=0,2。随后应用程序将分别将这些设备枚举为设备0和设备1。
18. 建议与最佳实践
本章包含本文档中解释的优化建议摘要。
18.1. 整体性能优化策略
性能优化围绕三个基本策略展开：
最大化并行执行
优化内存使用以实现最大内存带宽
优化指令使用以实现最大指令吞吐量
最大化并行执行的第一步是以能够暴露尽可能多并行性的方式构建算法。一旦算法的并行性被暴露出来，就需要尽可能高效地将其映射到硬件上。这通过仔细选择每个内核启动的执行配置来实现。应用程序还应在更高层次上最大化并行执行，通过流显式暴露设备上的并发执行，并最大化主机与设备之间的并发执行。
优化内存使用首先要尽量减少主机与设备之间的数据传输，因为这些传输的带宽远低于设备内部的数据传输。还应通过最大化利用设备上的共享内存来减少内核对全局内存的访问。有时，最佳优化方案甚至可能是在一开始就避免任何数据传输，只需在需要时重新计算数据即可。
有效带宽会根据每种内存类型的访问模式而有数量级的变化。因此，优化内存使用的下一步是根据最佳内存访问模式来组织内存访问。这种优化对于全局内存访问尤为重要，因为访问延迟会耗费数百个时钟周期。相比之下，共享内存访问通常只有在存在高度存储体冲突时才值得优化。
关于优化指令使用，应避免使用吞吐量较低的算术指令。这意味着在不影响最终结果的情况下，可以牺牲精度来换取速度，例如使用内部函数替代常规函数，或采用单精度而非双精度计算。最后，由于设备的SIMT（单指令多线程）特性，必须特别注意控制流指令。
19. nvcc编译器选项
19.1. nvcc
NVIDIA的nvcc编译器驱动程序将.cu文件转换为适用于主机系统的C++代码以及适用于设备的CUDA汇编或二进制指令。它支持许多命令行参数，其中以下参数对于优化及相关最佳实践特别有用：
-maxrregcount=N指定内核在文件级别可以使用的最大寄存器数量。参见寄存器压力。（另请参阅CUDA C++编程指南执行配置章节中讨论的__launch_bounds__限定符，用于控制每个内核使用的寄存器数量。）--ptxas-options=-v或-Xptxas=-v会列出每个内核的寄存器、共享内存和常量内存使用情况。-ftz=true(非规格化数字会被清零)-prec-div=false(除法精度较低)-prec-sqrt=false(平方根计算精度较低)-use_fast_math编译器选项会将每个functionName()调用强制转换为等效的__functionName()调用。这会以降低精度和准确性为代价使代码运行得更快。详见 Math Libraries。
20. 通知
20.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对客户就本文所述产品的全部及累计责任应受产品销售条款的限制。
20.2. OpenCL
OpenCL是苹果公司的商标，经Khronos Group Inc.授权使用。
20.3. 商标
NVIDIA和NVIDIA标识是美国及其他国家NVIDIA公司的商标或注册商标。其他公司及产品名称可能是其各自关联公司的商标。