2.1. 本指南的目标读者

本指南中的所有讨论均使用C++编程语言，因此您应能熟练阅读C++代码。

本指南参考并依赖于其他几份文档，您应备有这些文档以供查阅，所有这些文档均可从CUDA网站 https://docs.nvidia.com/cuda/ 免费获取。以下文档是尤为重要的参考资料：

• CUDA 安装指南
• CUDA C++ 编程指南
• CUDA 工具包参考手册

特别是，本指南的优化部分假设您已经成功下载并安装了CUDA Toolkit（如未完成，请参考适用于您平台的相关CUDA安装指南），并且对CUDA C++编程语言和环境有基本了解（如未具备，请参考CUDA C++编程指南）。

2.2. 评估、并行化、优化、部署

本指南介绍了应用程序的评估、并行化、优化、部署（APOD）设计周期，旨在帮助应用开发者快速识别其代码中最易受益于GPU加速的部分，迅速实现该效益，并尽早将由此带来的加速效果应用于生产环境。

APOD是一个循环往复的过程：只需投入最少的时间成本，即可实现、测试并部署初步的性能加速；此后，该循环可重新启动，通过识别更多优化机会、获得额外加速效果，最终将更快的应用程序版本投入生产环境。

2.3. 建议与最佳实践

在本指南中，针对CUDA C++代码的设计与实现提出了具体建议。这些建议按优先级分类，优先级综合考量了建议的影响力和适用范围。能为大多数CUDA应用带来显著改进的措施具有最高优先级，而仅影响特定场景的细微优化则被赋予较低优先级。

在实施较低优先级的建议之前，良好的做法是确保所有相关的高优先级建议均已得到应用。这种方法往往能以投入的时间获得最佳结果，并避免陷入过早优化的陷阱。

确定优化优先级时，其收益标准和适用范围会因程序性质而异。在本指南中，它们代表典型情况。您的代码可能反映不同的优先级因素。无论存在何种可能性，良好的实践是在处理低优先级项目前，先确认没有遗漏更高优先级的优化建议。

2.4. 评估您的应用程序

从超级计算机到移动电话，现代处理器日益依赖并行性来提供性能。其核心计算单元——包括控制单元、算术单元、寄存器及通常包含的部分缓存——被复制多次，并通过网络与内存相连。因此，所有现代处理器都需要并行代码才能充分利用其计算能力。

随着处理器不断演进，向程序员暴露更细粒度的并行性，许多现有应用程序仍以串行代码或粗粒度并行代码的形式存在（例如，数据被分解为多个区域并行处理，子区域通过MPI共享）。为了从包括GPU在内的任何现代处理器架构中获益，首要步骤是评估应用程序以识别热点，判断它们是否能够并行化，并理解当前及未来的相关工作量。

3.1. 主机与设备的差异

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

线程资源

主机系统的执行流水线仅能支持有限数量的并发线程。例如，配备两个32核处理器的服务器最多只能同时运行64个线程（若CPU支持同步多线程技术，此数值可能略高）。相比之下，CUDA设备上最小的并行执行单元包含32个线程（称为一个线程束）。现代NVIDIA GPU每个多处理器可支持多达2048个并发活跃线程（参见《CUDA C++编程指南》的特性与规格说明）。在拥有80个多处理器的GPU上，这意味着可同时活跃的线程数超过16万个。

线程

CPU上的线程通常是重量级实体。操作系统必须在CPU执行通道上切换线程以提供多线程能力。因此，上下文切换（当两个线程被交换时）缓慢且代价高昂。相比之下，GPU上的线程极其轻量。在典型系统中，成千上万的线程排队等待工作（以每束32个线程的线程束形式）。如果GPU必须等待某个线程束，它会直接开始执行另一个线程束的工作。由于所有活动线程都分配了独立的寄存器，在GPU线程间切换时无需交换寄存器或其他状态。资源会一直分配给每个线程，直到其执行完成。简而言之，CPU核心旨在最小化少量线程的延迟，而GPU则设计用于处理大量并发、轻量级的线程，以最大化吞吐量。

RAM

主机系统和设备各自拥有独立的物理内存。由于主机内存与设备内存是分离的，因此如《在支持CUDA的设备上运行什么？》所述，主机内存中的数据项偶尔需要在设备内存与主机内存之间进行通信。

这是CPU主机与GPU设备在并行编程方面的主要硬件差异。本文档其他部分将讨论其他差异。基于这些差异构建的应用程序可以将主机和设备视为一个统一的异构系统，其中每个处理单元都发挥其最擅长的工作类型：主机处理串行工作，设备处理并行工作。

3.2. 什么在支持CUDA的设备上运行

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

• 该设备非常适合可同时在众多数据元素上并行运行的计算。这通常涉及对大型数据集（如矩阵）进行算术运算，其中相同的操作可以同时跨越成千上万乃至数百万个元素执行。这是CUDA实现良好性能的必要条件：软件必须使用大量（通常为数千或数万）并发线程。支持并行运行大量线程的能力源于CUDA采用的上述轻量级线程模型。
• 要使用CUDA，必须将数据值从主机传输到设备。这些传输在性能方面代价高昂，应尽量减少。（参见《主机与设备间的数据传输》。）这一成本带来若干影响：操作的复杂性应能证明将数据移入移出设备的成本是合理的。为少量线程短暂使用而传输数据的代码将几乎看不到性能提升。理想情况是许多线程执行大量工作。例如，将两个矩阵传输到设备执行矩阵加法，然后将结果传回主机，不会带来太多性能收益。这里的关键在于每个传输数据元素执行的操作数量。对于上述过程，假设矩阵大小为N×N，共有N²次操作（加法）和3N²个元素被传输，因此操作数与传输元素数之比为1:3，即O(1)。当这个比例更高时，更容易实现性能提升。例如，相同矩阵的矩阵乘法需要N³次操作（乘加），因此操作数与传输元素数之比为O(N)，此时矩阵越大，性能收益越显著。操作类型是另一个因素，因为加法与三角函数等具有不同的复杂度特征。在决定操作应在主机还是设备上执行时，必须考虑数据在设备与主机间传输的开销。数据应尽可能长时间保留在设备上。由于传输应最小化，对相同数据运行多个内核的程序，应倾向于在内核调用之间将数据保留在设备上，而不是将中间结果传输到主机再传回设备进行后续计算。因此，在前述示例中，如果两个待加矩阵已因先前计算而存在于设备上，或者加法结果将用于后续计算，则矩阵加法应在设备本地执行。即使一系列计算中的某个步骤在主机上执行可能更快，也应采用这种方法。即使一个相对较慢的内核，若能避免一次或多次主机与设备内存间的传输，也可能是有利的。《主机与设备间的数据传输》提供了更多细节，包括主机与设备之间以及设备内部带宽的测量数据。
• 为获得最佳性能，设备上运行的相邻线程在内存访问中应保持一定的连贯性。特定的内存访问模式使硬件能够将多个数据项的读写操作合并为单次操作。若数据布局无法实现合并访问，或缺乏足够的局部性以有效利用L1或纹理缓存，则在GPU计算中往往难以获得显著的加速效果。值得注意的例外是完全随机的内存访问模式。通常应避免此类模式，因为相较于峰值性能，任何架构处理此类内存访问模式的效率都较低。然而，与基于缓存的架构（如CPU）相比，具有延迟隐藏特性的架构（如GPU）往往能更好地应对完全随机的内存访问模式。

