CUDA C++ 最佳实践指南

CUDA C++ 最佳实践指南 (nvidia.com)https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html#getting-started2. 异构计算

CUDA 编程涉及在两个不同的平台上同时运行代码：具有一个或多个 CPU 的主机系统和一个或多个支持 CUDA 的 NVIDIA GPU 设备。

虽然 NVIDIA GPU 经常与图形相关联，但它们也是强大的算术引擎，能够并行运行数千个轻量级线程。此功能使它们非常适合可以利用并行执行的计算。

但是，该设备基于与主机系统截然不同的设计，因此了解这些差异以及它们如何确定 CUDA 应用程序的性能以便有效地使用 CUDA 非常重要。

2.1. 主机和设备的区别

主要区别在于线程模型和单独的物理内存：

线程处理资源

主机系统上的执行管道可以支持有限数量的并发线程。例如，具有两个 32 核处理器的服务器只能同时运行 64 个线程（如果 CPU 支持同时多线程处理，则该线程的倍数较小）。相比之下，CUDA 设备上的最小可执行并行单元由 32 个线程组成（称为线程翘曲）。现代 NVIDIA GPU 每个多处理器可以同时支持多达 2048 个活动线程（请参阅 CUDA C++ 编程指南的功能和规格）在具有 80 个多处理器的 GPU 上，这会导致超过 160,000 个并发活动线程。

线程

CPU 上的线程通常是重量级的实体。操作系统必须在 CPU 执行通道上和下交换线程，以提供多线程功能。因此，上下文切换（当交换两个线程时）既慢又昂贵。相比之下，GPU 上的线程非常轻量级。在典型的系统中，有数千个线程排队等待工作（每个线程有 32 个线程的交替）。如果 GPU 必须等待一个线程窖窖，它就会开始在另一个线程窖上执行工作。由于将单独的寄存器分配给所有活动线程，因此在 GPU 线程之间切换时无需交换寄存器或其他状态。资源将一直分配给每个线程，直到它完成执行。简而言之，CPU 内核旨在最大程度地减少一次少量线程的延迟，而 GPU 旨在处理大量并发的轻量级线程，以最大限度地提高吞吐量。

公羊

主机系统和设备都有自己独特的附加物理内存1.由于主机内存和设备内存是分开的，因此主机内存中的项目有时必须在设备内存和主机内存之间进行通信，如在启用 CUDA 的设备上运行的内容中所述。

这些是 CPU 主机和 GPU 设备之间在并行编程方面的主要硬件差异。本文档其他部分将讨论其他差异。考虑到这些差异而组成的应用程序可以将主机和设备一起视为一个有凝聚力的异构系统，其中每个处理单元都被用于执行其最擅长的工作类型：主机上的顺序工作和设备上的并行工作。

2.2. 什么在启用了 CUDA 的设备上运行？

在确定要在设备上运行的应用程序的哪些部分时，应考虑以下问题：

该设备非常适合可以同时在多个数据元素上并行运行的计算。这通常涉及对大型数据集（如矩阵）的算术运算，在这些数据集中，可以同时对数千甚至数百万个元素执行相同的操作。这是在 CUDA 上实现良好性能的要求：软件必须使用大量（通常是数千或数万）并发线程。对并行运行多个线程的支持源于 CUDA 对上述轻量级线程模型的使用。
要使用 CUDA，必须将数据值从主机传输到设备。这些转移在性能方面代价高昂，应尽量减少这些转移。（请参阅主机和设备之间的数据传输。此成本有几个后果：
- 操作的复杂性应该证明将数据移入和移出设备的成本是合理的。传输数据供少数线程短暂使用的代码将看到很少或根本没有性能优势。理想的情况是许多线程执行大量工作。
  
  例如，将两个矩阵传输到设备以执行矩阵加法，然后将结果传输回主机，不会实现太大的性能优势。这里的问题是每个传输的数据元素执行的操作数。对于上述过程，假设矩阵的大小为 NxN，则有 N2操作（添加）和 3N2转移的元素，因此操作与转移的元素的比率为 1：3 或 O（1）。当这个比率较高时，可以更容易地实现性能优势。例如，相同矩阵的矩阵乘法需要 N3运算（乘加），因此运算与传输元素的比率为 O（N），在这种情况下，矩阵越大，性能优势越大。运算类型是一个额外的因素，因为加法与三角函数等具有不同的复杂度。在确定是在主机上还是在设备上执行操作时，将向设备传输数据以及从设备传输数据的开销包括在内，这一点很重要。
- 数据应尽可能长时间地保留在设备上。由于应尽量减少传输，因此在相同数据上运行多个内核的程序应倾向于在内核调用之间将数据留在设备上，而不是将中间结果传输到主机，然后再将它们发送回设备进行后续计算。因此，在前面的示例中，如果由于先前的一些计算，要相加的两个矩阵已经存在于设备上，或者如果相加的结果将用于后续的计算，则应在设备上本地执行矩阵相加。即使可以在主机上更快地执行一系列计算中的某个步骤，也应使用此方法。如果一个相对较慢的内核避免了主机和设备内存之间的一次或多次传输，那么它也可能是有利的。主机和设备之间的数据传输提供了更多详细信息，包括主机和设备之间的带宽测量值以及设备内部的带宽测量值。
为了获得最佳性能，设备上运行的相邻线程在内存访问方面应具有一定的一致性。某些内存访问模式使硬件能够将多个数据项的读取或写入组合并到一个操作中。无法进行布局以实现合并的数据，或者没有足够的位置性来有效使用 L1 或纹理缓存的数据，在用于 GPU 上的计算时，速度往往会降低。一个值得注意的例外是完全随机的内存访问模式。通常，应避免使用它们，因为与峰值功能相比，任何架构都以较低的效率处理这些内存访问模式。但是，与基于缓存的架构（如 CPU）相比，延迟隐藏架构（如 GPU）往往能更好地应对完全随机的内存访问模式。

在具有集成 GPU 的片上系统（例如 NVIDIA® Tegra®）上，主机和设备内存在物理上相同，但主机和设备内存之间仍然存在逻辑上的区别。有关详细信息，请参阅有关 CUDA for Tegra 的应用说明。

3. 应用程序分析

3.1. 配置文件

许多代码使用相对较少的代码量完成了大部分工作。使用分析器，开发人员可以识别此类热点，并开始编译并行化的候选列表。

3.1.1. 创建配置文件

有许多可能的方法来分析代码，但在所有情况下，目标都是相同的：确定应用程序花费大部分执行时间的一个或多个函数。

任何分析活动最重要的考虑因素是确保工作负载是现实的，即从测试中获得的信息和基于该信息的决策与真实数据相关。使用不切实际的工作负载可能会导致次优结果和浪费精力，这会导致开发人员针对不切实际的问题大小进行优化，并导致开发人员专注于错误的功能。

有许多工具可用于生成配置文件。以下示例基于，它是 GNU Binutils 集合中 Linux 平台的开源分析器。gprof

$ 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()

值得注意的是，上述示例中的其他几个函数也占用了总运行时间的很大一部分，例如和。并行化这些功能也应该会增加我们的加速潜力。但是，由于 APOD 是一个循环过程，因此我们可能会选择在后续的 APOD 传递中并行化这些函数，从而将我们在任何给定传递中的工作范围限制为一组较小的增量更改。calcStats()calcSummaryData()

4. 并行化您的应用程序

在确定了热点并完成了设定目标和期望的基本练习之后，开发人员需要并行化代码。根据原始代码，这可以像调用现有的 GPU 优化库（如、、或）一样简单，也可以像添加一些预处理器指令作为并行编译器的提示一样简单。cuBLAScuFFTThrust

另一方面，某些应用程序的设计将需要进行一定程度的重构，以揭示其固有的并行性。由于即使是 CPU 架构也需要公开这种并行性，以便改进或简单地保持顺序应用程序的性能，因此 CUDA 并行编程语言系列（CUDA C++、CUDA Fortran 等）旨在使这种并行性的表达尽可能简单，同时支持在支持 CUDA 的 GPU 上运行，以实现最大的并行吞吐量。

5. 入门

并行化顺序代码有几种关键策略。虽然如何将这些策略应用于特定应用程序的细节是一个复杂且特定于问题的主题，但无论我们是并行化代码以运行多核 CPU 还是在 CUDA GPU 上使用，此处列出的一般主题都适用。

5.1. 并行库

并行化应用程序的最直接方法是利用现有的库，这些库代表我们利用并行架构。CUDA 工具包包含许多已针对 NVIDIA CUDA GPU 进行微调的此类库，例如、等。cuBLAScuFFT

这里的关键是，当库与应用程序的需求很好地匹配时，它们最有用。例如，已经使用其他 BLAS 库的应用程序通常可以很容易地切换到，而几乎不做线性代数的应用程序对几乎没有用处。其他 CUDA 工具包库也是如此：具有类似于等的接口。cuBLAScuBLAScuFFTFFTW

另外值得注意的是 Thrust 库，它是一个并行的 C++ 模板库，类似于 C++ 标准模板库。Thrust 提供了丰富的并行基元集合，例如 scan、sort 和 reduce，这些基元可以组合在一起以实现具有简洁、可读源代码的复杂算法。通过用这些高级抽象来描述你的计算，你为Thrust提供了自动选择最有效实现的自由。因此，Thrust 可用于 CUDA 应用程序的快速原型设计，其中程序员的生产力最为重要，以及生产中，鲁棒性和绝对性能至关重要。

5.2. 并行化编译器

顺序代码并行化的另一种常用方法是使用并行编译器。这通常意味着使用基于指令的方法，其中程序员使用 pragma 或其他类似的表示法向编译器提供提示，说明可以在何处找到并行性，而无需修改或调整底层代码本身。通过向编译器公开并行性，指令允许编译器执行将计算映射到并行架构的详细工作。

OpenACC 标准提供了一组编译器指令，用于指定标准 C、C++ 和 Fortran 中的代码循环和区域，这些代码应从主机 CPU 卸载到附加的加速器（如 CUDA GPU）。管理加速器设备的详细信息由启用了 OpenACC 的编译器和运行时隐式处理。

有关详细信息，请参见 Homepage | OpenACC。

5.3. 编码以暴露并行性

对于需要超出现有并行库或并行编译器所能提供的额外功能或性能的应用程序，与现有顺序代码无缝集成的并行编程语言（如 CUDA C++）是必不可少的。

一旦我们在应用程序的配置文件评估中找到了一个热点，并确定自定义代码是最佳方法，我们就可以使用 CUDA C++ 将该部分代码中的并行性作为 CUDA 内核公开。然后，我们可以在 GPU 上启动此内核并检索结果，而无需对应用程序的其余部分进行重大重写。

当应用程序的大部分总运行时间都花在代码的几个相对孤立的部分时，这种方法最为简单。更难并行化的是具有非常扁平配置文件的应用程序 - 即，花费的时间相对均匀地分布在代码库的广泛部分的应用程序。对于后一种类型的应用程序，可能需要进行一定程度的代码重构以揭示应用程序中固有的并行性，但请记住，这种重构工作将倾向于使所有未来的架构（CPU 和 GPU）受益，因此，如果有必要，付出努力是值得的。

6. 获得正确答案

获得正确的答案显然是所有计算的主要目标。在并行系统上，可能会遇到传统串行编程中通常不会遇到的困难。这些问题包括线程问题、由于浮点值的计算方式导致的意外值，以及 CPU 和 GPU 处理器运行方式差异带来的挑战。本章探讨可能影响返回数据正确性的问题，并指出适当的解决方案。

6.1. 验证

6.1.1. 引用比较

对任何现有程序的修改进行正确性验证的一个关键方面是建立某种机制，通过该机制，可以将来自代表性输入的先前已知良好的参考输出与新结果进行比较。每次更改后，请确保使用适用于特定算法的任何条件的结果都匹配。有些人会期望按位相同的结果，这并不总是可能的，尤其是在涉及浮点运算时;有关数值精度，请参阅数值精度和精度。对于其他算法，如果实现与某个小 epsilon 中的引用匹配，则可能认为它们是正确的。

请注意，用于验证数值结果的过程也可以很容易地扩展为验证性能结果。我们希望确保我们所做的每一项更改都是正确的，并且它能提高性能（以及提高多少）。作为我们周期性 APOD 流程的一个组成部分，经常检查这些事情将有助于确保我们尽快实现预期的结果。

6.1.2. 单元测试

与上述引用比较相对应的一个有用方法是，以一种在单元级别易于验证的方式构建代码本身。例如，我们可以将 CUDA 内核编写为许多短函数的集合，而不是一个大型单体函数;在将它们全部连接在一起之前，可以独立测试每个设备的功能。__device____global__

例如，除了实际计算之外，许多内核还具有用于访问内存的复杂寻址逻辑。如果我们在引入大量计算之前单独验证寻址逻辑，那么这将简化任何后续的调试工作。（请注意，CUDA 编译器将任何不有助于写入全局内存的设备代码视为需要消除的死代码，因此我们至少必须根据寻址逻辑向全局内存写入一些东西，以便成功应用此策略。

更进一步，如果大多数函数被定义为而不仅仅是函数，那么这些函数可以在 CPU 和 GPU 上进行测试，从而增加我们对函数正确并且结果不会有任何意外差异的信心。如果存在差异，那么这些差异将及早看到，并且可以在简单函数的上下文中理解。__host__ __device____device__

作为一个有用的副作用，如果我们希望在应用程序中包含 CPU 和 GPU 执行路径，这种策略将允许我们减少代码重复：如果我们的 CUDA 内核的大部分工作是在函数中完成的，我们可以轻松地从主机代码和设备代码中调用这些函数，而无需重复。__host__ __device__

6.2. 调试

CUDA-GDB 是在 Linux 和 Mac 上运行的 GNU 调试器的一个端口;参见：CUDA-GDB | NVIDIA Developer。

适用于 Microsoft Windows 7、Windows HPC Server 2008、Windows 8.1 和 Windows 10 的 NVIDIA Nsight Visual Studio Edition 可作为 Microsoft Visual Studio 的免费插件使用;参见：https://developer.nvidia.com/nsight-visual-studio-edition。

一些第三方调试器也支持 CUDA 调试;有关详细信息，请参阅：Debugging Solutions | NVIDIA Developer。

6.3. 数值准确度和精度

错误或意外的结果主要是由于浮点值的计算和存储方式导致的浮点精度问题。以下各节介绍了感兴趣的主要项目。《CUDA C++编程指南的功能和技术规格》以及 Log in | NVIDIA Developer 提供的有关浮点精度和性能的白皮书和随附的网络研讨会中介绍了浮点运算的其他特性。

6.3.1. 单精度与双精度

计算能力为 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 位）。这是通过 x86 汇编指令或等效的操作系统 API 完成的。FLDCW

7. 优化 CUDA 应用程序

在每一轮应用程序并行化完成后，开发人员可以开始优化实现以提高性能。由于可以考虑许多可能的优化，因此充分理解应用程序的需求有助于使过程尽可能顺利。然而，与整个 APOD 一样，程序优化是一个迭代过程（确定优化机会，应用和测试优化，验证实现的加速，并重复），这意味着程序员没有必要花费大量时间记住所有可能的优化策略的大部分，然后才能看到良好的加速。取而代之的是，策略可以在学习过程中逐步应用。

优化可以应用于各个级别，从重叠的数据传输与计算一直到微调浮点运算序列。可用的分析工具对于指导此过程非常宝贵，因为它们可以帮助为开发人员的优化工作建议下一个最佳行动方案，并为本指南的优化部分的相关部分提供参考。

8. 性能指标

在尝试优化 CUDA 代码时，了解如何准确测量性能并了解带宽在性能测量中的作用是值得的。本章讨论如何使用 CPU 计时器和 CUDA 事件正确测量性能。然后，它探讨了带宽如何影响性能指标，以及如何缓解它带来的一些挑战。

8.1. 时序

可以使用 CPU 或 GPU 计时器对 CUDA 调用和内核执行进行计时。本节将介绍这两种方法的功能、优点和缺陷。

8.1.1. 使用 CPU 定时器

任何 CPU 计时器都可用于测量 CUDA 调用或内核执行的运行时间。各种 CPU 计时方法的详细信息超出了本文档的范围，但开发人员应始终了解其计时调用提供的分辨率。

使用 CPU 计时器时，请务必记住，许多 CUDA API 函数是异步的;也就是说，在完成工作之前，他们将控制权交还给调用的 CPU 线程。所有内核启动都是异步的，名称上有后缀的内存复制函数也是如此。因此，为了准确测量特定调用或 CUDA 调用序列的运行时间，有必要通过在启动和停止 CPU 计时器之前立即调用来将 CPU 线程与 GPU 同步。阻止调用 CPU 线程，直到该线程之前发出的所有 CUDA 调用都完成。AsynccudaDeviceSynchronize()cudaDeviceSynchronize()

虽然也可以将 CPU 线程与 GPU 上的特定流或事件同步，但这些同步函数不适用于默认流以外的流中的计时代码。阻塞 CPU 线程，直到之前向给定流发出的所有 CUDA 调用都已完成。阻止，直到 GPU 记录特定流中的给定事件。由于驱动程序可能会交错执行来自其他非默认流的 CUDA 调用，因此其他流中的调用可能会包含在计时中。cudaStreamSynchronize()cudaEventSynchronize()

由于默认流（流 0）在设备上工作时表现出序列化行为（默认流中的操作只有在任何流中的所有先前调用完成后才能开始;并且在完成之前，任何流中的后续操作都不能开始）），因此可以可靠地使用这些函数在默认流中计时。

请注意，CPU 到 GPU 的同步点（如本节中提到的同步点）意味着 GPU 的处理管道中会出现停滞，因此应谨慎使用，以最大程度地减少其性能影响。

8.1.2. 使用 CUDA GPU 计时器

CUDA 事件 API 提供用于创建和销毁事件、记录事件（包括时间戳）以及将时间戳差异转换为以毫秒为单位的浮点值的调用。如何使用 CUDA 事件对 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 );

这里用于将 and 事件放入默认流 0 中。当设备到达流中的该事件时，它将记录该事件的时间戳。该函数返回记录和事件之间经过的时间。该值以毫秒为单位，分辨率约为半微秒。与此列表中的其他调用一样，它们的特定操作、参数和返回值在 CUDA 工具包参考手册中进行了描述。请注意，时序是在 GPU 时钟上测量的，因此时序分辨率与操作系统无关。cudaEventRecord()startstopcudaEventElapsedTime()startstop

8.2. 带宽

带宽 - 数据传输的速率 - 是影响性能的最重要门控因素之一。几乎所有对代码的更改都应该在它们如何影响带宽的背景下进行。如本指南的内存优化中所述，带宽可能会受到数据存储内存选择、数据布局方式和访问顺序以及其他因素的显著影响。

为了准确测量性能，计算理论带宽和有效带宽非常有用。当后者远低于前者时，设计或实现细节可能会降低带宽，增加带宽应该是后续优化工作的主要目标。

8.2.1. 理论带宽计算

理论带宽可以使用产品资料中提供的硬件规格来计算。例如，NVIDIA Tesla V100 使用 HBM2（双倍数据速率）RAM，内存时钟频率为 877 MHz 和 4096 位宽的内存接口。

使用这些数据项，NVIDIA Tesla V100 的峰值理论内存带宽为 898 GB/s：

$(0.877\times 10^{9}\times(4096/8)\times2)\div 10^{9}=898GB/s$

在此计算中，内存时钟速率转换为 Hz，乘以接口宽度（除以 8，将位转换为字节），然后乘以 2，得到双倍数据速率。最后，将此产品除以 109将结果转换为 GB/s。

8.2.2. 有效带宽计算

有效带宽是通过对特定程序活动进行计时和了解程序如何访问数据来计算的。为此，请使用以下等式：

$Effective\setminus bandwidth=((B_{r}+B_{w})\div10^{9})\div time$

这里，有效带宽以GB/s为单位，B。r是每个内核读取的字节数，Bw是每个内核写入的字节数，时间以秒为单位。

例如，要计算 2048 x 2048 矩阵副本的有效带宽，可以使用以下公式：

$Effective\setminus bandwidth=((2048^{2}\times 4 \times 2)\div 10^{9})\div time$

元素的数量乘以每个元素的大小（浮点数为 4 个字节），再乘以 2（由于读取和写入），再除以 109（或 1,0243）以获取传输的内存 GB。该数字除以时间（以秒为单位），得到 GB/s。

8.2.3. Visual Profiler报告的吞吐量

对于计算能力为 2.0 或更高的设备，Visual Profiler 可用于收集多个不同的内存吞吐量度量值。可以在“详细信息”或“详细信息图表”视图中显示以下吞吐量指标：

请求的全局负载吞吐量
请求的全球存储吞吐量
全局负载吞吐量
全球存储吞吐量
DRAM 读取吞吐量
DRAM 写入吞吐量

“请求的全局负载吞吐量”和“请求的全局存储吞吐量”值表示内核请求的全局内存吞吐量，因此对应于“有效带宽计算”下显示的计算获得的有效带宽。

由于最小内存事务大小大于大多数字大小，因此内核所需的实际内存吞吐量可能包括内核未使用的数据传输。对于全局内存访问，此实际吞吐量由“全局负载吞吐量”和“全局存储吞吐量”值报告。

需要注意的是，这两个数字都很有用。实际内存吞吐量显示了代码与硬件限制的接近程度，将有效带宽或请求的带宽与实际带宽进行比较，可以很好地估计内存访问的次优合并浪费了多少带宽（请参阅对全局内存的合并访问）。对于全局内存访问，请求的内存带宽与实际内存带宽的比较由“全局内存负载效率”和“全局内存存储效率”指标报告。

作为一个例外，对 HBM2 的分散写入会从 ECC 中看到一些开销，但远低于在受 ECC 保护的 GDDR5 内存上具有类似访问模式的开销。

9. 内存优化

内存优化是提高性能的最重要方面。目标是通过最大化带宽来最大限度地利用硬件。使用尽可能多的快速内存和尽可能少的慢速访问内存可以最好地提供带宽。本章讨论主机和设备上的各种内存，以及如何最好地设置数据项以有效地使用内存。

9.1. 主机和设备之间的数据传输

设备内存和 GPU 之间的峰值理论带宽（例如，NVIDIA Tesla V100 上的 898 GB/s）远高于主机内存和设备内存之间的峰值理论带宽（PCIe x16 Gen3 上的 16 GB/s）。因此，为了获得最佳的整体应用程序性能，必须尽量减少主机和设备之间的数据传输，即使这意味着在 GPU 上运行的内核与在主机 CPU 上运行内核相比没有表现出任何加速。

中间数据结构应在设备内存中创建，由设备操作，并在不被主机映射或复制到主机内存的情况下销毁。

此外，由于每次传输都会产生开销，因此将许多小传输批量转换为一个较大的传输比单独进行每次传输的性能要好得多，即使这样做需要将不连续的内存区域打包到连续的缓冲区中，然后在传输后解压缩。

最后，当使用页面锁定（或固定）内存时，主机和设备之间可以实现更高的带宽，如本文档的 CUDA C++ 编程指南和固定内存部分所述。

9.1.1. 固定内存

页面锁定或固定的内存传输可在主机和设备之间获得最高带宽。例如，在 PCIe x16 Gen3 卡上，固定内存可以达到大约 12 GB/s 的传输速率。

固定内存是使用运行时 API 中的函数分配的。CUDA 示例演示如何使用这些函数以及如何测量内存传输性能。cudaHostAlloc()bandwidthTest

对于已预先分配的系统内存区域，可用于实时固定内存，而无需分配单独的缓冲区并将数据复制到其中。cudaHostRegister()

不应过度使用固定内存。过度使用会降低整体系统性能，因为固定内存是一种稀缺资源，但多少是太多很难提前知道。此外，与大多数正常的系统内存分配相比，系统内存的固定是一项重量级的操作，因此，与所有优化一样，测试应用程序及其运行的系统以获得最佳性能参数。

9.1.2. 异步和重叠的计算传输

主机和设备之间的数据传输正在阻止传输;也就是说，只有在数据传输完成后，控制权才会返回给主机线程。该函数是一种非阻塞变体，其中控制权立即返回给主机线程。与相比，异步传输版本需要固定主机内存（请参阅固定内存），并且它包含一个附加参数，即流 ID。流只是在设备上按顺序执行的一系列操作。不同流中的操作可以交错，在某些情况下可以重叠 - 该属性可用于隐藏主机和设备之间的数据传输。cudaMemcpy()cudaMemcpyAsync()cudaMemcpy()cudaMemcpy()

异步传输以两种不同的方式实现数据传输与计算的重叠。在所有启用了 CUDA 的设备上，都可以将主机计算与异步数据传输和设备计算重叠。例如，重叠计算和数据传输演示了在将数据传输到设备并执行使用该设备的内核时，如何在例程中执行主机计算。cpuFunction()

重叠的计算和数据传输

cudaMemcpyAsync(a_d, a_h, size, cudaMemcpyHostToDevice, 0);
kernel<<<grid, block>>>(a_d);
cpuFunction();

该函数的最后一个参数是流 ID，在本例中，它使用默认流 stream 0。内核也使用默认流，在内存复制完成之前不会开始执行;因此，不需要显式同步。由于内存副本和内核都立即将控制权返回给主机，因此主机函数的执行会重叠。cudaMemcpyAsync()cpuFunction()

在重叠计算和数据传输中，内存复制和内核执行是按顺序进行的。在能够并发复制和计算的设备上，可以将设备上的内核执行与主机和设备之间的数据传输重叠。设备是否具有此功能由结构字段指示（或在 CUDA 示例的输出中列出）。在具有此功能的设备上，重叠再次需要固定主机内存，此外，数据传输和内核必须使用不同的非默认流（具有非零流 ID 的流）。此重叠需要非默认流，因为内存复制、内存集函数和使用默认流的内核调用仅在设备（在任何流中）上的所有先前调用完成后才开始，并且在设备（在任何流中）上的任何操作都未开始，直到它们完成。asyncEngineCountcudaDevicePropdeviceQuery

并发复制和执行说明了基本技术。