4.1. 性能分析指南

许多代码通过相对较少的代码量完成了大部分工作。开发者可以利用性能分析工具识别此类热点，并开始编译并行化的候选列表。

$ 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

6.1. 并行库

并行化应用程序最直接的方法是利用现有的库，这些库能代表我们利用并行架构的优势。CUDA工具包包含了许多针对NVIDIA CUDA GPU进行过精细调优的此类库，例如cuBLAS、cuFFT等等。

关键在于库与应用程序需求的高度匹配时最为实用。例如，已使用其他BLAS库的应用程序通常可以轻松切换至cuBLAS，而几乎不涉及线性代数的应用程序则对cuBLAS需求甚少。其他CUDA工具包库亦是如此：cuFFT的接口与FFTW相似，等等。

同样值得注意的是Thrust库，这是一个类似于C++标准模板库的并行C++模板库。Thrust提供了丰富的数据并行原语，如扫描、排序和归约，这些原语可以组合在一起，以简洁、可读的源代码实现复杂算法。通过使用这些高级抽象来描述计算，您赋予Thrust自动选择最高效实现的自由度。因此，Thrust既可用于CUDA应用程序的快速原型开发（此时程序员的生产力最为关键），也可用于生产环境（此时鲁棒性和绝对性能至关重要）。

6.2. 并行化编译器

另一种并行化顺序代码的常见方法是利用并行化编译器。这通常意味着采用基于指令的方法，程序员使用pragma或其他类似符号向编译器提示何处存在并行性，而无需修改或调整底层代码本身。通过向编译器暴露并行性，指令允许编译器执行将计算映射到并行架构上的详细工作。

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

详情请参见 http://www.openacc.org/。

6.3. 编写代码以暴露并行性

对于需要超越现有并行库或并行化编译器所能提供功能或性能的应用而言，与现有串行代码无缝集成的并行编程语言（如CUDA C++）至关重要。

一旦我们在应用程序的性能分析中定位到热点区域，并确定自定义代码是最佳方案，便可以使用CUDA C++将该代码段的并行性以CUDA内核的形式呈现。随后，我们可以将此内核加载到GPU上执行并获取结果，而无需对应用程序的其他部分进行大规模重写。

当应用程序的总运行时间大部分集中在少数相对独立的代码部分时，这种方法最为直接。对于具有非常平坦性能剖面的应用程序——即时间消耗相对均匀分布在代码库广泛部分的情况——并行化则更为困难。对于后一类应用程序，可能需要进行一定程度的代码重构以暴露其内在并行性，但请记住，这种重构工作往往会使所有未来架构（无论是CPU还是GPU）受益，因此若有必要，这项努力是非常值得的。

7.1. 验证

7.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。

7.3. 数值精度与准确度

不正确或意外的结果主要源于浮点精度问题，这是由于浮点数值的计算和存储方式导致的。以下章节将解释主要关注点。浮点运算的其他特性在《CUDA C++编程指南》的功能与技术规范中有所说明，同时也可在https://developer.nvidia.com/content/precision-performance-floating-point-and-ieee-754-compliance-nvidia-gpus 上获取关于浮点精度与性能的白皮书及配套网络研讨会。

9.1. 计时

CUDA调用与内核执行既可使用CPU计时器也可使用GPU计时器进行测量。本节将探讨两种方法的功能特性、优势及潜在缺陷。

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

9.2. 带宽

带宽——数据传输的速率——是性能最重要的制约因素之一。几乎所有代码修改都应在考虑其对带宽影响的背景下进行。正如本指南的"内存优化"部分所述，带宽会受到数据存储内存的选择、数据布局方式、访问顺序以及其他因素的显著影响。

为了准确测量性能，计算理论带宽和有效带宽是很有用的。当后者远低于前者时，设计或实现细节很可能降低了带宽，而提高带宽应成为后续优化工作的首要目标。

10.1. 主机与设备之间的数据传输

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

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

此外，由于每次传输都会产生开销，将许多小批量传输合并为一次较大的传输，其性能显著优于单独进行每次传输，即使这样做需要将非连续的内存区域打包到连续缓冲区中，并在传输后解包。

最后，如《CUDA C++编程指南》及本文档的“固定内存”部分所述，使用页锁定（或称固定）内存可实现主机与设备间更高的带宽。

10.1.2. 计算与异步及重叠传输

使用cudaMemcpy()在主机和设备之间进行数据传输属于阻塞传输；也就是说，只有在数据传输完成后控制权才会返回给主机线程。cudaMemcpyAsync()函数是cudaMemcpy()的非阻塞变体，它会立即将控制权返回给主机线程。与cudaMemcpy()相比，异步传输版本需要固定的主机内存（参见固定内存部分），并且包含一个额外的参数——流ID。流仅仅是在设备上按顺序执行的一系列操作。不同流中的操作可以交错执行，在某些情况下甚至可以重叠——这一特性可用于隐藏主机与设备之间的数据传输。

异步传输通过两种不同的方式实现数据传输与计算的重叠。在所有支持CUDA的设备上，主机计算可以与异步数据传输及设备计算实现重叠。例如，"异步传输与计算的重叠"示例展示了在数据传输至设备并执行使用该设备的内核时，如何同时执行例程cpuFunction()中的主机计算。