并发复制和执行

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 可以被 .由于流中的执行是按顺序进行的，因此在各自流中的数据传输完成之前，不会启动任何内核。当前的 GPU 可以同时处理异步数据传输并执行内核。具有单个复制引擎的 GPU 可以执行一个异步数据传输并执行内核，而具有两个复制引擎的 GPU 可以同时执行一个从主机到设备的异步数据传输，一个从设备到主机的异步数据传输，并执行内核。GPU 上的复制引擎数量由结构字段给出，该字段也列在 CUDA 示例的输出中。（需要提到的是，阻塞传输不能与异步传输重叠，因为阻塞传输发生在默认流中，所以在所有之前的CUDA调用完成后才会开始。在完成之前，它不会允许任何其他 CUDA 调用开始。图 1 显示了一个图表，描绘了两个代码段的执行时间线，图的下半部分等于 4，表示暂存并发复制和执行。nThreads*nStreamsasyncEngineCountcudaDevicePropdeviceQuerynStreams

对于此示例，假设数据传输时间和内核执行时间相当。在这种情况下，当执行时间（tE）超过传输时间（tT）时，暂存版本的总时间粗略估计为 tE + tT/nStreams，而顺序版本为 tE + tT。如果传输时间超过执行时间，则对总时间的粗略估计为 tT + tE/nStreams。

9.1.3. 零拷贝

零复制是 CUDA 工具包 2.2 版中添加的一项功能。它使 GPU 线程能够直接访问主机内存。为此，它需要映射的固定（不可分页）内存。在集成 GPU（即 CUDA 设备属性结构的集成字段设置为 1 的 GPU）上，映射的固定内存始终是一种性能提升，因为它避免了多余的副本，因为集成 GPU 和 CPU 内存在物理上是相同的。在离散 GPU 上，映射的固定内存仅在某些情况下才具有优势。由于数据未缓存在 GPU 上，因此映射的固定内存应仅读取或写入一次，并且应合并读取和写入内存的全局加载和存储。可以使用零拷贝代替流，因为内核源自的数据传输会自动与内核执行重叠，而无需设置和确定最佳流数的开销。

零拷贝主机代码中的主机代码显示了通常如何设置零拷贝。

零拷贝主机代码

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);

在此代码中，返回的结构体字段用于检查设备是否支持将主机内存映射到设备的地址空间。页锁定内存映射是通过使用调用来启用的。请注意，必须在设置设备或进行需要状态的 CUDA 调用之前调用（即，基本上，在创建上下文之前）。页面锁定的映射主机内存是使用分配的，并且指向映射设备地址空间的指针是通过函数获取的。在代码中的零拷贝主机代码中，可以使用指针引用映射的固定主机内存，就像a_map引用设备内存中的某个位置一样。canMapHostMemorycudaGetDeviceProperties()cudaSetDeviceFlags()cudaDeviceMapHostcudaSetDeviceFlags()cudaHostAlloc()cudaHostGetDevicePointer()kernel()a_map

9.1.4. 统一虚拟寻址

计算能力 2.0 及更高版本的设备在 64 位 Linux 和 Windows 上支持一种称为统一虚拟寻址 （UVA）的特殊寻址模式。使用 UVA 时，所有已安装的受支持设备的主机内存和设备内存共享一个虚拟地址空间。

在 UVA 之前，应用程序必须跟踪哪些指针引用设备内存（以及针对哪个设备），哪些指针引用主机内存作为每个指针的单独元数据位（或作为程序中的硬编码信息）。另一方面，使用 UVA，只需使用检查指针的值，就可以确定指针指向的物理内存空间。cudaPointerGetAttributes()

在 UVA 下，分配的固定主机内存将具有相同的主机和设备指针，因此无需调用此类分配。但是，事后通过固定的主机内存分配将继续具有与其主机指针不同的设备指针，因此在这种情况下仍然是必要的。cudaHostAlloc()cudaHostGetDevicePointer()cudaHostRegister()cudaHostGetDevicePointer()

UVA 也是绕过主机内存，直接通过 PCIe 总线或 NVLink 实现点对点（P2P）数据传输的必要先决条件，适用于受支持配置中的受支持 GPU。

请参阅 CUDA C++ 编程指南，了解 UVA 和 P2P 的进一步说明和软件要求。

9.2. 设备内存空间

CUDA 设备使用多个内存空间，这些内存空间具有不同的特征，反映了它们在 CUDA 应用程序中的不同用法。这些内存空间包括全局内存空间、本地内存空间、共享内存空间、纹理内存空间和寄存器空间，如图 2 所示。

在这些不同的内存空间中，全局内存是最丰富的;请参阅 CUDA C++ 编程指南的功能和技术规格，了解每个计算能力级别的每个内存空间中的可用内存量。全局内存、本地内存和纹理内存的访问延迟最大，其次是常量内存、共享内存和寄存器文件。

表 1 显示了内存类型的各种主要特征。

表 1.设备内存的显著特点
记忆	位置：片上/片外	缓存	访问	范围	辈子
注册	上	不适用	读/写	1 个线程	线
当地	关闭	是的††	读/写	1 个线程	线
共享	上	不适用	读/写	块中的所有线程	块
全球	关闭	†	读/写	所有线程 + 主机	主机分配
不断	关闭	是的	R	所有线程 + 主机	主机分配
质地	关闭	是的	R	所有线程 + 主机	主机分配
†默认情况下，在计算能力为 6.0 和 7.x 的设备上缓存在 L1 和 L2 中;默认情况下，在计算能力较低的设备上仅在 L2 中缓存，但有些设备也允许通过编译标志选择在 L1 中缓存。
††默认情况下缓存在 L1 和 L2 中，计算能力为 5.x 的设备除外;计算能力为 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 时，以合并方式访问内存更为重要。分散访问会增加 ECC 内存传输开销，尤其是在将数据写入全局内存时。

以下简单示例说明了合并概念。除非另有说明，否则这些示例假定计算能力为 6.0 或更高版本，并且访问是针对 4 字节字的。

9.2.1.1. 简单的访问模式

第一种也是最简单的合并情况可以通过任何计算能力为 6.0 或更高版本的启用 CUDA 的设备来实现：第 k 个线程访问 32 字节对齐数组中的第 k 个字。并非所有线程都需要参与。

例如，如果 warp 的线程访问相邻的 4 字节字（例如，相邻的值），则四个合并的 32 字节事务将为该内存访问提供服务。这种模式如图 3 所示。float

此访问模式产生四个 32 字节的事务，由红色矩形表示。

如果从四个 32 字节段中的任何一个段中只请求单词的子集（例如，如果多个线程访问了同一个单词，或者如果某些线程没有参与访问），则无论如何都会获取整个段。此外，如果 warp 线程的访问在四个段内或跨四个段进行置换，则计算能力为 6.0 或更高的设备仍然只能执行四个 32 字节的事务。

9.2.1.2. 顺序但未对齐的访问模式

如果 warp 访问内存中的顺序线程是顺序的，但未与 32 字节段对齐，则将请求五个 32 字节的段，如图 4 所示。

通过 CUDA 运行时 API 分配的内存（例如 via ）保证至少对齐为 256 个字节。因此，选择合理的线程块大小，例如翘曲大小的倍数（即，在当前 GPU 上为 32 倍），有助于通过正确对齐的翘曲访问内存。（例如，考虑如果线程块大小不是 warp 大小的倍数，则第二个、第三个和后续线程块访问的内存地址会发生什么情况。cudaMalloc()

9.2.1.3. 未对齐访问的影响

使用简单的复制内核（例如 A copy kernel 中说明未对齐访问的复制内核）来探索未对齐访问的后果既简单又有益。

一个复制内核，用于说明未对齐的访问

__global__ void offsetCopy(float *odata, float* idata, int offset)
{
    int xid = blockIdx.x * blockDim.x + threadIdx.x + offset;
    odata[xid] = idata[xid];
}

在说明未对齐访问的复制内核中，数据从输入数组复制到输出数组，这两个数组都存在于全局内存中。内核在主机代码的循环中执行，该循环将参数从 0 更改为 32。（例如，图 4 对应于这种错位）图 5 显示了 NVIDIA Tesla V100 （计算能力 7.0）上具有各种偏移量的副本的有效带宽。idataoffset

对于 NVIDIA Tesla V100，没有偏移量或偏移量是 8 个字的倍数的全局内存访问会导致四个 32 字节的事务。实现的带宽约为 790 GB/s。否则，每个 warp 加载 5 个 32 字节的段，我们预计大约有 4/5 个第在不偏移的情况下实现的内存吞吐量。

然而，在此特定示例中，实现的偏移内存吞吐量约为 9/10第，因为相邻的 Warp 会重用其邻居获取的缓存行。因此，虽然影响仍然很明显，但它并不像我们预期的那么大。如果相邻的扭曲没有表现出对超取缓存行的如此高度的重用，情况会更糟。

9.2.1.4. 跨步访问

如上所述，在顺序访问未对齐的情况下，缓存有助于减轻性能影响。但是，它与非单元跨步访问可能有所不同，这是在处理多维数据或矩阵时经常出现的模式。因此，确保实际使用每个缓存行中尽可能多的数据是这些设备上内存访问性能优化的重要部分。

为了说明跨步访问对有效带宽的影响，请参阅 A 内核以说明非单位跨步数据复制中的内核，该内核在线程之间复制具有跨步元素跨步的数据从到。strideCopy()idataodata

用于说明非单位步幅数据复制的内核

__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）上每个扭曲加载 8 个 L2 缓存段。

步长为 2 会导致 50% 的负载/存储效率，因为事务中的一半元素没有被使用，代表浪费的带宽。随着步幅的增加，有效带宽会减小，直到在扭曲中为 32 个线程加载 32 个 32 字节段，如图 7 所示。

如图 7 所示，应尽可能避免非单位步幅全局内存访问。执行此操作的一种方法是利用共享内存，这将在下一节中讨论。

9.2.2. L2缓存

从 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 */

可以使用 CUDA 流或 CUDA 图形内核节点上的访问策略窗口来控制用户数据到 L2 预留部分的映射。以下示例显示了如何在 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);

访问策略窗口需要和的值。根据参数的值和 L2 缓存的大小，可能需要调整的值以避免 L2 缓存行的抖动。hitRationum_bytesnum_byteshitRatio

9.2.2.2. 调整访问窗口命中率

该参数可用于指定接收属性的访问比例。例如，如果值为 0.6，则 60% 的内存访问全局内存区域 [ptr..PTR+num_bytes）具有 Persisting 属性，40% 的内存访问具有 Streaming 属性。为了理解 and 的效果，我们使用了一个滑动窗口微基准。hitRatiohitProphitRatiohitRationum_bytes

此微基准测试在 GPU 全局内存中使用 1024 MB 区域。首先，我们留出 30 MB 的 L2 缓存用于使用的持久访问，如上所述。然后，如下图所示，我们指定对内存区域的第一个字节的访问是持久的。因此，此数据将使用 L2 预留部分。在我们的实验中，我们将此持久性数据区域的大小从 10 MB 更改为 60 MB，以模拟数据适合或超过 30 MB 的可用 L2 预留部分的各种场景。请注意，NVIDIA Tesla A100 GPU 的总 L2 缓存容量为 40 MB。对内存区域的剩余数据（即流数据）的访问被视为正常访问或流式访问，因此将使用剩余的 10 MB 非预留 L2 部分（除非 L2 预留部分的一部分未使用）。cudaDeviceSetLimit()freqSize * sizeof(int)

考虑以下内核代码和访问窗口参数，作为滑动窗口实验的实现。

__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 缓存的 30 MB 预留部分时，可以观察到高达 50% 的性能提升。但是，一旦此持久性数据区域的大小超过 L2 预留缓存部分的大小，由于 L2 缓存行的抖动，性能会下降约 10%。

为了优化性能，当持久化数据的大小大于预留的 L2 缓存部分的大小时，我们在访问窗口中调整 and 参数，如下所示。num_byteshitRatio

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.

我们将访问窗口中的 20 MB 固定为 20 MB，并调整该值，使得总持久性数据的随机 20 MB 驻留在 L2 预留缓存部分。此持久性数据的其余部分将使用 streaming 属性进行访问。这有助于减少缓存抖动。结果如下图所示，无论持久性数据是否适合 L2 预留，我们都能看到良好的性能。num_byteshitRatio

9.2.3. 共享内存

因为它是片上的，所以共享内存比本地和全局内存具有更高的带宽和更低的延迟 - 前提是线程之间没有银行冲突，如下一节所述。

9.2.3.1. 共享内存和内存库

为了实现并发访问的高内存带宽，共享内存被划分为可以同时访问的相同大小的内存模块（bank）。因此，跨 n 个不同内存组的任何内存负载或 n 个地址的存储都可以同时提供服务，从而产生比单个内存组的带宽高 n 倍的有效带宽。

但是，如果内存请求的多个地址映射到同一内存组，则访问将被序列化。硬件根据需要将具有银行冲突的内存请求拆分为尽可能多的单独无冲突请求，从而将有效带宽减少等于单独内存请求数的系数。这里的一个例外是，当 warp 中的多个线程处理相同的共享内存位置时，从而导致广播。在这种情况下，来自不同库的多个广播从请求的共享内存位置合并为一个组播到线程。

为了最大程度地减少内存库冲突，了解内存地址如何映射到内存库以及如何以最佳方式安排内存请求非常重要。

在计算能力为 5.x 或更高版本的设备上，每个时钟周期的 bank 带宽为 32 位，并且连续的 32 位字分配给连续的 bank。翘曲大小为 32 个线程，库数也是 32，因此翘曲中的任何线程之间都可能发生库冲突。有关详细信息，请参阅 CUDA C++ 编程指南中的计算能力 5.x。

9.2.3.2. 矩阵乘法中的共享内存（C=AB）

共享内存使块中的线程之间能够进行协作。当块中的多个线程使用全局内存中的相同数据时，共享内存只能用于访问全局内存中的数据一次。共享内存还可用于避免未合并的内存访问，方法是从全局内存中以合并模式加载和存储数据，然后在共享内存中对其进行重新排序。除了内存库冲突之外，共享内存中的扭曲不会对非顺序或未对齐的访问进行惩罚。

共享内存的使用通过矩阵乘法 C = AB 的简单示例进行了说明，其中 A 的维度为 Mxw，B 的维度为 wxN，C 的维度为 MxN。为了保持内核的简单性，M 和 N 是 32 的倍数，因为当前设备的翘曲大小（w）为 32。

这个问题的一个自然分解是使用 wxw 线程的块和瓦片大小。因此，就wxw瓦片而言，A是列矩阵，B是行矩阵，C是它们的外积;参见图 11。启动一个由 M/w 块组成的 N/w 网格，其中每个线程块从 A 的单个瓦片和 B 的单个瓦片计算 C 中不同瓦片的元素。

为此，内核（未优化矩阵乘法）计算矩阵 C 瓦片的输出元素。simpleMultiply

未优化的矩阵乘法

__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 的全局内存的指针;、和都等于 w。wxw 线程块中的每个线程都计算 C 瓦片中的一个元素。并且是 C 中元素的行和列，由特定线程计算。循环将 A 的行乘以 B 的一列，然后将其写入 C。abcblockDim.xblockDim.yTILE_DIMrowcolfori

在 NVIDIA Tesla V100 上，此内核的有效带宽为 119.9 GB/s。为了分析性能，有必要考虑扭曲如何在循环中访问全局内存。每个线程经线计算 C 瓦片的一行，该瓦片取决于 A 的单行和 B 的整个瓦片，如图 12 所示。for

对于循环的每次迭代 i，warp 中的线程都会读取 B tiles 的一行，这是所有计算功能的顺序和合并访问。for

但是，对于每次迭代 i，warp 中的所有线程都从矩阵 A 的全局内存中读取相同的值，因为索引在 warp 中是恒定的。尽管在计算能力为 2.0 或更高的设备上，这种访问只需要 1 个事务，但事务中会浪费带宽，因为在 32 字节缓存段的 8 个字中只使用一个 4 字节字。我们可以在循环的后续迭代中重用此缓存行，最终将使用所有 8 个单词;但是，当许多 Warp 同时在同一多处理器上执行时（通常情况下），缓存行很容易在迭代 I 和 I+1 之间从缓存中逐出。row*TILE_DIM+i

通过将 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 瓦片中的每个元素仅以完全合并的方式（不浪费带宽）从全局内存中读取一次到共享内存。在循环的每次迭代中，共享内存中的值会以扭曲的形式广播到所有线程。在将 A 的瓦片读取到共享内存中后，a 就足够了，而不是同步屏障调用，因为只有将数据写入共享内存的 warp 内的线程才会读取此数据。该内核在 NVIDIA Tesla V100 上的有效带宽为 144.4 GB/s。这说明了当硬件 L1 缓存逐出策略与应用程序的需求不匹配时，或者当 L1 缓存不用于从全局内存中读取时，将共享内存用作用户管理的缓存。for__syncthreads()__syncwarp()

在矩阵乘法中，使用共享内存提高全局内存加载效率如何处理矩阵 B，可以进一步改进。在计算矩阵 C 的瓦片的每一行时，将读取 B 的整个瓦片。通过向共享内存中读取一次 B 图块，可以消除对 B tile的重复读取（通过将额外数据读取到共享内存中的改进）。

通过将其他数据读取到共享内存中进行改进

__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 图块后需要调用，因为 warp 从共享内存中读取由不同 warp 写入共享内存的数据。在 NVIDIA Tesla V100 上，此例程的有效带宽为 195.5 GB/s。请注意，性能的提高不是由于在这两种情况下的合并得到改善，而是由于避免了来自全局内存的冗余传输。__syncthreads()

表 2 总结了各种优化的结果。

表 2.优化 C = AB 矩阵乘法的性能改进：class table-no-stripes
优化	英伟达 Tesla V100
无优化	119.9 千兆字节/s
使用共享内存合并以存储 A 的 tile	144.4 千兆字节/s
使用共享内存消除对 B 瓦片的冗余读取	195.5 千兆字节/s

9.2.3.3. 矩阵乘法中的共享内存（C=AAT）

可以使用前一个矩阵乘法的变体来说明如何处理对全局内存的跨步访问以及共享内存库冲突。这种变体只是使用 A 的转置代替 B，因此 C = AAT.

C = AA 的简单实现T显示在对全局内存的跨步访问的未优化处理中

对全局内存的跨步访问的未优化处理

__global__ void simpleMultiply(float *a, float *c, int M)
{
    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] * a[col*TILE_DIM+i];
    }
    c[row*M+col] = sum;
}

在对全局内存跨步访问的未优化处理中，通过取 A 的第 1 行和第 1 行的点积来获得 C 的第 1 行第 1 行第 1 行元素。在 NVIDIA Tesla V100 上，此内核的有效带宽为 12.8 GB/s。这些结果大大低于 C = AB 内核的相应测量值。区别在于，对于每次迭代，半翘曲中的线程如何在第二项中访问 A 的元素。对于线程的翘曲，表示 A 转置的顺序列，因此表示全局内存的跨步访问，步长为 w，导致大量带宽浪费。a[col*TILE_DIM+i]icolcol*TILE_DIM

避免跨步访问的方法是像以前一样使用共享内存，但在本例中，warp 将 A 行读取到共享内存瓦片的一列中，如使用全局内存中的合并读取优化的跨步访问处理中所示。

使用从全局内存中合并读取的跨步访问的优化处理

__global__ void coalescedMultiply(float *a, float *c, int M)
{
    __shared__ float aTile[TILE_DIM][TILE_DIM],
                     transposedTile[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];
    transposedTile[threadIdx.x][threadIdx.y] =
        a[(blockIdx.x*blockDim.x + threadIdx.y)*TILE_DIM +
        threadIdx.x];
    __syncthreads();
    for (int i = 0; i < TILE_DIM; i++) {
        sum += aTile[threadIdx.y][i]* transposedTile[i][threadIdx.x];
    }
    c[row*M+col] = sum;
}

使用来自全局内存的合并读取对跨步访问的优化处理，在点积的第二项中使用共享技术来避免未合并的访问，并使用上一个示例中的共享技术来避免第一项的未合并访问。在 NVIDIA Tesla V100 上，此内核的有效带宽为 140.2 GB/s。这些结果低于 C = AB 时最终内核获得的结果。造成这种差异的原因是共享内存库冲突。transposedTileaTile

for 循环中元素的读取没有冲突，因为每个半扭曲的线程在瓦片的行上读取，导致单元跨过岸的跨步。但是，在将瓦片从全局内存复制到共享内存时，会发生银行冲突。为了能够合并来自全局内存的负载，将按顺序从全局内存中读取数据。但是，这需要在列中写入共享内存，并且由于在共享内存中使用了 wxw tiles，这会导致 w bank 的线程之间出现跨步 - warp 的每个线程都命中同一bank（回想一下，w 被选为 32）。这些多方银行冲突的代价非常高昂。简单的补救措施是填充共享内存数组，使其具有额外的列，如以下代码行所示。transposedTile

__shared__ float transposedTile[TILE_DIM][TILE_DIM+1];

这种填充完全消除了冲突，因为现在线程之间的步幅是 w+1 bank（即，对于当前设备为 33），由于用于计算 bank 索引的模算术，这相当于一个单位步幅。此更改后，NVIDIA Tesla V100 上的有效带宽为 199.4 GB/s，与上一个 C = AB 内核的结果相当。

表 3 总结了这些优化的结果。

这些结果应与表2中的结果进行比较。从这些表中可以看出，明智地使用共享内存可以显著提高性能。

本节中的示例说明了使用共享内存的三个原因：

启用对全局内存的合并访问，特别是为了避免大步长（对于常规矩阵，步长远大于 32）
从全局内存中消除（或减少）冗余负载
为避免浪费带宽

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);
}

内核的同步版本将元素从全局内存加载到中间寄存器，然后将中间寄存器值存储到共享内存。在内核的异步版本中，一旦调用函数，就会发出从全局内存加载并直接存储到共享内存中的指令。将等到管道对象中的所有指令都已执行完毕。使用异步副本不使用任何中间寄存器。不使用中间寄存器有助于减轻寄存器压力，并可以提高内核占用率。使用异步复制指令从全局内存复制到共享内存的数据可以缓存在 L1 缓存中，也可以选择性地绕过 L1 缓存。如果单个 CUDA 线程正在复制 16 个字节的元素，则可以绕过 L1 缓存。这种差异如图 13 所示。__pipeline_memcpy_async()__pipeline_wait_prior(0)

我们使用每个线程大小为 4B、8B 和 16B 的元素来评估两个内核的性能，即使用、和模板参数。我们调整内核中的，使每个线程块从 512 字节复制到 48 MB。内核的性能如图 14 所示。intint2int4copy_count

从性能图表中，可以对此实验进行以下观察。

当参数是所有三个元素大小的 4 的倍数时，可实现同步复制的最佳性能。编译器可以优化 4 个加载和存储指令的组。这从锯齿曲线中可以明显看出。copy_count
异步复制在几乎所有情况下都能实现更好的性能。
async-copy 不要求参数是 4 的倍数，以便通过编译器优化最大限度地提高性能。copy_count
总体而言，使用元素大小为 8 或 16 字节的异步副本时，可实现最佳性能。

本地内存之所以如此命名，是因为它的作用域是线程的本地内存，而不是因为它的物理位置。事实上，本地内存是片外的。因此，访问本地内存与访问全局内存一样昂贵。换言之，名称中的“本地”一词并不意味着访问速度更快。

本地内存仅用于保存自动变量。这是由编译器在确定没有足够的寄存器空间来容纳变量时完成的。可能放置在本地内存中的自动变量是会占用过多寄存器空间的大型结构或数组，以及编译器确定可能动态索引的数组。nvcc

检查 PTX 汇编代码（通过使用或命令行选项进行编译获得）可揭示在第一个编译阶段是否已将变量放置在本地内存中。如果有，它将使用助记符声明，并使用和助记符进行访问。如果没有，如果后续编译阶段发现变量为目标架构占用了太多寄存器空间，则仍可能做出其他决定。无法检查特定变量的此情况，但是当使用该选项运行时，编译器会报告每个内核（lmem）的总本地内存使用情况。-ptx-keepnvcc.localld.localst.local--ptxas-options=-v

9.2.5. 纹理内存

只读纹理内存空间被缓存。因此，纹理提取仅在缓存未命中时读取一个设备内存;否则，只需从纹理缓存中读取一次。纹理缓存针对 2D 空间局部性进行了优化，因此读取彼此靠近的纹理地址的相同扭曲线程将获得最佳性能。纹理内存还设计用于具有恒定延迟的流式提取;也就是说，缓存命中会减少 DRAM 带宽需求，但不会减少提取延迟。

在某些寻址情况下，通过纹理提取读取设备内存可能是从全局或恒定内存中读取设备内存的有利替代方法。

9.2.5.1. 其他纹理功能

如果使用、、或而不是获取纹理，则硬件会提供其他功能，这些功能可能对某些应用程序（如图像处理）有用，如表 4 所示。tex1D()tex2D()tex3D()tex1Dfetch()

表 4.tex1D（）、tex2D（）和 tex3D（）提取的有用功能
特征	用	警告
滤波	纹素之间的快速、低精度插值	仅当纹理引用返回浮点数据时才有效
归一化纹理坐标	与分辨率无关的编码	没有
寻址模式	边界情况的自动处理1	只能与归一化纹理坐标一起使用
1表 4 底部行中边界情况的自动处理是指当纹理坐标超出有效寻址范围时如何解析纹理坐标。有两种选择：夹紧和包裹。如果 x 是坐标，N 是一维纹理的纹素数，则使用 clamp 时，如果 x < 0，则 x 替换为 0，如果 1 <x，则替换为 1-1/N。使用包装时，x 被 frac（x）替换，其中 frac（x） = x - floor（x）。Floor 返回小于或等于 x 的最大整数。因此，在 N = 1 的钳位模式下，1.3 的 x 被钳位为 1.0;而在 wrap 模式下，它被转换为 0.3

在内核调用中，纹理缓存不会与全局内存写入保持一致，因此从同一内核调用中通过全局存储写入的地址提取纹理会返回未定义的数据。也就是说，如果内存位置已被先前的内核调用或内存副本更新过，则线程可以通过纹理安全地读取该位置，但如果该位置之前已被同一线程或同一内核调用中的另一个线程更新过，则则无法读取该位置。