重叠计算与数据传输

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

cudaMemcpyAsync()函数的最后一个参数是流ID，此处使用默认流（流0）。该内核同样使用默认流，且其执行将等待内存复制完成后才开始；因此无需显式同步。由于内存复制与内核都会立即将控制权交还给主机，主机函数cpuFunction()实现了两者的执行重叠。

在异步与计算重叠传输中，内存复制和内核执行按顺序发生。在支持并发复制与计算的设备上，可以实现设备上的内核执行与主机和设备间数据传输的重叠。设备是否具备此能力由asyncEngineCount字段指示（或在deviceQuery CUDA示例的输出中列出）。对于具备此能力的设备，实现重叠仍需使用固定主机内存，并且数据传输与内核必须使用不同的非默认流（流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字段获取，该信息同样列于deviceQuery CUDA示例的输出中。（需特别说明的是，阻塞传输无法与异步传输重叠执行，因为阻塞传输发生在默认流中，必须等待所有先前的CUDA调用完成后才能开始，且在其完成前不会允许任何其他CUDA调用启动。）图1展示了两段代码执行时间线的示意图，图中下半部分的分阶段并发复制与执行示例中nStreams等于4。

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

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

10.2. 设备内存空间

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

在这些不同的内存空间中，全局内存的容量最为充裕；各计算能力级别下各内存空间的具体容量请参见《CUDA C++编程指南》的功能与技术规格章节。全局内存、局部内存和纹理内存的访问延迟最高，其次是常量内存、共享内存以及寄存器文件。

各类内存的主要特性如表1所示。

在纹理访问的情况下，如果纹理引用被绑定到全局内存中的线性数组，则设备代码可以写入底层数组。绑定到CUDA数组的纹理引用可以通过表面写入操作进行写入（方法是将一个表面绑定到相同的底层CUDA数组存储）。应当避免在同一内核启动中从纹理读取数据的同时向其底层全局内存数组写入数据，因为纹理缓存是只读的，并且在关联的全局内存被修改时不会失效。

10.2.1. 对全局内存的合并访问

在针对支持CUDA的GPU架构进行编程时，一个非常重要的性能考量是全局内存访问的合并。设备会将一个线程束内各线程的全局内存加载与存储操作，尽可能合并为最少次数的内存事务。

ℹ️ 注意

高优先级：确保全局内存访问尽可能实现合并访问。

合并访问的要求取决于设备的计算能力，具体细节在《CUDA C++编程指南》中有详细说明。

对于计算能力6.0或更高的设备，其要求可以相当简洁地概括为：一个线程束中线程的并发访问将合并为若干次事务，这些事务的数量等于服务该线程束所有线程所需的32字节事务的数量。

对于计算能力5.2的特定设备，可以选择启用对全局内存访问的L1缓存。若在这些设备上启用L1缓存，所需的事务数量将等于所需的128字节对齐段的数量。

ℹ️ 注意

在计算能力6.0或更高的设备上，L1缓存默认启用，但无论全局加载是否缓存在L1中，数据访问单元均为32字节。

在配备GDDR内存的设备上，启用ECC时以合并访问方式访问内存尤为重要。分散的访问会增加ECC内存传输开销，尤其是在向全局内存写入数据时。

合并访问的概念在以下简单示例中得以说明。除非另有说明，这些示例假设计算能力为6.0或更高，且访问对象为4字节字。

10.2.1.1. 一种简单的访问模式

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

如果一个线程束的线程访问相邻的4字节字（例如，相邻的 float 值），四次合并的32字节事务将处理该内存访问。这种模式如图3 所示。

这种访问模式导致了四次32字节的事务，如红色矩形所示。

如果四个32字节段中的任何一个仅请求了部分字（例如，若多个线程访问了同一字，或部分线程未参与访问），系统仍会完整获取整个段。此外，即使线程束内线程的访问在四个段内部或跨段发生了重排，对于计算能力6.0或更高的设备，也仅会执行四次32字节事务。

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

如果一个线程束中的连续线程访问的内存是连续的但未与32字节段对齐，将会请求五个32字节段，如图4所示。

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

10.2.1.3. 未对齐访问的影响

使用一个简单的复制内核（例如在A copy kernel that illustrates misaligned accesses中展示的复制内核）来探究未对齐访问的影响，既简单又富有启发性。

一个展示未对齐访问的复制内核

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

在一个展示未对齐访问的复制内核中，数据从全局内存中的输入数组 in 复制到输出数组 out。该内核在主机代码的循环内执行，循环中参数 offset 从 0 变化到 32（例如，图 4 对应这些未对齐情况）。 在 NVIDIA Tesla V100（计算能力 7.0）上，不同偏移量下复制操作的有效带宽如图 5 所示。

对于NVIDIA Tesla V100，无偏移或偏移量为8字（word）倍数的全局内存访问会产生四次32字节事务。实现的带宽约为790 GB/s。否则，每个线程束将加载五个32字节段，我们预计其内存吞吐量约为无偏移情况下的4/5。

在此特定示例中，所实现的偏移内存吞吐量约为理论值的十分之九，这是因为相邻的线程束复用了其邻近线程束所获取的缓存行。因此，尽管影响仍然显著，但并未达到我们预期的严重程度。若相邻线程束未能对超额获取的缓存行表现出如此高的复用程度，其影响将会更为显著。

10.2.1.4. 跨步访问

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

为了说明跨步访问对有效带宽的影响，请参见《用于演示非单位跨步数据复制的内核》中的内核 strideCopy()，该内核以跨步 stride 个元素的方式在线程间将数据从 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展示了这样一种情况：在此场景中，线程束内的线程以步长为2的方式访问内存中的字。这一操作导致在Tesla V100（计算能力7.0）上每个线程束需要加载八个L2缓存段。

步长为2时，加载/存储效率仅为50%，因为事务中一半的元素未被使用，这代表了带宽的浪费。随着步长增加，有效带宽持续下降，直至达到一个临界点：此时为一个线程束中的32个线程加载了32个32字节的段，如图7所示。

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

10.2.2. L2 缓存

自CUDA 11.0起，计算能力8.0及以上的设备具备影响L2缓存中数据持久性的能力。由于L2缓存位于芯片上，它可能为全局内存访问提供更高的带宽和更低的延迟。

更多详细信息，请参阅《CUDA C++ 编程指南》中的 L2 访问管理章节。