9.2.6. 常数内存

设备上总共有 64 KB 的恒定内存。缓存恒定内存空间。因此，从常量内存中读取仅在缓存未命中时从设备内存中读取一次内存的成本;否则，它只需从常量缓存中读取一次。Warp 内线程对不同地址的访问是序列化的，因此成本与 Warp 内所有线程读取的唯一地址数量成线性比例。因此，当同一 warp 中的线程仅访问几个不同的位置时，常量缓存是最好的。如果 warp 的所有线程都访问同一位置，则恒定内存的速度可以与寄存器访问一样快。

9.2.7. 寄存器

通常，访问寄存器每条指令消耗零额外的时钟周期，但由于寄存器先写后读依赖性和寄存器内存组冲突，可能会发生延迟。

编译器和硬件线程调度器将尽可能优化地调度指令，以避免寄存器内存组冲突。应用程序无法直接控制这些银行冲突。特别是，没有寄存器相关的原因将数据打包到向量数据类型（如 OR 类型）中。float4int4

9.2.7.1. 寄存器压力

当没有足够的寄存器可用于给定任务时，就会发生寄存器压力。尽管每个多处理器都包含数千个 32 位寄存器（请参阅 CUDA C++ 编程指南的功能和技术规格），但这些寄存器在并发线程之间进行分区。为防止编译器分配过多的寄存器，请使用编译器命令行选项（请参阅 nvcc）或启动边界内核定义限定符（请参阅 CUDA C++ 编程指南的执行配置）来控制每个线程分配的最大寄存器数。-maxrregcount=N

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. 执行配置优化

良好性能的关键之一是使设备上的多处理器尽可能繁忙。如果设备在多处理器之间工作不平衡，则性能将欠佳。因此，将应用程序设计为以最大化硬件利用率的方式使用线程和块，并限制阻碍工作自由分发的做法非常重要。这项工作中的一个关键概念是占用率，以下各节将对此进行说明。

在某些情况下，还可以通过设计应用程序来提高硬件利用率，以便可以同时执行多个独立的内核。多个内核同时执行称为并发内核执行。并发内核执行如下所述。

另一个重要概念是管理为特定任务分配的系统资源。本章的最后几节将讨论如何管理此资源利用率。

10.1. 入住率

线程指令在 CUDA 中按顺序执行，因此，当一个 Warp 暂停或停止时执行其他 Warp 是隐藏延迟并保持硬件繁忙的唯一方法。因此，与多处理器上的活动翘曲数量相关的一些指标对于确定硬件保持繁忙的效率非常重要。这个指标是入住率。

占用率是每个多处理器的活动扭曲数与可能的最大活动扭曲数之比。（要确定后一个数字，请参阅 CUDA 示例或参阅 CUDA C++ 编程指南中的计算能力。查看占用率的另一种方法是硬件处理翘曲的能力中正在积极使用的百分比。deviceQuery

更高的入住率并不总是等同于更高的性能——超过一个点，额外的入住率并不能提高性能。但是，低占用率总是会干扰隐藏内存延迟的能力，从而导致性能下降。

CUDA 内核所需的每线程资源可能会以不需要的方式限制最大块大小。为了保持与未来硬件和工具包的向前兼容性，并确保至少有一个线程块可以在 SM 上运行，开发人员应该包含一个参数，该参数指定内核将启动的最大块大小。如果不这样做，可能会导致“启动时请求的资源过多”错误。在某些情况下，提供的两个参数版本可以提高性能。应使用详细的每个内核分析来确定正确的值。__launch_bounds__(maxThreadsPerBlock)__launch_bounds__(maxThreadsPerBlock,minBlocksPerMultiprocessor)minBlocksPerMultiprocessor

10.1.1. 计算占用率

决定入住率的几个因素之一是收银机的可用性。寄存器存储使线程能够将局部变量保持在附近，以实现低延迟访问。但是，寄存器集（称为寄存器文件）是驻留在多处理器上的所有线程必须共享的有限商品。寄存器一次分配给整个模块。因此，如果每个线程块使用许多寄存器，则可以驻留在多处理器上的线程块数量会减少，从而降低多处理器的占用率。每个线程的最大寄存器数可以在编译时手动设置，每个文件使用选项或每个内核使用限定符设置（请参阅寄存器压力）。-maxrregcount__launch_bounds__

为了计算占用率，每个线程使用的寄存器数量是关键因素之一。例如，在计算能力为 7.0 的设备上，每个多处理器都有 65,536 个 32 位寄存器，并且最多可以驻留 2048 个并发线程（64 个扭曲 x 每个扭曲 32 个线程）。这意味着，在其中一个设备中，要使多处理器具有 100% 的占用率，每个线程最多可以使用 32 个寄存器。但是，这种确定寄存器计数如何影响占用率的方法并未考虑寄存器分配粒度。例如，在计算能力为 7.0 的设备上，具有 128 线程块的内核（每个线程使用 37 个寄存器）会导致占用率为 75%，每个多处理器有 12 个活动的 128 线程块，而具有 320 线程块的内核（每个线程使用相同的 37 个寄存器）会导致占用率为 63%，因为只有四个 320 线程块可以驻留在多处理器上。此外，寄存器分配四舍五入到每个经线最接近的 256 个寄存器。

可用的寄存器数量、每个多处理器上驻留的最大并发线程数以及寄存器分配粒度因计算能力的不同而变化。由于寄存器分配中的这些细微差别，以及多处理器的共享内存也在常驻线程块之间分配的事实，因此很难确定寄存器使用和占用率之间的确切关系。该选项详细说明了每个内核的每个线程使用的寄存器数量。请参阅 CUDA C++ 编程指南的硬件多线程，了解各种计算能力设备的寄存器分配公式，以及 CUDA C++ 编程指南的特性和技术规格，了解这些设备上可用的寄存器总数。或者，NVIDIA 以 Excel 电子表格的形式提供了一个占用率计算器，使开发人员能够磨练最佳平衡并更轻松地测试不同的可能场景。此电子表格如图 15 所示，名为 CUDA Toolkit 安装的 tools 子目录，位于该子目录中。--ptxas options=vnvccCUDA_Occupancy_Calculator.xls

除了计算器电子表格外，还可以使用 NVIDIA Nsight 计算分析器确定占用率。有关占用的详细信息显示在“占用”部分中。

应用程序还可以使用 CUDA 运行时中的 Occupancy API，例如，根据运行时参数动态选择启动配置。cudaOccupancyMaxActiveBlocksPerMultiprocessor

10.2. 隐藏寄存器依赖关系

当指令使用存储在寄存器中的结果时，就会产生寄存器依赖性，寄存器由指令在其之前写入。在计算能力为 7.0 的设备上，大多数算术指令的延迟通常为 4 个周期。因此，线程必须等待大约 4 个周期才能使用算术结果。但是，这种延迟可以通过在其他 warps 中执行线程来完全隐藏。有关详细信息，请参阅寄存器。

10.3. 线程和块启发式

每个网格的块的尺寸和大小以及每个块的线程的尺寸和大小都是重要因素。这些参数的多维方面允许更容易地将多维问题映射到 CUDA，并且不会对性能产生影响。因此，本部分讨论的是大小，但不讨论尺寸。

延迟、隐藏和占用取决于每个多处理器的活动扭曲数量，而活动扭曲的数量是由执行参数以及资源（寄存器和共享内存）约束隐式确定的。选择执行参数是在延迟隐藏（占用）和资源利用率之间取得平衡的问题。

选择执行配置参数应同时进行;但是，有一些启发式方法可以单独应用于每个参数。在选择第一个执行配置参数（每个网格的块数或网格大小）时，主要关注的是保持整个 GPU 繁忙。网格中的块数应大于多处理器的数量，以便所有多处理器至少有一个块要执行。此外，每个多处理器应该有多个活动块，以便不等待的块可以使硬件保持繁忙。此建议视资源可用性而定;因此，它应该在第二个执行参数的上下文中确定 - 每个块的线程数或块大小 - 以及共享内存使用情况。为了扩展到未来的设备，每次内核启动的块数应为数千个。__syncthreads()

在选择块大小时，重要的是要记住，多个并发块可以驻留在多处理器上，因此占用率不仅仅由块大小决定。特别是，更大的区块大小并不意味着更高的入住率。

如“入住率”中所述，更高的入住率并不总是等同于更好的性能。例如，将入住率从 66% 提高到 100% 通常不会转化为类似的绩效提升。占用率较低的内核将比占用率较高的内核在每个线程上有更多的可用寄存器，这可能会导致较少的寄存器溢出到本地内存;特别是，通过高度公开的指令级并行性（ILP），在某些情况下，可以以低占用率完全覆盖延迟。

在选择块大小时，涉及到许多这样的因素，不可避免地需要进行一些实验。但是，应遵循一些经验法则：

每个块的线程数应该是 warp 大小的倍数，以避免在填充不足的 warp 上浪费计算，并促进合并。
每个块至少应使用 64 个线程，并且仅当每个多处理器有多个并发块时才应使用。
每个块 128 到 256 个线程是试验不同块大小的良好初始范围。
如果延迟影响性能，则使用几个较小的线程块，而不是每个多处理器的一个大线程块。这对于经常调用 .__syncthreads()

请注意，当线程块分配的寄存器多于多处理器上的可用寄存器时，内核启动将失败，就像当请求过多的共享内存或过多线程时一样。

10.4. 共享内存的影响

共享内存在多种情况下可能会有所帮助，例如帮助合并或消除对全局内存的冗余访问。但是，它也可以作为对入住率的限制。在许多情况下，内核所需的共享内存量与所选的块大小有关，但线程到共享内存元素的映射不需要是一对一的。例如，可能需要在内核中使用 64x64 元素的共享内存数组，但由于每个块的最大线程数为 1024，因此无法启动每个块具有 64x64 个线程的内核。在这种情况下，可以启动具有 32x32 或 64x16 线程的内核，每个线程处理共享内存数组的四个元素。使用单个线程处理共享内存阵列的多个元素的方法可能是有益的，即使每个块的线程数等限制不是问题。这是因为线程可以执行一次对每个元素通用的某些操作，从而将成本摊销到线程处理的共享内存元素的数量上。

确定性能对占用率敏感度的一种有用技术是通过试验动态分配的共享内存量，如执行配置的第三个参数中指定的那样。通过简单地增加此参数（不修改内核），可以有效地减少内核的占用率并衡量其对性能的影响。

10.5. 并发内核执行

如异步和重叠传输与计算中所述，CUDA 流可用于将内核执行与数据传输重叠。在能够并发执行内核的设备上，流还可用于同时执行多个内核，以更充分地利用设备的多处理器。设备是否具有此功能由结构字段指示（或在 CUDA 示例的输出中列出）。并发执行需要非默认流（流 0 以外的流），因为使用默认流的内核调用仅在设备（在任何流中）上的所有先前调用完成后开始，并且在设备（在任何流中）上的任何操作都开始，直到它们完成。concurrentKernelscudaDevicePropdeviceQuery

以下示例说明了基本技术。由于和在不同的非默认流中执行，因此有能力的设备可以同时执行内核。kernel1kernel2

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 中是完全隐式的，该 API 会自动创建和管理上下文。

使用 CUDA 驱动程序 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);

11. 指令优化

了解指令的执行方式通常允许进行低级优化，这在频繁运行的代码中（即程序中的所谓热点）非常有用。最佳做法建议在完成所有更高级别的优化后执行此优化。

11.1. 算术指令

单精度浮子可提供最佳性能，强烈建议使用单精度浮子。CUDA C++ 编程指南中详细介绍了单个算术运算的吞吐量。

11.1.1. 除法模运算

整数除法和模运算的成本特别高，应尽可能避免或用按位运算代替：If 是 2 的幂，（）等价于（），（）等价于（），（）等效于（）。ni/ni≫log2(n)i%ni&(n−1)

如果 n 是文本，编译器将执行这些转换。（有关详细信息，请参阅 CUDA C++ 编程指南中的性能指南）。

11.1.2. 有符号与无符号的循环计数器

在 C 语言标准中，无符号整数溢出语义定义明确，而有符号整数溢出会导致未定义的结果。因此，与使用无符号算术相比，编译器可以使用有符号算术进行更积极的优化。对于循环计数器来说，这一点尤其要注意：由于循环计数器的值始终为正是很常见的，因此将计数器声明为无符号可能很诱人。但是，为了获得稍微好一点的性能，应将它们声明为已签名。

例如，请考虑以下代码：

for (i = 0; i < n; i++) {
    out[i] = in[offset + stride*i];
}

在这里，子表达式可能会溢出 32 位整数，因此，如果声明为无符号，则溢出语义会阻止编译器使用一些可能已应用的优化，例如强度降低。如果改为声明为 signed，其中溢出语义未定义，则编译器有更多的余地来使用这些优化。stride*iii

11.1.3. 倒数平方根

对于单精度和双精度，应始终显式调用倒数平方根。编译器仅在不违反 IEEE-754 语义的情况下进行优化。rsqrtf()rsqrt()1.0f/sqrtf(x)rsqrtf()