10.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缓存行的抖动。

10.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。对该内存区域剩余数据（即流式数据）的访问被视为常规或流式访问，因此将使用非预留L2部分的剩余10 MB容量（除非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缓存预留的30 MB空间时，性能提升最高可达50%。然而，一旦该持久数据区域的大小超过L2缓存预留部分的大小，由于L2缓存行的频繁换出，性能会下降约10%。

为了优化性能，当持久化数据的大小超过预留的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固定为20 MB，并调整hitRatio，使得总持久数据中随机20 MB的部分驻留在L2预留缓存区域中。其余部分的持久数据将使用流式属性进行访问。这有助于减少缓存抖动。结果如下图表所示，无论持久数据是否完全放入L2预留区域，我们都观察到了良好的性能表现。

10.2.3. 共享内存

由于位于芯片上，共享内存相比本地内存和全局内存具有更高的带宽和更低的延迟——前提是线程之间不存在存储体冲突，具体细节将在下一节详述。

10.2.3.1. 共享内存与存储体

为实现并发访问时的高内存带宽，共享内存被划分为多个大小相等的内存模块（存储体），这些模块可被同时访问。因此，任何跨越n个不同存储体的n地址内存加载或存储操作均可同时进行，从而产生相当于单个存储体带宽n倍的有效带宽。

然而，如果一次内存请求的多个地址映射到同一个存储体，这些访问将被串行化。硬件会将存在存储体冲突的内存请求拆分为多个独立的无冲突请求，有效带宽会因此降低，降低的倍数等于独立内存请求的数量。唯一的例外是当一个线程束中的多个线程访问同一个共享内存地址时，此时会触发广播机制。在这种情况下，来自不同存储体的多个广播会被合并为一次从请求的共享内存位置到线程的单播多播传输。

为最小化存储体冲突，理解内存地址如何映射到存储体以及如何优化调度内存请求至关重要。

在计算能力5.x或更高版本的设备上，每个存储体每个时钟周期具有32位的带宽，且连续的32位字被分配到连续的存储体中。线程束大小为32个线程，存储体数量也为32，因此线程束内的任意线程之间都可能发生存储体冲突。更多详细信息请参阅计算能力5.x章节。

10.2.3.2. 矩阵乘法C=AB中的共享内存

共享内存支持线程块内线程间的协作。当块内多个线程需要使用来自全局内存的相同数据时，可通过共享内存仅从全局内存中读取一次数据。共享内存还可用于避免非合并内存访问：先以合并模式从全局内存加载/存储数据，随后在共享内存中重新排序。除存储体冲突外，线程束在共享内存中进行非连续或未对齐访问不会产生额外开销。

共享内存的使用通过一个简单的矩阵乘法示例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 中不同瓦片的元素。

为此，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所示。

在 for 循环的每次迭代 i 中，线程束内的线程会读取 B 图块的一行，这对于所有计算能力而言都是顺序且合并的访问。

然而，对于每次迭代 i，线程束中的所有线程都会从全局内存中读取矩阵 A 的相同值，因为索引 row*TILE_DIM+i 在线程束内是恒定的。尽管在计算能力 2.0 或更高的设备上，此类访问仅需 1 次事务，但该事务中存在带宽浪费，因为 32 字节缓存段中的 8 个字仅使用了 1 个 4 字节字。我们可以在循环的后续迭代中重用该缓存行，并最终利用全部 8 个字；然而，当多个线程束同时在同一个多处理器上执行时（通常情况如此），缓存行很可能在迭代 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中。

表2 优化C = AB矩阵乘法的性能提升

优化	NVIDIA Tesla V100
无优化	119.9 GB/s
通过共享内存存储A的一个分片来实现合并访问	144.4 GB/s
使用共享内存来消除对B矩阵分块的冗余读取。	195.5 GB/s

ℹ️ 注意

请注意，在通过将额外数据读入共享内存进行优化的方案中，读取 B 数据块后需要调用 __syncthreads()，因为一个线程束会读取由其他线程束写入共享内存的数据。此例程在 NVIDIA Tesla V100 上的有效带宽为 195.5 GB/s。请注意，性能的提升并非由于任何一种情况下的合并访问得到改善，而是因为避免了从全局内存的冗余传输。

中等优先级：使用共享内存以避免从全局内存进行冗余传输。

10.2.3.3. 矩阵乘法中的共享内存C=AAT

先前矩阵乘法的一个变体可用于说明如何处理对全局内存的跨步访问以及共享内存的存储体冲突。该变体简单地使用A的转置代替B，因此C = AA^T。

一个简单的 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;
}

在上述示例中，C矩阵的第row行、第col列元素通过取A矩阵第row行与第col行的点积获得。该内核在NVIDIA Tesla V100上的有效带宽为12.8 GB/s。这些结果显著低于C = AB内核的对应测量值。差异在于半线程束中的线程在每次迭代i时，对第二项a[col*TILE_DIM+i]中A矩阵元素的访问方式。对于一个线程束而言，col代表A矩阵转置的连续列，因此col*TILE_DIM表示以步长w对全局内存进行跨步访问，从而导致大量带宽浪费。

避免跨步访问的方法与之前一样使用共享内存，但在此情况下，一个线程束将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;
}

通过使用全局内存的合并读取来优化跨步访问，利用共享transposedTile以避免点积第二项中的非合并访问，并采用先前示例中的共享aTile技术以避免第一项中的非合并访问。该内核在NVIDIA Tesla V100上的有效带宽为140.2 GB/s。这些结果低于C = AB的最终内核所达到的数值。差异的原因在于共享内存的存储体冲突。

在for循环中对transposedTile内元素的读取是无冲突的，因为每个半线程束的线程按瓦片的行方向读取，从而实现了跨存储体的单位步长。然而，当将瓦片从全局内存复制到共享内存时会发生存储体冲突。为使全局内存的加载能够合并访问，数据需按顺序从全局内存读取。但这要求按列写入共享内存，并且由于共享内存中使用了w×w瓦片，这导致线程间产生w个存储体的步长——线程束中的每个线程都会访问同一存储体（请注意w被选为32）。这种多路存储体冲突的代价非常高昂。简单的解决方案是对共享内存数组进行填充，使其增加一个额外的列，如下列代码所示。

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