11.1.4. 其他算术指令

编译器有时必须插入转换指令，从而引入额外的执行周期。这是以下情况：

在操作数上运行的函数或其操作数通常需要转换为charshortint
用作单精度浮点计算输入的双精度浮点常数（定义不带任何类型后缀）

后一种情况可以通过使用单精度浮点常数来避免，该浮点常数使用后缀（如、、等）定义。f3.141592653589793f1.0f0.5f

对于单精度代码，强烈建议使用 float 类型和单精度数学函数。

还应该注意的是，CUDA 数学库的互补误差函数，在完全单精度精度下特别快。erfcf()

11.1.5. 使用小分数参数的幂

对于一些分数指数，与使用平方根、立方根及其逆函数相比，幂可以显着加速。对于那些指数不能完全表示为浮点数的幂（例如 1/3），这也可以提供更准确的结果，因为使用会放大初始表示误差。pow()pow()

下表中的公式对有效，即。x >= 0, x != -0signbit(x) == 0

11.1.6. 数学库

支持两种类型的运行时数学运算。它们可以通过它们的名称来区分：有些名称带有前置下划线，而另一些则没有（例如，与）。遵循命名约定的函数直接映射到硬件级别。它们速度更快，但精度稍低（例如，和）。遵循命名约定的函数速度较慢，但具有更高的准确性（例如，和）。、和的吞吐量远大于、和的吞吐量。如果需要减小论点的幅度，后者的成本会变得更加昂贵（大约慢一个数量级）。此外，在这种情况下，参数减少代码使用本地内存，由于本地内存的高延迟，这可能会对性能产生更大的影响。有关更多详细信息，请参阅 CUDA C++ 编程指南。__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

另请注意，每当计算同一参数的正弦和余弦时，都应使用指令系列来优化性能：sincos

__sincosf()对于单精度快速数学运算（见下一段）
sincosf()用于常规单精度
sincos()用于双精度

编译器选项强制将每次调用都转换为等效调用。它还禁用了单精度非正态支持，并降低了单精度除法的精度。这是一种激进的优化，既可以降低数值精度，又可以改变特殊情况的处理。一种更稳健的方法是，只有在性能提升方面有选择地引入对快速内在函数的调用，并且可以容忍改变的行为。请注意，此开关仅在单精度浮点上有效。-use_fast_mathnvccfunctionName()__functionName()

中等优先级：如果可能的话，更喜欢更快、更专业的数学函数，而不是更慢、更通用的数学函数。

对于小整数幂（例如 x2 或 x3），显式乘法几乎可以肯定比使用一般幂例程（如）更快。虽然编译器优化的改进不断寻求缩小这一差距，但显式乘法（或使用等效的专用内联函数或巨集）可以具有显著的优势。当需要相同基数的多次幂时（例如，x2 和 x5 的计算都非常接近），这种优势会增加，因为这有助于编译器进行公共子表达式消除（CSE）优化。pow()

对于使用以 2 或 10 为基数的求幂，请使用函数 or 和或代替函数或。由于在一般幂中会出现许多特殊情况，并且难以在基数和指数的整个范围内实现良好的精度，因此在寄存器压力和指令计数方面，两者都是重量级函数。另一方面，函数、、和在性能方面与和相似，并且可能比它们的 / 等效物快十倍。exp2()expf2()exp10()expf10()pow()powf()pow()powf()exp2()exp2f()exp10()exp10f()exp()expf()pow()powf()

对于指数为 1/3 的幂，请使用 or 函数而不是通用幂函数 or，因为前者明显快于后者。同样，对于指数为 -1/3 的指数，请使用或。cbrt()cbrtf()pow()powf()rcbrt()rcbrtf()

替换为、和。这在准确性和性能方面都是有利的。作为一个特定示例，要以度数而不是弧度计算正弦函数，请使用。同样，当函数参数的形式为时，单精度函数，和应替换对、和的调用。（性能优势是由于简化的参数减少;准确性优势是因为仅隐式地乘以，有效地使用了无限精确的数学，而不是其单精度或双精度近似。sin(π*<expr>)sinpi(<expr>)cos(π*<expr>)cospi(<expr>)sincos(π*<expr>)sincospi(<expr>)sinpi(x/180.0)sinpif()cospif()sincospif()sinf()cosf()sincosf()π*<expr>sinpi()sin()sinpi()ππ

11.1.7. 精度相关的编译器标志

默认情况下，编译器会生成符合 IEEE 的代码，但它也提供了用于生成准确度稍低但速度更快的代码的选项：nvcc

-ftz=true（非规范化数字被刷新为零）
-prec-div=false（不太精确的划分）
-prec-sqrt=false（不太精确的平方根）

另一个更具侵略性的选项是，它将每个调用强制到等效调用。这使得代码运行得更快，但代价是精度和准确性降低。请参阅数学库。-use_fast_mathfunctionName()__functionName()

11.2. 内存指令

注意

高优先级：尽量减少全局内存的使用。在可能的情况下，首选共享内存访问。

内存指令包括从共享内存、本地内存或全局内存读取或写入的任何指令。访问未缓存的本地或全局内存时，内存延迟有数百个时钟周期。

例如，以下示例代码中的赋值运算符具有较高的吞吐量，但至关重要的是，从全局内存读取数据时存在数百个时钟周期的延迟：

__shared__ float shared[32];
__device__ float device[32];
shared[threadIdx.x] = device[threadIdx.x];

如果在等待全局内存访问完成时可以发出足够的独立算术指令，则线程调度器可以隐藏大部分全局内存延迟。但是，最好尽可能避免访问全局内存。

12. 控制流程

12.1. 分支和背离

流量控制指令（、、、）会导致相同翘速的线程发散，从而显著影响指令吞吐量;也就是说，遵循不同的执行路径。如果发生这种情况，则必须单独执行不同的执行路径;这会增加为此翘曲执行的指令总数。ifswitchdoforwhile

为了在控制流依赖于线程 ID 的情况下获得最佳性能，应编写控制条件，以最大程度地减少发散翘曲的数量。

这是可能的，因为翘曲在块中的分布是确定性的，如 CUDA C++ 编程指南的 SIMT 架构中所述。一个简单的例子是，当控制条件仅取决于（ / ）时，翘曲尺寸在哪里。threadIdxWSIZEWSIZE

在这种情况下，没有经线发散，因为控制条件与经线完全对齐。

对于仅包含几条指令的分支，扭曲背离通常会导致边际性能损失。例如，编译器可以使用谓词来避免实际的分支。相反，所有指令都是调度的，但每个线程的条件代码或谓词控制哪些线程执行指令。具有虚假谓词的线程不会写入结果，也不会计算地址或读取操作数。

从 Volta 架构开始，独立线程调度允许翘曲在依赖于数据的条件块之外保持发散。显式可用于保证翘曲已重新收敛以进行后续指令。__syncwarp()

12.2. 分支谓词

有时，编译器可能会使用分支谓词来展开循环或优化 or 语句。在这些情况下，任何经线都不可能发散。编程器还可以使用以下命令控制循环展开ifswitch

#pragma unroll

有关此编译指示的更多信息，请参阅 CUDA C++ 编程指南。

使用分支谓词时，不会跳过执行取决于控制条件的任何指令。取而代之的是，每个这样的指令都与每个线程的条件代码或谓词相关联，该代码或谓词根据控制条件设置为 true 或 false。尽管这些指令中的每一个都被安排执行，但实际上只有具有真实谓词的指令才会被执行。带有假谓词的指令不会写入结果，也不会计算地址或读取操作数。

仅当分支条件控制的指令数量小于或等于某个阈值时，编译器才会用谓词指令替换分支指令。

13. 部署 CUDA 应用程序

在完成了应用程序的一个或多个组件的 GPU 加速后，可以将结果与原始预期进行比较。回想一下，初始评估步骤允许开发人员确定通过加速给定热点可实现的潜在加速的上限。

在处理其他热点以提高总加速之前，开发人员应考虑采用部分并行化的实现，并将其贯彻到生产环境中。这很重要，原因有很多;例如，它允许用户尽早从他们的投资中获利（加速可能是部分的，但仍然很有价值），并且它通过为应用程序提供一组进化而非革命性的更改，将开发人员和用户的风险降至最低。

14. 了解编程环境

在每一代 NVIDIA 处理器中，CUDA 都可以利用的 GPU 中添加了新功能。因此，了解架构的特征非常重要。

程序员应该知道两个版本号。第一个是计算能力，第二个是 CUDA 运行时和 CUDA 驱动程序 API 的版本号。

14.1. CUDA计算能力

计算能力描述了硬件的功能，并反映了设备支持的指令集以及其他规格，例如每个块的最大线程数和每个多处理器的寄存器数。较高的计算能力版本是较低（即早期）版本的超集，因此它们向后兼容。

可以通过编程方式查询设备中 GPU 的计算能力，如 CUDA 示例中所示。该程序的输出如图 16 所示。此信息是通过调用和访问它返回的结构中的信息来获取的。deviceQuerycudaGetDeviceProperties()

计算能力的主要修订号和次要修订号显示在图 16 的第七行。此系统的设备 0 的计算能力为 7.0。

有关各种 GPU 计算能力的更多详细信息，请参阅 CUDA 启用的 GPU 和 CUDA C++ 编程指南的计算能力。特别是，开发人员应注意设备上的多处理器数量、寄存器数量和可用内存量，以及设备的任何特殊功能。

14.2. 其他硬件数据

计算功能未描述某些硬件功能。例如，无论计算能力如何，大多数 GPU 都支持将内核执行与主机和设备之间的异步数据传输重叠的功能，但并非所有 GPU 都可用。在这种情况下，请拨打以确定设备是否支持特定功能。例如，设备属性结构的字段指示是否可以重叠的内核执行和数据传输（如果可以，则可能有多少个并发传输）;同样，该字段指示是否可以执行零拷贝数据传输。cudaGetDeviceProperties()asyncEngineCountcanMapHostMemory

14.3. 哪个计算能力目标

要面向 NVIDIA 硬件和 CUDA 软件的特定版本，请使用的、和选项。例如，使用warp shuffle操作的代码必须使用（或更高的计算能力）进行编译。-arch-code-gencodenvcc-arch=sm_30

有关用于同时为多代支持 CUDA 的设备构建代码的标志的进一步讨论，请参阅构建以实现最大兼容性。

14.4. CUDA 运行时

CUDA 软件环境的主机运行时组件只能由主机函数使用。它提供用于处理以下内容的函数：

设备管理
上下文管理
内存管理
代码模块管理
执行控制
纹理引用管理
与 OpenGL 和 Direct3D 的互操作性

与较低级别的 CUDA 驱动程序 API 相比，CUDA 运行时通过提供隐式初始化、上下文管理和设备代码模块管理，大大简化了设备管理。生成的 C++ 主机代码利用了 CUDA 运行时，因此链接到此代码的应用程序将依赖于 CUDA 运行时;同样，任何使用、、和其他 CUDA 工具包库的代码也将依赖于 CUDA 运行时，该运行时由这些库在内部使用。nvcccuBLAScuFFT

CUDA 工具包参考手册中介绍了构成 CUDA 运行时 API 的函数。

CUDA 运行时在启动内核之前处理内核加载和设置内核参数和启动配置。隐式驱动程序版本检查、代码初始化、CUDA 上下文管理、CUDA 模块管理（cubin 到函数映射）、内核配置和参数传递都由 CUDA 运行时执行。

它由两个主要部分组成：

C 样式函数接口（）。cuda_runtime_api.h
C++ 风格的便利包装器（）构建在 C 样式函数之上。cuda_runtime.h

有关运行时 API 的更多信息，请参阅 CUDA C++ 编程指南的 CUDA 运行时。

15. CUDA 兼容性开发者指南

CUDA 工具包按月发布频率发布，以提供新功能、性能改进和关键错误修复。CUDA 兼容性允许用户更新最新的 CUDA Toolkit 软件（包括编译器、库和工具），而无需更新整个驱动程序堆栈。

CUDA软件环境由三个部分组成：

CUDA Toolkit（库、CUDA 运行时和开发者工具） - 用于开发人员构建 CUDA 应用程序的 SDK。
CUDA 驱动程序 - 用于运行 CUDA 应用程序的用户模式驱动程序组件（例如，在 Linux 系统上 libcuda.so）。
NVIDIA GPU 设备驱动程序 - NVIDIA GPU 的内核模式驱动程序组件。

在 Linux 系统上，CUDA 驱动程序和内核模式组件一起在 NVIDIA 显示驱动程序包中提供。如图 1 所示。

CUDA的组件

CUDA 编译器（nvcc）提供了一种处理 CUDA 和非 CUDA 代码（通过拆分和引导编译）的方法，以及 CUDA 运行时，是 CUDA 编译器工具链的一部分。CUDA 运行时 API 为开发人员提供了高级 C++ 接口，用于简化设备管理、内核执行等，而 CUDA 驱动程序 API 为应用程序提供（CUDA 驱动程序 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 平台的兼容性旨在解决以下几种情况：