这种填充方式完全消除了冲突，因为现在线程间的步幅为w+1个存储体（即当前设备上的33个存储体），由于计算存储体索引时采用取模运算，这等效于单位步幅。经过此修改后，在NVIDIA Tesla V100上的有效带宽达到199.4 GB/s，这与上一个C = AB内核的结果相当。

这些优化结果总结于表3。

表3 优化C = AA^T矩阵乘法的性能提升

优化	NVIDIA Tesla V100
无优化	12.8 GB/s
使用共享内存来合并全局内存读取	140.2 GB/s
消除存储体冲突	199.4 GB/s

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

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

• 为了实现对全局内存的合并访问，特别是为了避免大跨度访问（对于通用矩阵，跨度通常远大于32）
• 为了消除（或减少）来自全局内存的冗余加载
• 为避免带宽浪费

10.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所示。

我们使用每个线程处理4B、8B和16B大小的元素来评估两个内核的性能，即分别使用int、int2和int4作为模板参数。我们调整内核中的copy_count，使得每个线程块复制的数据量从512字节到48 MB不等。内核的性能如图14所示。

从性能图表中，可以得出本次实验的以下观察结果。

• 当所有三种元素大小的copy_count参数为4的倍数时，同步复制可获得最佳性能。编译器能够优化成组的4条加载与存储指令。这一点从锯齿状曲线中可以明显看出。
• 异步复制在几乎所有情况下都能实现更好的性能。
• 异步复制操作不要求copy_count参数必须是4的倍数，以便通过编译器优化实现最佳性能。
• 总体而言，使用元素大小为8或16字节的异步拷贝操作时，可获得最佳性能。

10.3. 分配

通过 cudaMalloc() 和 cudaFree() 进行设备内存的分配与释放是开销高昂的操作。建议使用 cudaMallocAsync() 和 cudaFreeAsync() 这类流序池分配器来管理设备内存。

10.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。

11.1. 占用率

在CUDA中，线程指令是按顺序执行的，因此当一个线程束暂停或停顿时，执行其他线程束是隐藏延迟并保持硬件繁忙的唯一方式。因此，与多处理器上活动线程束数量相关的某些指标对于确定硬件保持繁忙的有效性至关重要。这一指标即为占用率。

占用率是指每个多处理器上活跃线程束数量与可能的最大活跃线程束数量之比。（要确定后者，请参阅 deviceQuery CUDA 示例或参考计算能力文档。）理解占用率的另一种视角是：硬件处理线程束的能力中实际被使用的百分比。

更高的占用率并不总是等同于更高的性能——存在一个临界点，超过该点后额外的占用率将不再提升性能。然而，低占用率总会干扰内存延迟的隐藏能力，从而导致性能下降。

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

11.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%占用率，因为每个多处理器只能容纳四个320线程块。此外，寄存器分配会按每线程束向上取整至最接近的256个寄存器。

可用寄存器的数量、每个多处理器上可同时驻留的最大线程数以及寄存器分配粒度因计算能力的不同而异。由于寄存器分配的这些细微差别，以及多处理器的共享内存也在驻留线程块之间进行分区的事实，寄存器使用量与占用率之间的确切关系可能难以确定。nvcc的--ptxas options=v选项详细说明了每个内核每个线程使用的寄存器数量。有关不同计算能力设备的寄存器分配公式，请参阅《CUDA C++编程指南》的硬件多线程部分；有关这些设备上可用寄存器的总数，请参阅《CUDA C++编程指南》的特性与技术规格部分。此外，NVIDIA在Nsight Compute中提供了一个占用率计算器；请参考https://docs.nvidia.com/nsight-compute/NsightCompute/index.html#occupancy-calculator。

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

11.2. 隐藏寄存器依赖关系

寄存器依赖发生在一条指令使用其前一条指令写入寄存器的结果时。在计算能力7.0的设备上，大多数算术指令的延迟通常为4个周期。因此线程在使用算术结果前需等待约4个周期。然而，通过执行其他线程束中的线程，这一延迟可被完全隐藏。详见寄存器章节。

11.3. 线程与线程块启发式策略

每个网格的线程块维度与大小，以及每个线程块的线程维度与大小均是关键因素。这些参数的多维特性便于将多维问题映射至CUDA，但其本身对性能并无影响。因此，本节将讨论规模而非维度。

延迟隐藏与占用率取决于每个多处理器上活跃的线程束数量，这一数量由执行参数及资源（寄存器和共享内存）限制共同隐式决定。选择执行参数实质上是在延迟隐藏（占用率）与资源利用率之间寻求平衡。

选择执行配置参数时应协同进行；然而，某些启发式方法可独立应用于各个参数。在选择第一个执行配置参数——每个网格的线程块数量（即网格规模）时，主要考虑因素是保持整个GPU处于忙碌状态。网格中的线程块数量应大于流多处理器数量，以确保所有流多处理器至少有一个线程块可执行。此外，每个流多处理器应有多个活跃线程块，使得那些未等待__syncthreads()的线程块能够维持硬件持续工作。此建议需考虑资源可用性限制，因此应在第二个执行参数（每个线程块的线程数，即线程块规模）及共享内存使用情况的背景下综合确定。为适配未来设备，每次内核启动的线程块数量应达到数千规模。

在选择线程块大小时，必须记住多个并发线程块可以驻留在一个多处理器上，因此占用率并非仅由线程块大小决定。特别需要注意的是，更大的线程块尺寸并不一定意味着更高的占用率。

如占用率所述，更高的占用率并不总是等同于更好的性能。例如，将占用率从66%提升至100%通常不会带来同比例的性能增长。较低占用率的内核每个线程可用的寄存器数量会比较高占用率的内核更多，这可能减少寄存器溢出到本地内存的情况；特别是在具有高度显式指令级并行（ILP）的情况下，有时甚至可以通过低占用率完全掩盖延迟。

在选择线程块大小时涉及诸多此类因素，不可避免地需要进行一些实验。然而，应遵循以下几条经验法则：

• 每个线程块中的线程数应为线程束大小的整数倍，以避免在未满载的线程束上浪费计算资源，并促进合并访问。
• 每个线程块至少应使用64个线程，且仅当每个多处理器上存在多个并发线程块时适用。
• 每个线程块包含128至256个线程是尝试不同线程块尺寸的良好初始范围。
• 如果延迟影响性能，请为每个多处理器使用多个较小的线程块，而非单个大型线程块。这对于频繁调用 __syncthreads() 的内核尤为有益。

请注意，当线程块分配的寄存器数量超过多处理器可用数量时，内核启动将失败——这与请求过多共享内存或过多线程时的情况相同。

11.4. 共享内存的影响

共享内存在多种场景下都能发挥作用，例如帮助实现或消除对全局内存的冗余访问。然而，它也可能成为占用率的制约因素。在许多情况下，内核所需的共享内存量与所选的线程块大小相关，但线程到共享内存元素的映射并不需要是一对一的。例如，可能希望在内核中使用64x64元素的共享内存数组，但由于每个线程块的最大线程数为1024，无法启动每块包含64x64线程的内核。在这种情况下，可以启动每块包含32x32或64x16线程的内核，让每个线程处理共享内存数组的四个元素。即使不存在每块线程数等限制，采用单线程处理共享内存数组多个元素的方法也可能带来益处。这是因为每个元素共有的某些操作可由线程一次性执行，其开销可分摊到该线程处理的共享内存元素数量上。

一种通过实验确定性能对占用率敏感度的有效方法是调整动态分配的共享内存量，该参数在执行配置的第三个参数中指定。通过简单地增加此参数（无需修改内核），可以有效降低内核的占用率，并测量其对性能的影响。

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

11.6. 多上下文

CUDA工作在一个特定GPU的进程空间内进行，该空间被称为上下文。上下文封装了针对该GPU的内核启动和内存分配，同时支持页表等辅助结构。上下文在CUDA驱动API中是显式存在的，但在CUDA运行时API中则完全隐式，后者会自动创建和管理上下文。

使用CUDA Driver API，一个CUDA应用程序进程可以为给定的GPU创建多个上下文。如果多个CUDA应用程序进程同时访问同一GPU，这几乎总是意味着存在多个上下文，因为上下文与特定的主机进程绑定，除非正在使用多进程服务（Multi-Process Service）。

在给定的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);

12.1. 算术指令

为最大化指令吞吐量，应用程序应：

• 最小化使用低吞吐量的算术指令；这包括在不影响最终结果的情况下，以精度换取速度，例如使用内联函数而非常规函数（内联函数列于《CUDA C++编程指南》中）、使用单精度而非双精度、或将非规格化数刷新为零；
• 最小化由控制流指令引起的线程束分化，详见控制流指令部分。
• 减少指令数量，例如通过尽可能优化掉同步点（如“同步指令”部分所述）或使用受限指针（如《CUDA C++编程指南》所述）。

在本节中，吞吐量以每个多处理器每时钟周期的操作数量给出。对于线程束大小为32的情况，一条指令对应32次操作，因此若N为每时钟周期的操作数量，则指令吞吐量为每时钟周期N/32条指令。

所有吞吐量均针对单个多处理器。必须将其乘以设备中的多处理器数量，才能获得整个设备的总吞吐量。

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

12.2. 内存指令

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

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

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

通过线程调度器可以隐藏大部分全局内存延迟，前提是在等待全局内存访问完成期间有足够多的独立算术指令可供发射。然而，最佳实践仍是尽可能避免访问全局内存。

13.1. 分支与发散

流控制指令（if、switch、do、for、while）可能导致同一线程束内的线程发生分支，即遵循不同的执行路径，从而显著影响指令吞吐量。若发生这种情况，不同的执行路径必须被分别执行；这会增加该线程束执行指令的总数。

在控制流依赖于线程ID的情况下，为获得最佳性能，控制条件的编写应尽可能减少发散线程束的数量。

这是可能的，因为线程束在线程块中的分布是确定性的，正如《CUDA C++编程指南》的SIMT架构部分所述。一个简单的例子是当控制条件仅依赖于（threadIdx / WSIZE）时，其中WSIZE为线程束大小。

在这种情况下，没有线程束发生分支发散，因为控制条件与线程束的边界完全对齐。

对于仅包含少量指令的分支，线程束发散通常只会导致轻微的性能损失。例如，编译器可能会使用谓词化技术来避免实际的分支跳转。取而代之的是，所有指令都被调度执行，但每个线程的条件码或谓词会控制哪些线程执行这些指令。谓词为假的线程不会写入结果，也不会计算地址或读取操作数。

自Volta架构起，独立线程调度允许线程束在数据依赖的条件块之外保持分支执行状态。可通过显式使用__syncwarp()来确保该线程束在后续指令执行前重新收敛。

13.2. 分支预测

有时，编译器可能会通过使用分支预测而非实际分支来展开循环或优化掉 if 或 switch 语句。在这些情况下，任何线程束都不会发生分支。程序员也可以使用

#pragma unroll

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

在使用分支预测时，所有执行依赖于控制条件的指令都不会被跳过。相反，每条此类指令都与一个按线程的条件码或谓词相关联，该谓词根据控制条件被设置为真或假。尽管这些指令都被调度执行，但只有谓词为真的指令才会实际执行。谓词为假的指令不会写入结果，也不会计算地址或读取操作数。

编译器仅在分支条件控制的指令数量小于或等于特定阈值时，才会将分支指令替换为谓词化指令。

15.1. CUDA 计算能力

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

可以通过编程方式查询设备中GPU的计算能力，如deviceQuery CUDA示例所示。该程序的输出如图16所示。这些信息通过调用cudaGetDeviceProperties()并访问其返回结构体中的信息获得。

计算能力的主版本号和次版本号如图16第七行所示。该系统的设备0具有计算能力7.0。

有关各种GPU计算能力的更多详细信息，请参阅《CUDA C++编程指南》中的“支持CUDA的GPU”和“计算能力”章节。特别需要注意的是，开发者应关注设备上的多处理器数量、寄存器数量与可用内存容量，以及设备的任何特殊功能。

15.2. 其他硬件数据

某些硬件特性并未通过计算能力进行描述。例如，无论计算能力如何，大多数（但非全部）GPU都支持内核执行与主机和设备间异步数据传输的重叠操作。在此类情况下，应调用cudaGetDeviceProperties()以确定设备是否支持特定功能。例如，设备属性结构中的asyncEngineCount字段指示是否可能实现内核执行与数据传输的重叠（若支持，则指明可能的并发传输数量）；同样，canMapHostMemory字段指示是否可执行零拷贝数据传输。

15.3. 选择哪个计算能力目标