NVIDIA 驱动程序升级到在企业或数据中心生产环境中运行的具有 GPU 的系统可能很复杂，可能需要提前规划。推出新的 NVIDIA 驱动程序的延迟可能意味着此类系统的用户可能无法访问 CUDA 版本中可用的新功能。不需要对新的 CUDA 版本进行驱动程序更新意味着用户可以更快地获得软件的新版本。
许多基于 CUDA 构建的软件库和应用程序（例如数学库或深度学习框架）并不直接依赖于 CUDA 运行时、编译器或驱动程序。在这种情况下，用户或开发人员仍然可以从而无需升级整个 CUDA 工具包或驱动程序来使用这些库或框架中受益。
升级依赖项容易出错且耗时，在某些极端情况下，甚至会改变程序的语义。不断使用最新的 CUDA 工具包重新编译意味着强制应用程序产品的最终客户进行升级。包管理器促进了此过程，但仍然可能出现意外问题，如果发现错误，则需要重复上述升级过程。

CUDA 支持多种兼容性选择：

CUDA Forward Compatible Upgrade 首次在 CUDA 10 中引入，旨在允许用户访问新的 CUDA 功能，并在安装了旧版 NVIDIA 数据中心驱动程序的系统上运行使用新 CUDA 版本构建的应用程序。
CUDA 增强兼容性在 CUDA 11.1 中首次引入，它提供了两个好处：
- 通过利用 CUDA 工具包中跨组件的语义版本控制，可以为一个 CUDA 次要版本（例如 11.1）构建应用程序，并在主要系列（即 11.x）中的所有未来次要版本中运行。
- CUDA 运行时放宽了最低驱动程序版本检查，因此在迁移到新的次要版本时不再需要驱动程序升级。
CUDA 驱动程序可确保为编译的 CUDA 应用程序保持向后二进制兼容性。使用 CUDA 工具包版本 3.2 编译的应用程序将在较新的驱动程序上运行。

15.2. 源代码兼容性

我们将源兼容性定义为库提供的一组保证，其中，当安装了较新版本的 SDK 时，针对特定版本的库（使用 SDK）构建的格式良好的应用程序将继续构建和运行，而不会出错。

CUDA 驱动程序和 CUDA 运行时在不同的 SDK 版本之间不兼容源代码。可以弃用和删除 API。因此，在较旧版本的工具包上成功编译的应用程序可能需要进行更改，以便针对较新版本的工具包进行编译。

通过弃用和文档机制通知开发人员任何当前或即将发生的更改。这并不意味着使用旧工具包编译的应用程序二进制文件将不再受支持。应用程序二进制文件依赖于 CUDA 驱动程序 API 接口，即使 CUDA 驱动程序 API 本身可能在工具包版本之间也发生了变化，但 CUDA 保证了 CUDA 驱动程序 API 接口的二进制兼容性。

15.3. 二进制兼容性

我们将二进制兼容性定义为库提供的一组保证，其中针对所述库的应用程序在动态链接到库的不同版本时将继续工作。

CUDA 驱动程序 API 具有版本控制的 C 样式 ABI，它保证针对旧驱动程序（例如 CUDA 3.2）运行的应用程序仍将针对现代驱动程序（例如 CUDA 11.0 附带的驱动程序）正常运行。这意味着，即使必须针对较新的 CUDA 工具包重新编译应用程序源以使用较新的功能，也可能需要更改应用程序源，但将系统中安装的驱动程序组件替换为较新版本将始终支持现有应用程序及其功能。

因此，CUDA 驱动程序 API 是二进制兼容的（操作系统加载器可以选择较新的版本，应用程序继续工作），但不兼容源代码（针对较新的 SDK 重新构建应用程序可能需要更改源代码）。

CUDA 工具包和最低驱动程序版本

CUDA 工具包和最低驱动程序版本

在我们进一步讨论此主题之前，开发人员必须了解最低驱动程序版本的概念以及这可能如何影响他们。

CUDA 工具包的每个版本（和运行时）都需要 NVIDIA 驱动程序的最低版本。针对 CUDA 工具包版本编译的应用程序将仅在具有该工具包版本的指定最低驱动程序版本的系统上运行。在 CUDA 11.0 之前，工具包的最低驱动程序版本与该版本的 CUDA 工具包附带的驱动程序相同。

因此，当使用 CUDA 11.0 构建应用程序时，它只能在具有 R450 或更高版本驱动程序的系统上运行。如果此类应用程序在安装了 R418 驱动程序的系统上运行，CUDA 初始化将返回错误，如下例所示。

在此示例中，deviceQuery 示例使用 CUDA 11.1 编译，并在具有 R418 的系统上运行。在这种情况下，由于最低驱动程序要求，CUDA 初始化将返回错误。

请参阅 CUDA 工具包发行说明，了解最低驱动程序版本和工具包附带的驱动程序版本的详细信息。

15.3.1. CUDA 二进制（cubin）兼容性

一个稍微相关但重要的主题是 CUDA 中跨 GPU 架构的应用程序二进制兼容性之一。

CUDA C++ 为熟悉 C++ 编程语言的用户提供了一条简单的路径，可以轻松编写供设备执行的程序。可以使用称为 PTX 的 CUDA 指令集架构编写内核，PTX 参考手册中对此进行了描述。但是，使用高级编程语言（如 C++）通常更有效。在这两种情况下，内核都必须由 nvcc（称为 cubins）编译成二进制代码才能在设备上执行。

cubins 是特定于体系结构的。从一个计算能力次要修订版到下一个计算能力修订版，可以保证 cubin 的二进制兼容性，但不能从一个计算能力次要修订版到上一个修订版，也不会保证主要计算能力修订版之间的二进制兼容性。换言之，为计算能力 X.y 生成的立方体对象将仅在计算能力为 X.z 的设备上执行，其中 z≥y。

若要在具有特定计算能力的设备上执行代码，应用程序必须加载与此计算能力兼容的二进制代码或 PTX 代码。为了便携性，即为了能够在具有更高计算能力的未来 GPU 架构上执行代码（尚无法生成二进制代码），应用程序必须加载 PTX 代码，这些代码将由 NVIDIA 驱动程序为这些未来设备实时编译。

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

有关 cubin、PTX 和应用程序兼容性的更多信息，请参阅 CUDA C++ 编程指南。

15.4. CUDA在次要版本之间的兼容性

通过利用语义版本控制，从 CUDA 11 开始，CUDA 工具包中的组件将在工具包的次要版本之间保持二进制兼容。为了在次要版本之间保持二进制兼容性，CUDA 运行时不再提高每个次要版本所需的最低驱动程序版本 - 这仅在主要版本发布时发生。

新工具链需要新的最小驱动程序的主要原因之一是处理 PTX 代码的 JIT 编译和二进制代码的 JIT 链接。

在本节中，我们将回顾在利用 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 的一部分不需要新的驱动程序，它们都可以在没有任何驱动程序依赖项的情况下使用。例如，API 或 CUDA 11.0 之前引入的任何 API（例如）不需要驱动程序升级。要使用次要版本中引入的其他 CUDA API（需要新的驱动程序），必须实施回退或优雅地失败。这种情况与现在可用的情况没有什么不同，在当今可用的情况下，开发人员使用宏来编译基于 CUDA 版本的功能。用户应参阅 CUDA 标头和文档，了解发布中引入的新 CUDA API。cuMemMapcudaDeviceSynchronize

使用工具包的次要版本中公开的功能时，如果应用程序正在针对较旧的 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 定义了虚拟机和 ISA 以执行通用并行线程。PTX 程序在加载时通过 JIT 编译器转换为目标硬件指令集，JIT 编译器是 CUDA 驱动程序的一部分。由于 PTX 是由 CUDA 驱动程序编译的，因此新的工具链将生成与旧版 CUDA 驱动程序不兼容的 PTX。当 PTX 用于未来的设备兼容性（最常见的情况）时，这不是问题，但在用于运行时编译时可能会导致问题。

对于继续使用 PTX 的代码，为了支持在较旧的驱动程序上进行编译，必须首先通过静态 ptxjitcompiler 库或 NVRTC 将您的代码转换为设备代码，并可选择为特定架构（例如 sm_80）而不是虚拟架构（例如 compute_80）生成代码。对于此工作流，CUDA 工具包附带了一个新的 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 驱动程序中的 * 组 API 将无法与增强的兼容性一起工作。这些 API 目前需要与 CUDA 运行时版本匹配的升级驱动程序。cuLink

如 PTX 部分所述，将 PTX 编译为设备代码与 CUDA 驱动程序一起存在，因此生成的 PTX 可能比部署系统上的驱动程序支持的 PTX 更新。使用 NVRTC 时，建议首先通过 PTX 用户工作流概述的步骤将生成的 PTX 代码转换为最终设备代码。这样可以确保您的代码兼容。或者，NVRTC 可以从 CUDA 11.1 开始直接生成 cubin。使用新 API 的应用程序可以直接使用驱动程序 API 和加载最终设备代码。cuModuleLoadDatacuModuleLoadDataEx

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 （）：CUDA 驱动程序 API ：： CUDA 工具包文档。cuGetProcAddress

使用共享库或静态库时，请按照所述库的发行说明确定该库是否支持次要版本兼容性。

16. 准备部署

16.1. 测试 CUDA 可用性

在部署 CUDA 应用程序时，通常需要确保应用程序将继续正常运行，即使目标计算机没有安装 CUDA 的 GPU 和/或足够版本的 NVIDIA 驱动程序。（面向具有已知配置的单台计算机的开发人员可以选择跳过此部分。

检测支持 CUDA 的 GPU

当应用程序将部署到任意/未知配置的目标计算机时，应用程序应明确测试是否存在支持 CUDA 的 GPU，以便在没有此类设备可用时采取适当的操作。该函数可用于查询可用设备的数量。与所有 CUDA 运行时 API 函数一样，如果没有支持 CUDA 的 GPU 或未安装适当版本的 NVIDIA 驱动程序，则此函数将正常失败并返回到应用程序。如果报告错误，应用程序应回退到备用代码路径。cudaGetDeviceCount()cudaErrorNoDevicecudaErrorInsufficientDrivercudaGetDeviceCount()

具有多个 GPU 的系统可能包含不同硬件版本和功能的 GPU。使用来自同一应用程序的多个 GPU 时，建议使用相同类型的 GPU，而不是混合使用各代硬件。该功能可用于选择与所需功能集最匹配的设备。cudaChooseDevice()

检测硬件和软件配置

当应用程序依赖于某些硬件或软件功能的可用性来启用某些功能时，可以查询 CUDA API 以获取有关可用设备配置和已安装软件版本的详细信息。

该函数报告可用设备的各种功能，包括设备的 CUDA 计算能力（另请参阅 CUDA C++ 编程指南的计算能力部分）。有关如何查询可用的 CUDA 软件 API 版本的详细信息，请参阅版本管理。cudaGetDeviceProperties()

16.2. 错误处理

所有 CUDA 运行时 API 调用都返回一个类型的错误代码;如果未发生任何错误，则返回值将等于。（例外情况是内核启动，它返回 void，以及，它返回一个字符串，描述传递给其中的代码。CUDA 工具包库（、等）同样返回自己的错误代码集。cudaError_tcudaSuccesscudaGetErrorString()cudaError_tcuBLAScuFFT

由于某些 CUDA API 调用和所有内核启动相对于主机代码都是异步的，因此错误也可能异步报告给主机;通常，这种情况发生在主机和设备下次相互同步时，例如在调用 to 或 to 期间。cudaMemcpy()cudaDeviceSynchronize()

始终检查所有 CUDA API 函数的错误返回值，即使是预计不会失败的函数，因为这将允许应用程序在发生错误时尽快检测并从错误中恢复。要使用不返回任何错误代码的语法检查内核启动期间发生的错误，应在内核启动后立即检查的返回代码。不检查 CUDA API 错误的应用程序有时可能会运行完成，而不会注意到 GPU 计算的数据不完整、无效或未初始化。<<<...>>>cudaGetLastError()

16.3. 构建最大兼容性

每一代支持 CUDA 的设备都有一个关联的计算能力版本，该版本指示设备支持的功能集（请参阅 CUDA 计算能力）。在构建文件时，可以为 nvcc 编译器指定一个或多个计算功能版本;编译应用程序目标 GPU 的本机计算功能对于确保应用程序内核实现最佳性能并能够使用给定一代 GPU 上可用的功能非常重要。

当同时为多个计算功能构建应用程序时（使用 nvcc 标志的多个实例），指定计算功能的二进制文件将合并到可执行文件中，并且 CUDA 驱动程序在运行时根据当前设备的计算能力选择最合适的二进制文件。如果没有合适的原生二进制文件（cubin），但中间 PTX 代码（面向抽象虚拟指令集并用于向前兼容性）可用，则内核将从 PTX 编译为设备的原生 cubin。如果 PTX 也不可用，则内核启动将失败。-gencode

window

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

但是，虽然命令行选项默认情况下确实会导致包含 PTX 后端目标（由于它暗示了目标），但它一次只能指定单个目标体系结构，并且不可能在同一命令行上使用多个选项，这就是上述示例显式使用的原因。-arch=sm_XXcode=compute_XXcubin-arch=nvcc-gencode=

16.4. 分发 CUDA 运行时和库