要针对特定版本的NVIDIA硬件和CUDA软件，请使用-arch、-code和-gencode选项的nvcc。例如，使用线程束洗牌操作的代码必须使用-arch=sm_30（或更高计算能力）进行编译。

请参阅《构建最大兼容性》章节，以进一步了解用于为多代支持CUDA的设备同时构建代码时所使用的标志。

15.4. CUDA Runtime

CUDA软件环境的主机运行时组件仅可由主机函数使用。它提供处理以下功能的函数：

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

与较低级别的CUDA Driver API相比，CUDA Runtime通过提供隐式初始化、上下文管理和设备代码模块管理，极大地简化了设备管理。由nvcc生成的C++主机代码使用CUDA Runtime，因此链接此代码的应用程序将依赖于CUDA Runtime；类似地，任何使用cuBLAS、cuFFT及其他CUDA Toolkit库的代码也将依赖于CUDA Runtime，因为这些库内部均使用该运行时。

构成CUDA运行时API的函数在《CUDA工具包参考手册》中有详细说明。

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

它包含两个主要部分：

• 一个C语言风格的函数接口（cuda_runtime_api.h）。
• 基于C风格函数构建的C++风格便捷封装（cuda_runtime.h）。

有关Runtime API的更多信息，请参阅《CUDA C++编程指南》中的CUDA Runtime部分。

16.1. CUDA 工具包版本管理

自CUDA 11起，工具包版本采用行业标准的语义化版本控制方案：.X.Y.Z，其中：

• .X 代表主版本号——API 已发生变更，二进制兼容性被破坏。
• .Y 代表次要版本 - 引入新的 API，弃用旧的 API，可能会破坏源代码兼容性，但会保持二进制兼容性。
• .Z 代表发布/补丁版本号——新的更新和补丁将递增此数字。

工具包中的每个组件均推荐采用语义化版本控制。自CUDA 11.3起，NVRTC也开始采用语义化版本控制。我们将在文档后续部分对其中部分内容进行说明。工具包中各组件的版本信息可在此表格中查阅。

因此，CUDA平台的兼容性旨在应对以下几种场景：

1. 对于企业或数据中心中运行GPU的生产系统而言，NVIDIA驱动程序的升级可能较为复杂，且可能需要提前规划。推迟部署新的NVIDIA驱动程序可能意味着此类系统的用户无法使用CUDA版本中提供的新功能。若新CUDA版本无需更新驱动程序，则意味着能够更快地向用户提供新版本的软件。
2. 许多基于CUDA构建的软件库和应用程序（例如数学库或深度学习框架）并不直接依赖于CUDA运行时、编译器或驱动程序。在这种情况下，用户或开发者仍能从中获益，无需升级整个CUDA工具包或驱动程序即可使用这些库或框架。
3. 升级依赖项容易出错且耗时，在某些极端情况下甚至可能改变程序的语义。持续使用最新的CUDA Toolkit重新编译，意味着强制应用程序产品的终端客户进行升级。包管理器虽能简化此流程，但意外问题仍可能出现；一旦发现缺陷，就需要重复上述升级过程。

CUDA支持多种兼容性选项：

1. 首次在CUDA 10中引入的CUDA前向兼容升级功能，旨在让用户能够访问新的CUDA特性，并在安装旧版NVIDIA数据中心驱动程序的系统上运行基于新版CUDA构建的应用程序。
2. 首次在CUDA 11.1中引入的CUDA增强兼容性提供两大优势：通过利用CUDA工具包中各组件的语义版本控制，应用程序可针对某一CUDA次要版本（例如11.1）构建，并能在主版本系列（即11.x）的所有未来次要版本中运行。CUDA运行时放宽了最低驱动程序版本检查，因此在升级到新的次要版本时不再需要升级驱动程序。
3. CUDA驱动程序确保为已编译的CUDA应用程序保持向后二进制兼容性。使用旧至3.2版本的CUDA工具包编译的应用程序可在更新的驱动程序上运行。

16.2. 源代码兼容性

我们将源兼容性定义为库提供的一系列保证，即针对特定版本库（使用SDK）构建的正确应用程序，在安装新版本SDK后仍能继续构建和运行而不出错。

CUDA驱动程序和CUDA运行时在不同SDK版本间均不保证源代码兼容性。API可能被弃用或移除。因此，在旧版本工具链上成功编译的应用程序可能需要修改才能适配新版本工具链。

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

16.3. 二进制兼容性

我们将二进制兼容性定义为库提供的一组保证，即针对该库开发的应用程序在动态链接到不同版本的库时仍能继续正常工作。

CUDA Driver API 采用版本化的C风格ABI，这保证了针对旧版驱动程序（例如CUDA 3.2）运行的应用程序仍能在新版驱动程序（例如随CUDA 11.0发布的驱动程序）上正常运行并保持功能正确。这意味着，虽然应用程序若需使用新特性而必须基于新版CUDA Toolkit重新编译时，其源代码可能需要修改，但将系统中安装的驱动程序组件替换为新版本后，新版驱动程序将始终兼容现有应用程序及其功能。

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

在深入探讨此主题之前，开发者理解最低驱动程序版本的概念及其可能产生的影响至关重要。

每个版本的CUDA Toolkit（及其运行时）都需要一个最低版本的NVIDIA驱动程序。针对特定CUDA Toolkit版本编译的应用程序，仅能在满足该工具包版本所要求的最低驱动程序版本的系统上运行。在CUDA 11.0之前，工具包所需的最低驱动程序版本与该版本CUDA Toolkit附带的驱动程序版本相同。

因此，当应用程序使用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工具包发布说明，以了解最低驱动程序版本以及工具包附带的驱动程序版本详情。

16.4. 跨次要版本的CUDA兼容性

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

新工具链要求新最低驱动程序的主要原因之一，是为了处理PTX代码的即时编译以及二进制代码的即时链接。

在本节中，我们将回顾利用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                                                 |
+-----------------------------------------------------------------------------+

$ 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

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

#define MIN_VERSION 11010
cudaError_t foo()
{
    int version = 0;
    cudaGetDriverVersion(&version);
    if (version < MIN_VERSION) {
        return CUDA_ERROR_INSUFFICIENT_DRIVER;
    }
    // proceed as normal
}

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

#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, NULL, 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;
}

17.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版本的详细信息，请参阅版本管理部分。

17.2. 错误处理