CUDA 应用程序是针对 CUDA 运行时库构建的，该库处理设备、内存和内核管理。与 CUDA 驱动程序不同，CUDA 运行时既不保证版本之间的向前兼容性，也不保证向后二进制兼容性。因此，在使用动态链接时，最好将 CUDA 运行时库与应用程序一起重新分发，或者静态链接到 CUDA 运行时。这将确保即使用户没有安装构建应用程序所针对的相同 CUDA 工具包，可执行文件也能够运行。

静态链接的 CUDA 运行时

最简单的选择是静态链接到 CUDA 运行时。如果在 CUDA 5.5 及更高版本中使用链接，这是默认设置。静态链接使可执行文件稍大一些，但它可以确保应用程序二进制文件中包含正确版本的运行时库函数，而无需单独重新分发 CUDA 运行时库。nvcc

动态链接的 CUDA 运行时

如果由于某种原因对 CUDA 运行时的静态链接是不切实际的，那么也可以使用动态链接版本的 CUDA 运行时库。（这是 CUDA 版本 5.0 及更早版本中提供的默认且唯一选项。

在使用 from CUDA 5.5 或更高版本链接应用程序时，要对 CUDA 运行时使用动态链接，请将标志添加到链接命令行;否则，默认情况下使用静态链接的 CUDA 运行时库。nvcc--cudart=shared

在应用程序与 CUDA 运行时动态链接后，此版本的运行时库应与应用程序捆绑在一起。它可以复制到与应用程序可执行文件相同的目录中，也可以复制到该安装路径的子目录中。

其他 CUDA 库

尽管 CUDA 运行时提供了静态链接的选项，但 CUDA 工具包中包含的某些库仅以动态链接形式提供。与 CUDA 运行时库的动态链接版本一样，在分发该应用程序时，这些库应与应用程序可执行文件捆绑在一起。

16.4.1. CUDA Toolkit库重新分发

CUDA 工具包的最终用户许可协议（EULA）允许在某些条款和条件下重新分发许多 CUDA 库。这使得依赖于这些库的应用程序能够重新分发构建和测试它们的库的确切版本，从而避免了可能在其计算机上安装了不同版本的CUDA工具包（或者可能根本没有安装）的最终用户的任何麻烦。有关详细信息，请参阅 EULA。

16.4.1.1. 要重新分发的文件

在重新分发一个或多个 CUDA 库的动态链接版本时，确定需要重新分发的确切文件非常重要。以下示例使用 CUDA Toolkit 5.5 中的 cuBLAS 库作为说明：

Linux操作系统

在 Linux 上的共享库中，有一个名为 the 的字符串字段，表示库的二进制兼容级别。生成应用程序所针对的库的文件名必须与随应用程序一起重新分发的库的文件名匹配。SONAMESONAME

例如，在标准 CUDA 工具包安装中，文件和都是指向 cuBLAS 特定构建的符号链接，其命名类似于，其中 x 是构建号（例如，）。但是，此库的 “”表示为“：libcublas.solibcublas.so.5.5libcublas.so.5.5.xlibcublas.so.5.5.17SONAMElibcublas.so.5.5

$ objdump -p /usr/local/cuda/lib64/libcublas.so | grep SONAME
   SONAME               libcublas.so.5.5

因此，即使在链接应用程序时使用（未指定版本号），在链接时找到也意味着“”是动态加载器在加载应用程序时将查找的文件的名称，因此必须是与应用程序一起重新分发的文件的名称（或指向相同的符号链接）。-lcublasSONAMElibcublas.so.5.5

该工具可用于识别应用程序希望在运行时找到的库的确切文件名，以及动态加载器在给定当前库搜索路径的情况下加载应用程序时将选择的该库副本的路径（如果有）：ldd

$ ldd a.out | grep libcublas
   libcublas.so.5.5 => /usr/local/cuda/lib64/libcublas.so.5.5

苹果电脑

在 Mac OS X 上的共享库中，有一个名为 the 的字段，用于指示预期的安装路径和库的文件名;CUDA 库也使用此文件名来表示二进制兼容性。此字段的值将传播到针对该库构建的应用程序中，并用于在运行时查找正确版本的库。install name

例如，如果 cuBLAS 库的安装名称为，则该库的版本为 5.5，并且必须命名与应用程序一起重新分发的此库的副本，即使在链接时仅使用（未指定版本号）。此外，此文件应安装到应用程序中;请参阅在何处安装重新分发的 CUDA 库。@rpath/libcublas.5.5.dyliblibcublas.5.5.dylib-lcublas@rpath

要查看库的安装名称，请使用以下命令：otool -L

$ otool -L a.out
a.out:
        @rpath/libcublas.5.5.dylib (...)

window

Windows 上 CUDA 库的二进制兼容版本在文件名中表示。

例如，链接到 cuBLAS 5.5 的 64 位应用程序将在运行时查找，因此这是应与该应用程序一起重新分发的文件，即使它是应用程序链接的文件。对于 32 位应用程序，该文件将为 .cublas64_55.dllcublas.libcublas32_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/苹果电脑

nvcc -I $(CUDA_HOME)/include
  -Xlinker "-rpath '$ORIGIN'" --cudart=shared
  -o myprogram myprogram.cu

window

nvcc.exe -ccbin "C:\vs2008\VC\bin"
  -Xcompiler "/EHsc /W3 /nologo /O2 /Zi /MT" --cudart=shared
  -o "Release\myprogram.exe" "myprogram.cu"

若要指定将分发库的备用路径，请使用类似于下面的链接器选项：

Linux/苹果电脑

nvcc -I $(CUDA_HOME)/include
  -Xlinker "-rpath '$ORIGIN/lib'" --cudart=shared
  -o myprogram myprogram.cu

window

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，该选项将像以前一样使用。对于 Windows，使用该选项;这要求应用程序在第一次调用任何 CUDA API 函数之前调用，以指定包含 CUDA DLL 的目录。-rpath/DELAYSetDllDirectory()

17. 部署基础设施工具

17.1. 英伟达-SMI

NVIDIA 系统管理接口（）是一个命令行实用程序，可帮助管理和监视 NVIDIA GPU 设备。此实用程序允许管理员查询 GPU 设备状态，并允许管理员以适当的权限修改 GPU 设备状态。针对 Tesla 和某些 Quadro GPU，但其他 NVIDIA GPU 也提供有限的支持。在 Linux 以及 64 位 Windows Server 2008 R2 和 Windows 7 上附带 NVIDIA GPU 显示驱动程序。可以将查询的信息输出为 XML 或人类可读的纯文本到标准输出或文件。有关详细信息，请参阅 nvidia-smi 文档。请注意，nvidia-smi 的新版本不保证与以前的版本向后兼容。nvidia-sminvidia-sminvidia-sminvidia-smi

17.1.1. 可查询状态

ECC 错误计数

报告了可纠正的单位错误和可检测的双比特错误。提供了当前启动周期和 GPU 生命周期的错误计数。

GPU 利用率

报告了 GPU 的计算资源和内存接口的当前利用率。

活动计算过程

将报告在 GPU 上运行的活动进程列表，以及相应的进程名称/ID 和分配的 GPU 内存。

时钟和性能状态

报告了几个重要时钟域的最大和当前时钟速率，以及当前 GPU 性能状态（pstate）。

温度和风扇速度

报告了当前的 GPU 核心温度，以及具有主动冷却功能的产品的风扇速度。

电源管理

对于报告这些测量值的产品，将报告当前的电路板功耗和功率限制。

鉴定

报告各种动态和静态信息，包括板卡序列号、PCI 设备 ID、VBIOS/Inforom 版本号和产品名称。

17.1.2. 可修改状态

ECC模式

启用和禁用 ECC 报告。

ECC复位

清除单位和双位 ECC 错误计数。

计算模式

指示计算进程是否可以在 GPU 上运行，以及它们是以独占方式运行还是与其他计算进程同时运行。

持久化模式

指示当没有应用程序连接到 GPU 时，NVIDIA 驱动程序是否保持加载状态。在大多数情况下，最好启用此选项。

GPU 重置

通过辅助总线复位重新初始化 GPU 硬件和软件状态。

17.2. NVML的

NVIDIA 管理库（NVML）是一个基于 C 的接口，可直接访问通过旨在作为构建第三方系统管理应用程序的平台公开的查询和命令。NVML API 随 CUDA 工具包一起提供（从版本 8.0 开始），也可以作为 GPU 部署套件的一部分在 NVIDIA 开发者网站上独立提供，通过单个头文件附带 PDF 文档、存根库和示例应用程序;请参见 GPU Deployment Kit | NVIDIA Developer。NVML的每个新版本都向后兼容。nvidia-smi

为 NVML API 提供了一组额外的 Perl 和 Python 绑定。这些绑定公开了与基于 C 的接口相同的功能，并且还提供向后兼容性。Perl 绑定通过 CPAN 提供，Python 绑定通过 PyPI 提供。

所有这些产品（NVML 和 NVML 语言绑定）都会随着每个新的 CUDA 版本进行更新，并提供大致相同的功能。nvidia-smi

有关更多信息，请参阅 https://developer.nvidia.com/nvidia-management-library-nvml。

17.3. 集群管理工具

管理 GPU 集群将有助于实现最大的 GPU 利用率，并帮助您和您的用户获得最佳性能。许多业界最流行的集群管理工具都通过 NVML 支持 CUDA GPU。有关其中一些工具的列表，请参阅 Cluster Management | NVIDIA Developer。

17.4. 编译器 JIT 缓存管理工具

应用程序在运行时加载的任何 PTX 设备代码都会由设备驱动程序进一步编译为二进制代码。这称为实时编译（JIT）。实时编译会增加应用程序加载时间，但允许应用程序从最新的编译器改进中受益。这也是应用程序在编译应用程序时不存在的设备上运行的唯一方式。

使用 PTX 设备代码的 JIT 编译时，NVIDIA 驱动程序会将生成的二进制代码缓存在磁盘上。此行为的某些方面（如缓存位置和最大缓存大小）可以通过使用环境变量来控制;请参阅 CUDA C++ 编程指南的实时编译。

17.5. CUDA_VISIBLE_DEVICES

通过环境变量，可以重新排列已安装的 CUDA 设备的集合，这些设备将对 CUDA 应用程序可见并由 CUDA 应用程序枚举。CUDA_VISIBLE_DEVICES

要对应用程序可见的设备应作为逗号分隔的列表包含在系统范围的可枚举设备列表中。例如，要仅使用系统范围设备列表中的设备 0 和 2，请在启动应用程序之前设置。然后，应用程序会将这些设备分别枚举为设备 0 和设备 1。CUDA_VISIBLE_DEVICES=0,2

18. 建议和最佳实践

本章包含本文档中介绍的优化建议的摘要。

18.1. 整体性能优化策略

性能优化围绕三个基本策略展开：

最大化并行执行
优化内存使用以实现最大内存带宽
优化指令使用以实现最大指令吞吐量

要最大化并行执行，首先要以一种能够公开尽可能多的并行性的方式构建算法。一旦暴露了算法的并行性，就需要尽可能有效地将其映射到硬件。这是通过仔细选择每个内核启动的执行配置来完成的。应用程序还应该通过流显式公开设备上的并发执行，以及最大化主机和设备之间的并发执行，从而在更高级别上最大化并行执行。

优化内存使用首先要尽量减少主机和设备之间的数据传输，因为这些传输的带宽比内部设备数据传输的带宽低得多。此外，还应通过最大限度地使用设备上的共享内存来最小化对全局内存的内核访问。有时，最好的优化甚至可能是首先避免任何数据传输，只需在需要时重新计算数据即可。

有效带宽可能会相差一个数量级，具体取决于每种内存类型的访问模式。因此，优化内存使用的下一步是根据最佳内存访问模式组织内存访问。这种优化对于全局内存访问尤为重要，因为访问的延迟会消耗数百个时钟周期。与此相反，共享内存访问通常仅在存在高度的银行冲突时才值得优化。

至于优化指令的使用，应避免使用吞吐量低的算术指令。这表明在不影响最终结果的情况下，用精度换取速度，例如使用内部函数而不是常规函数，或者使用单精度而不是双精度。最后，由于设备的SIMT（单指令多线程）特性，必须特别注意控制流指令。

19. NVCC编译器开关

19.1. NVCC

NVIDIA 编译器驱动程序将文件转换为主机系统的 C++ 文件和设备的 CUDA 汇编或二进制指令。它支持许多命令行参数，其中以下参数对于优化和相关的最佳实践特别有用：nvcc.cu

-maxrregcount=N指定内核在每个文件级别可以使用的最大寄存器数。请参阅注册压力。（另请参阅 CUDA C++ 编程指南的执行配置中讨论的限定符，以控制每个内核使用的寄存器数量。__launch_bounds__
--ptxas-options=-v或列出每个内核寄存器、共享内存和常量内存使用情况。-Xptxas=-v
-ftz=true（非规范化数字被刷新为零）
-prec-div=false（不太精确的划分）
-prec-sqrt=false（不太精确的平方根）
-use_fast_math编译器选项强制将每次调用都强制到等效调用。这使得代码运行得更快，但代价是精度和准确性降低。请参阅数学库。nvccfunctionName()__functionName()

20. 通知

20.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 对客户在此处描述的产品的累计和累积责任应根据产品的销售条款进行限制。