所有CUDA运行时API调用均返回类型为cudaError_t的错误代码；若无错误发生，返回值将等于cudaSuccess。（例外情况包括内核启动——其返回void，以及cudaGetErrorString()——该函数返回描述传入的cudaError_t代码的字符串。）CUDA工具包库（cuBLAS、cuFFT等）同样会返回其各自的错误代码集。

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

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

17.3. 构建以实现最大兼容性

每一代支持CUDA的设备都有一个相关的计算能力版本，该版本指示设备支持的功能集（参见CUDA计算能力）。在构建文件时，可以向nvcc编译器指定一个或多个计算能力版本；针对应用程序目标GPU的原生计算能力进行编译至关重要，这能确保应用程序内核实现最佳性能，并能够利用特定代次GPU上可用的功能。

当应用程序为多个计算能力同时构建时（通过向nvcc使用多个-gencode标志实例），指定计算能力的二进制文件会被合并到可执行文件中，CUDA驱动程序在运行时根据当前设备的计算能力选择最合适的二进制文件。若未找到合适的原生二进制文件（cubin），但存在中间PTX代码（针对抽象虚拟指令集，用于向前兼容），则内核将通过即时编译（JIT）从PTX生成设备的原生cubin（参见编译器JIT缓存管理工具）。若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=。

17.4. 分发CUDA运行时和库

CUDA应用程序基于CUDA运行时库构建，该库负责设备、内存和内核管理。与CUDA驱动程序不同，CUDA运行时库不保证跨版本的向前或向后二进制兼容性。因此，在使用动态链接时，最好将CUDA运行时库与应用程序一同分发，或者选择静态链接CUDA运行时库。这将确保即使最终用户未安装与应用程序构建时相同版本的CUDA工具包，可执行文件仍能正常运行。

静态链接的CUDA运行时

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

动态链接的 CUDA 运行时

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

在使用CUDA 5.5或更高版本的nvcc链接应用程序时，若要与CUDA运行时进行动态链接，需在链接命令行中添加--cudart=shared标志；否则默认使用静态链接的CUDA运行时库。

应用程序动态链接至CUDA Runtime后，此版本的运行时库应随应用程序一同分发。可将其复制至应用程序可执行文件所在目录，或安装路径的子目录中。

其他 CUDA 库

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

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

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

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

$ 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
    ...

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

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

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

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

18.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不保证与旧版本向后兼容。

18.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 获取更多信息。

18.3. 集群管理工具

管理您的GPU集群将有助于实现GPU的最大利用率，并帮助您和您的用户获得最佳性能。许多业界最流行的集群管理工具通过NVML支持CUDA GPU。有关部分此类工具的列表，请参阅 https://developer.nvidia.com/cluster-management 。

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

应用程序在运行时加载的任何PTX设备代码，都会由设备驱动程序进一步编译为二进制代码。这一过程称为即时编译（JIT）。即时编译会增加应用程序的加载时间，但能让应用程序受益于最新的编译器改进。同时，这也是让应用程序在编译时尚未存在的设备上运行的唯一途径。

当使用PTX设备代码的即时编译时，NVIDIA驱动程序会将生成的二进制代码缓存至磁盘。此行为的某些方面，例如缓存位置和最大缓存大小，可通过环境变量进行控制；详见《CUDA C++编程指南》中的即时编译章节。

18.5. CUDA_VISIBLE_DEVICES

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

应用程序可见的设备应作为系统范围内可枚举设备列表的逗号分隔列表进行指定。例如，若仅使用系统设备列表中的设备0和设备2，应在启动应用程序前设置CUDA_VISIBLE_DEVICES=0,2。随后，应用程序将把这些设备分别枚举为设备0和设备1。

19.1. 整体性能优化策略

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

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

最大化并行执行始于以尽可能暴露更多并行性的方式构建算法。一旦算法的并行性得以暴露，就需要尽可能高效地将其映射到硬件上。这通过仔细选择每个内核启动的执行配置来实现。应用程序还应在更高层次上最大化并行执行，通过流（streams）显式暴露设备上的并发执行，并最大化主机与设备之间的并发执行。

优化内存使用始于最小化主机与设备间的数据传输，因为这些传输的带宽远低于设备内部的数据传输。内核访问全局内存也应通过最大化设备上共享内存的使用来尽量减少。有时，最佳的优化甚至可能是在一开始就避免任何数据传输，只需在需要时重新计算数据即可。

有效带宽可能因每种内存的访问模式不同而相差一个数量级。因此，优化内存使用的下一步是根据最优内存访问模式来组织内存访问。这一优化对于全局内存访问尤为重要，因为访问延迟可能高达数百个时钟周期。相比之下，共享内存访问通常仅在存在高度存储体冲突时才值得优化。

在优化指令使用方面，应避免使用低吞吐量的算术指令。这意味着在不影响最终结果的前提下，可以用精度换取速度，例如使用内部函数替代常规函数，或采用单精度而非双精度。最后，由于设备的SIMT（单指令多线程）特性，必须特别关注控制流指令。

20.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（精度较低的开方运算）
• nvcc 的 -use_fast_math 编译器选项强制将每个 functionName() 调用转换为等价的 __functionName() 调用。这以降低精度和准确性为代价，使代码运行得更快。请参阅数学库。

21.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 书面批准、未经改动完整复制、完全遵守所有适用的出口法律法规，并附带所有相关条件、限制和声明的情况下进行复制。

本文档及所有英伟达设计规格、参考板、文件、图纸、诊断工具、清单及其他文档（合称及单称“资料”）均“按现状”提供。英伟达对资料不作任何明示或暗示的法定或其他形式的保证，并明确否认所有关于不侵权、适销性及特定用途适用性的暗示保证。在法律允许的最大范围内，英伟达对因使用本文档引起的任何损害（包括但不限于任何直接、间接、特殊、附带、惩罚性或后果性损害，无论其成因及责任认定理论为何）概不负责，即使英伟达已被告知发生此类损害的可能性。尽管客户可能因任何原因遭受任何损害，英伟达就本文所述产品对客户承担的总计及累计责任应受产品销售条款的限制。

21.2. OpenCL

OpenCL 是 Apple Inc. 的商标，经许可由 Khronos Group Inc. 使用。

21.3. 商标

NVIDIA与NVIDIA标识是NVIDIA公司在美国及其他国家的商标或注册商标。其他公司及产品名称可能是其各自关联公司的商标。