接第一章 CUDA C++ 编程指北-第一章 入门以及编程模型

CUDA的编程接口

CUDA C++为熟悉C++编程语言的用户提供了一个简单的路径，使他们能够轻松地编写供设备执行的程序。

它由C++语言的最小扩展集和一个运行时库组成。

核心语言扩展已在编程模型中介绍。它们允许程序员将内核定义为C++函数，并使用一些新的语法每次调用函数时指定grid and block的维度。所有扩展的完整描述可以在C++语言扩展中[1]找到。包含这些扩展的任何源文件都必须按照NVCC编译中的说明使用nvcc进行编译[2]。

运行时在CUDA运行时[3]中被介绍。它提供在主机上执行的C和C++函数，用于分配和释放设备内存，传输主机内存和设备内存之间的数据，管理具有多个设备的系统等。运行时的完整描述可以在CUDA参考手册中找到。

运行时是基于一个更低级别的C API构建的，即CUDA驱动API，应用程序也可以访问。驱动API通过公开更低级别的概念（如CUDA上下文 – 设备的主机进程的类比，以及CUDA模块 – 设备的动态加载库的类比）提供了额外的控制级别。大多数应用程序不使用驱动API，因为它们不需要这个额外的控制级别，而当使用运行时时，上下文和模块管理是隐式的，从而产生更简洁的代码。由于运行时与驱动API是互操作的，大多数需要一些驱动API功能的应用程序可以默认使用运行时API，并仅在需要时使用驱动API。驱动程序API在Driver API[4]中介绍。

使用NVCC进行编译

内核可以使用CUDA指令集架构PTX编写，该架构在PTX参考手册中有描述。然而，通常更有效的方法是使用高级编程语言如C++。在两种情况下，内核必须由nvcc编译成二进制代码才能在设备上执行。

nvcc是一个编译器驱动程序，简化了编译C++或PTX代码的过程：它提供简单和熟悉的命令行选项，并通过调用实现不同编译阶段的工具集合来执行它们。本节概述了nvcc工作流程和命令选项。完整说明可在nvcc用户手册中找到。

Compilation Workflow

Offline Compilation 离线编译

使用nvcc编译的源文件可以包括host代码（即在主机上执行的代码）和设备代码（即在设备上执行的代码）的混合。nvcc的基本工作流程包括将设备代码与主机代码分开，然后:

• 将设备代码编译成汇编形式（PTX代码）和/或二进制形式（cubin对象）,
• 并通过替换在内核中引入的>>语法（并在执行配置中更详细地描述^[1]）来修改主机代码，以必要的CUDA运行时函数调用从PTX代码和/或cubin对象加载和启动每个编译的内核

修改后的主机代码输出为C++代码，可以使用另一个工具编译，或者直接通过让nvcc在最后的编译阶段调用主机编译器输出为对象代码。

Applications can then:

• 要么链接到编译后的主机代码（这是最常见的情况），
• 要么忽略修改后的主机代码（如果有的话），并使用CUDA驱动API（参见驱动API）加载和执行PTX代码或cubin对象。

Just-in-Time Compilation 即时编译

任何在运行时由应用程序加载的PTX代码都会被设备驱动程序进一步编译成二进制代码。这被称为即时编译。即时编译增加了应用程序的加载时间，但允许应用程序受益于每个新设备驱动程序带来的任何新的编译器改进。这也是应用程序在被编译时尚不存在的设备上运行的唯一方式，如在应用程序兼容性[6]中详细描述。

当设备驱动程序为某个应用程序即时编译一些PTX代码时，它会自动缓存生成的二进制代码的副本，以避免在后续调用应用程序时重复编译。该缓存 – 称为计算缓存 – 在升级设备驱动程序时会自动失效，这样应用程序就可以从新的即时编译器中受益，该编译器内置于设备驱动程序中。

环境变量可用于控制即时编译，如在CUDA环境变量中所述[7]。

作为使用nvcc编译CUDA C++设备代码的替代方法，可以在运行时使用NVRTC编译CUDA C++设备代码到PTX。NVRTC是一个用于CUDA C++的运行时编译库；更多信息可以在NVRTC用户指南中找到。

Binary Compatibility

二进制代码是特定于架构的。使用编译器选项-code 生成_cubin_ 对象，该选项指定目标架构：例如，使用-code=sm_80 编译会为计算能力[8] 8.0的设备生成二进制代码。从一个次要版本到下一个次要版本保证二进制兼容性，但从一个次要版本到前一个次要版本或跨主要版本则不保证。换句话说，为计算能力_X.y_ 生成的_cubin_ 对象只会在计算能力为_X.z_ 的设备上执行，其中_z≥y_ 。

Binary compatibility is supported only for the desktop. It is not supported for Tegra. Also, the binary compatibility between desktop and Tegra is not supported.

PTX Compatibility

某些PTX指令仅在具有更高计算能力的设备上受支持。例如，Warp Shuffle函数仅在计算能力为5.0及以上的设备上受支持。-arch编译器选项指定编译C++到PTX代码时假定的计算能力。因此，包含warp shuffle的代码，例如，必须使用-arch=compute_50（或更高）进行编译。

为某个特定计算能力生成的PTX代码始终可以编译为更大或相等的计算能力的二进制代码。请注意，从早期PTX版本编译的二进制文件可能不会使用某些硬件功能。例如，针对计算能力7.0（Volta）的设备编译的二进制文件，从为计算能力6.0（Pascal）生成的PTX编译，将不会使用Tensor Core指令，因为Pascal上没有这些指令。因此，最终的二进制文件的性能可能比使用PTX的最新版本生成的二进制文件差。

Application Compatibility

为了在具有特定计算能力的设备上执行代码，应用程序必须加载与此计算能力兼容的二进制或PTX代码，如在二进制兼容性[9]和PTX兼容性[10]中所述。特别地，为了能够在具有更高计算能力的未来架构上执行代码（为其尚未生成二进制代码），应用程序必须加载将为这些设备即时编译的PTX代码（参见即时编译[11]）。

CUDA C++应用程序中嵌入的PTX和二进制代码是由-arch和-code编译器选项或-gencode编译器选项控制的，如nvcc用户手册中详细描述的那样。例如，

nvcc x.cu -gencode arch=compute_50,code=sm_50 -gencode arch=compute_60,code=sm_60 -gencode arch=compute_70,code="compute_70,sm_70"

嵌入了与计算能力5.0和6.0兼容的二进制代码（第一和第二个-gencode选项）以及与计算能力7.0兼容的PTX和二进制代码（第三个-gencode选项）。

主机代码在运行时自动生成，以自动选择最适合的代码进行加载和执行，上述示例中将是：

• 5.0的二进制代码用于计算能力为5.0和5.2的设备，
• 6.0的二进制代码用于计算能力为6.0和6.1的设备，
• 7.0的二进制代码用于计算能力为7.0和7.5的设备，
• PTX代码在运行时编译为计算能力为8.0和8.6的设备的二进制代码。

x.cu可以有一个优化的代码路径，例如使用仅在计算能力为8.0及更高版本的设备上受支持的warp reduction操作。__CUDA_ARCH__宏可用于基于计算能力区分各种代码路径。它仅为设备代码定义。例如，当使用-arch=compute_80编译时，__CUDA_ARCH__等于800。

使用驱动API的应用程序必须将代码编译为单独的文件，并在运行时明确加载和执行最合适的文件。

Volta架构引入了独立线程调度，这改变了GPU上线程的调度方式。对于依赖于先前架构中SIMT调度的特定行为的代码，独立线程调度可能会更改参与线程的集合，导致结果不正确。为了在实施独立线程调度中详细描述的纠正措施时帮助迁移，Volta开发人员可以选择使用编译器选项组合-arch=compute_60 -code=sm_70来选择Pascal的线程调度。

nvcc用户手册列出了-arch、-code和-gencode编译器选项的各种简写。例如，-arch=sm_70是-arch=compute_70 -code=compute_70,sm_70的简写（与-gencode arch=compute_70,code=”compute_70,sm_70″相同）。

C++ Compatibility

编译器的前端根据C++语法规则处理CUDA源文件。主机代码支持完整的C++。但是，只有C++语言支持中描述的部分C++代码在device代码中得到完全支持。

64-位兼容

nvcc的64位版本以64位模式编译设备代码（即指针为64位）。以64位模式编译的设备代码只能与以64位模式编译的主机代码一起使用。

CUDA Runtime

运行时在cudart库中实现，该库链接到应用程序，可以通过cudart.lib或libcudart.a静态链接，或通过cudart.dll或libcudart.so动态链接。需要cudart.dll和/或cudart.so进行动态链接的应用程序通常将它们作为应用程序安装包的一部分包含在内。只有在链接到CUDA运行时的同一实例的组件之间传递CUDA运行时符号的地址才是安全的。

它的所有入口点都带有cuda前缀。

• 如异构编程[12]中所述，CUDA编程模型假定一个由主机和设备组成的系统，每个系统都有自己的独立内存。设备内存[13]概述了用于管理设备内存的运行时函数。
• 共享内存[14]说明了如何使用在线程层次结构[15]中引入的共享内存来最大化性能。
• 页锁定的主机内存[16]介绍了与设备内存之间的数据传输重叠的内核执行所需的页锁定的主机内存。
• 异步并发执行[17]描述了用于在系统的各个级别启用异步并发执行的概念和API。
• 多设备系统[18]展示了编程模型如何扩展到与同一主机连接的多个设备的系统。
• 错误检查[19]描述了如何正确检查运行时生成的错误。
• 调用堆栈[20]提到了用于管理CUDA C++调用堆栈的运行时函数。
• 纹理和表面内存[21]介绍了纹理和表面内存空间，这些空间提供了另一种访问设备内存的方法；它们还公开了GPU纹理硬件的一个子集。
• 图形互操作性[22]介绍了运行时提供的与两个主要图形API（OpenGL和Direct3D）互操作的各种函数。

Initialization

从CUDA 12.0开始，cudaInitDevice()和cudaSetDevice()调用会初始化与指定设备关联的运行时和主上下文。如果没有这些调用，运行时将隐式使用设备0，并根据需要自我初始化以处理其他运行时API请求。在计时运行时函数调用和解释第一次调用运行时的错误代码时，需要记住这一点。在12.0之前，cudaSetDevice()不会初始化运行时，应用程序通常会使用无操作运行时调用cudaFree(0)来隔离运行时初始化和其他api活动（为了计时和错误处理）。

运行时为系统中的每个设备创建一个CUDA上下文（有关CUDA上下文的更多详细信息，请参见上下文[23]）。这个上下文是该设备的主上下文，并在第一个需要在此设备上有活动上下文的运行时函数上初始化。它在应用程序的所有主机线程之间共享。作为此上下文创建的一部分，如果需要，设备代码会进行即时编译（参见即时编译[24]）并加载到设备内存中。这一切都是透明的。如果需要，例如，为了驱动API互操作性，可以从驱动API访问设备的主上下文，如运行时和驱动API之间的互操作性[25]中所述。

当主机线程调用cudaDeviceReset()时，这将销毁主机线程当前操作的设备的主上下文（即，在设备选择[26]中定义的当前设备）。任何将此设备作为当前设备的主机线程进行的下一个运行时函数调用都将为此设备创建一个新的主上下文。

CUDA接口使用在主机程序初始化期间初始化并在主机程序终止期间销毁的全局状态。CUDA运行时和驱动程序无法检测此状态是否无效，因此在程序初始化或终止期间（在main之后）使用这些接口（隐式或显式）将导致未定义的行为。

从CUDA 12.0开始，cudaSetDevice()现在将在为主机线程更改当前设备后明确初始化运行时。CUDA的先前版本在cudaSetDevice()之后延迟了新设备上的运行时初始化，直到进行了第一个运行时调用。这一变化意味着现在检查cudaSetDevice()的返回值以查找初始化错误非常重要。

参考手册的错误处理和版本管理部分的运行时函数不会初始化运行时。

Device Memory

如在异构编程[27]中所提及的，CUDA编程模型假设一个由主机和设备组成的系统，每个都有自己的独立内存。内核在设备内存中运行，因此运行时提供了分配、释放和复制设备内存的函数，以及在主机内存和设备内存之间传输数据。

设备内存可以分配为线性内存或CUDA数组。

CUDA数组是为纹理提取优化的不透明内存布局。它们在 Texture and Surface Memory[28]中有描述。

线性内存在单一统一的地址空间中分配，这意味着单独分配的实体可以通过指针相互引用，例如在二叉树或链表中。地址空间的大小取决于主机系统（CPU）和所使用的GPU的计算能力：

在计算能力为5.3（Maxwell）及更早版本的设备上，CUDA驱动程序会创建一个未提交的40位虚拟地址保留，以确保内存分配（指针）位于支持的范围内。这个保留会显示为保留的虚拟内存，但直到程序实际分配内存之前，不会占用任何物理内存。

线性内存通常使用cudaMalloc()进行分配，并使用cudaFree()进行释放，主机内存和设备内存之间的数据传输通常使用cudaMemcpy()完成。在内核的向量加法代码示例中，需要将向量从主机内存复制到设备内存：

// Device code  
__global__ void VecAdd(float* A, float* B, float* C, int N)  
{  
    int i = blockDim.x * blockIdx.x + threadIdx.x;  
    if (i         C[i] = A[i] + B[i];  
}  
  
// Host code  
int main()  
{  
    int N = ...;  
    size_t size = N * sizeof(float);  
  
    // Allocate input vectors h_A and h_B in host memory  
    float* h_A = (float*)malloc(size);  
    float* h_B = (float*)malloc(size);  
    float* h_C = (float*)malloc(size);  
  
    // Initialize input vectors  
    ...  
  
    // Allocate vectors in device memory  
    float* d_A;  
    cudaMalloc(&d_A, size);  
    float* d_B;  
    cudaMalloc(&d_B, size);  
    float* d_C;  
    cudaMalloc(&d_C, size);  
  
    // Copy vectors from host memory to device memory  
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);  
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);  
  
    // Invoke kernel  
    int threadsPerBlock = 256;  
    int blocksPerGrid =  
            (N + threadsPerBlock - 1) / threadsPerBlock;  
    VecAdd>>(d_A, d_B, d_C, N);  
  
    // Copy result from device memory to host memory  
    // h_C contains the result in host memory  
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);  
  
    // Free device memory  
    cudaFree(d_A);  
    cudaFree(d_B);  
    cudaFree(d_C);  
  
    // Free host memory  
    ...  
}  

GPU的硬件实现

NVIDIA GPU架构是围绕可扩展的多线程流式多处理器（SMs）阵列构建的。当主机CPU上的CUDA程序调用一个内核网格时，网格的块会被枚举并分配给具有可用执行能力的多处理器。一个线程块的线程在一个多处理器上并发执行，一个多处理器上可以并发执行多个线程块。当线程块终止时，新的块在空闲的多处理器上启动。

一个多处理器被设计为并发执行数百个线程。为了管理如此大量的线程，它采用了一种称为SIMT（单指令，多线程）的独特架构，该架构在SIMT架构[29]中有描述。指令是流水线式的，利用单个线程内的指令级并行性，以及通过同时的硬件多线程实现广泛的线程级并行性，如硬件多线程[30]中详细描述的那样。与CPU核心不同，它们是按顺序发出的，没有分支预测或推测执行。

SIMT架构[31]和硬件多线程[32]描述了流处理器的架构特性，这些特性对所有设备都是通用的。计算能力5.x[33]、计算能力6.x[34]和计算能力7.x[35]分别为计算能力5.x、6.x和7.x的设备提供了具体细节。

NVIDIA GPU架构使用小端表示法。

SIMT Architecture

多处理器创建、管理、调度并执行称为_warps_的32个并行线程组。组成warp的单个线程在同一程序地址上同时开始，但它们有自己的指令地址计数器和寄存器状态，因此可以自由地分支和独立执行。_warp_一词来源于编织（weaving），编织（weaving）是一种古老的技术，其中多个线程交织在一起创建一个布料。在这里，这个词被用作一个比喻，表示多个线程并行工作，就像在编织中的线程一样。_半warp_是warp的第一半或第二半。_四分之一warp_是warp的第一、第二、第三或第四部分（ A half-warp is either the first or second half of a warp. A quarter-warp is either the first, second, third, or fourth quarter of a warp）。

详细解释下上述这段内容：

• 在CUDA架构中，多处理器是执行线程的硬件单元。这里提到的“warp”是一个包含32个线程的组，这些线程在多处理器上并行执行。
• 尽管所有线程都从同一个程序地址开始，但每个线程都有自己的指令计数器和寄存器状态。这意味着，尽管它们开始时执行相同的指令，但随后可以根据各自的数据和条件独立地分支和执行。
• 一个warp包含32个线程，所以一个“半warp”就是其中的16个线程，可以是前16个或后16个线程。一个warp被分为四个部分，每部分包含8个线程。所以，“四分之一warp”可以是warp的前8个、第9到16个、第17到24个或最后8个线程。

当一个多处理器被赋予一个或多个要执行的线程块时，它将它们划分为warps，每个warp由warp调度器调度执行。一个块被划分为warps的方式总是相同的；每个warp包含连续的、递增的线程ID，第一个warp包含线程0。线程层次结构[36]描述了线程ID如何与块中的线程索引相关。

一个warp一次执行一个公共指令，所以当一个warp的所有32个线程执行路径都一致（没有分支），就会实现完全效率。如果warp的线程通过数据依赖的条件分支发散，warp会执行每个分支路径，禁用不在该路径上的线程。分支发散只在warp内部发生；不同的warps无论它们是否执行公共或不同的代码路径都是独立执行的。

SIMT架构类似于SIMD（单指令，多数据）向量组织，其中一个单独的指令控制多个处理元素。一个关键的区别是SIMD向量组织将SIMD宽度暴露给软件，而SIMT指令指定单个线程的执行和分支行为。与SIMD向量机器相比，SIMT允许程序员为独立的、标量的线程编写线程级并行代码，以及为协调的线程编写数据并行代码。为了正确性的目的，程序员基本上可以忽略SIMT行为；然而，通过确保代码很少需要warp中的线程发散，可以实现大量的性能提升（however, substantial performance improvements can be realized by taking care that the code seldom requires threads in a warp to diverge）。在实践中，这类似于传统代码中缓存行的作用：在设计正确性时可以安全地忽略缓存行大小，但在设计峰值性能时必须考虑代码结构。另一方面，向量架构要求软件将加载合并为向量，并手动管理发散。

详细解释下上述这段：

• 为了确保程序的正确性，程序员不必深入了解SIMT（单指令多线程）的行为。也就是说，即使程序员不完全理解SIMT的工作原理，他们编写的代码仍然可以正确运行
• 在GPU中，一个warp是一组线程，这些线程同时执行相同的指令。但是，如果某些线程需要执行不同的指令（例如，由于条件分支），那么这些线程就会“发散”。当warp中的线程发散时，它们不能同时执行，这会降低性能。因此，为了实现最佳性能，程序员应尽量避免这种发散
• 传统的CPU编程中，缓存行是内存的一个块，它在物理上连续，并且作为一个整体被加载到缓存中。程序员在编写代码时，为了确保程序的正确性，不需要考虑缓存行的大小。但是，为了获得最佳性能，他们需要考虑如何组织数据，以便有效地利用缓存行，减少缓存未命中
• 向量架构，如SIMD（单指令多数据）架构，要求程序员明确地管理数据的并行性。在这种架构中，数据被组织成向量，并且一次操作可以同时应用于整个向量。为了获得最佳性能，程序员需要确保数据加载是“合并”的，这意味着数据应该在内存中连续存放。此外，程序员还需要手动管理线程或数据元素之间的发散

在NVIDIA Volta之前，warps使用一个由warp中的所有32个线程共享的单一程序计数器，以及一个指定warp的活动线程的活动掩码。因此，来自同一warp的处于发散区域或不同执行状态的线程不能相互发送信号或交换数据，而且需要锁或互斥体保护的数据的细粒度共享的算法很容易导致死锁，这取决于争用线程来自哪个warp。

从NVIDIA Volta架构开始，独立线程调度允许线程之间完全并发，无论warp如何。通过独立线程调度，GPU为每个线程维护执行状态，包括程序计数器和调用堆栈，并可以以每个线程的粒度产生执行，要么更好地利用执行资源，要么允许一个线程等待另一个线程产生的数据。一个调度优化器确定如何将来自同一warp的活动线程组合成SIMT单元。这保留了与之前的NVIDIA GPU中的SIMT执行相同的高吞吐量，但具有更大的灵活性：线程现在可以在子warp粒度上发散和重新汇合。

独立线程调度可能导致参与执行代码的线程集合与开发人员对先前硬件架构的warp同步性[37]做出的假设大不相同。特别是，任何warp同步的代码（例如无同步、warp内规约）都应重新检查，以确保与NVIDIA Volta及更高版本的兼容性。有关更多详细信息，请参阅Compute Capability 7.x[38]。

参与当前指令的warp的线程被称为活动线程，而不在当前指令上的线程是非活动的（已禁用）。线程可能因各种原因处于非活动状态，包括比其warp的其他线程更早地退出，采取与warp当前执行的分支路径不同的分支路径，或者是线程数不是warp大小的倍数的块的最后线程。如果由warp执行的非原子指令为warp的多个线程的相同位置在全局或共享内存中写入，那么发生在该位置的序列化写入的数量取决于设备的计算能力（参见计算能力5.x、计算能力6.x和计算能力7.x），以及执行最终写入的线程是未定义的。如果由warp执行的原子指令为warp的多个线程在全局内存中的相同位置读取、修改和写入，那么对该位置的每个读/修改/写都会发生，它们都会被序列化，但它们发生的顺序是未定义的。

Hardware Multithreading

由多处理器处理的每个warp的执行上下文（程序计数器、寄存器等）在warp的整个生命周期中都在芯片上维护。因此，从一个执行上下文切换到另一个执行上下文没有成本，每次指令发出时，warp调度器选择一个有线程准备执行其下一指令的warp（warp的活动线程[39]）并向这些线程发出指令。

特别地，每个多处理器都有一组32位寄存器，这些寄存器在warp之间进行划分，以及一个并行数据缓存或共享内存，该内存在线程块之间进行划分。

对于给定的内核，可以在多处理器上驻留和一起处理的块和warp的数量取决于内核使用的寄存器和共享内存的数量以及多处理器上可用的寄存器和共享内存的数量。每个多处理器还有一个最大的驻留块数和一个最大的驻留warp数。这些限制以及多处理器上可用的寄存器和共享内存的数量是设备的计算能力的函数，并在计算能力[40]中给出。如果每个多处理器没有足够的寄存器或共享内存来处理至少一个块，内核将无法启动。

一个块中的warp总数如下：

• T 是每个块的线程数，
• Wsize 是warp大小，等于32，
• ceil(x, y) 等于x四舍五入到最接近的y的倍数。

为块分配的寄存器总数和共享内存总量在CUDA工具包中提供的CUDA占用率计算器中有记录。

术语warp-synchronous指的是隐式地假设在同一个warp中的线程在每个指令上都是同步的代码。

参考

• https://docs.nvidia.com/cuda/cuda-c-programming-guide/#hardware-implementation

参考资料

[1]
C++语言扩展中: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#c-language-extensions

[2]
使用nvcc进行编译: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compilation-with-nvcc

[3]
CUDA运行时: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#cuda-c-runtime

[4]
Driver API: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#driver-api

[5]
执行配置中更详细地描述: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#execution-configuration

[6]
应用程序兼容性: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#application-compatibility

[7]
CUDA环境变量中所述: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#env-vars

[8]
计算能力: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capability

[9]
二进制兼容性: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#binary-compatibility

[10]
PTX兼容性: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#ptx-compatibility

[11]
即时编译: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#just-in-time-compilation

[12]
异构编程: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#heterogeneous-programming

[13]
设备内存: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#device-memory

[14]
共享内存: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#shared-memory

[15]
线程层次结构: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#thread-hierarchy

[16]
页锁定的主机内存: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#page-locked-host-memory

[17]
异步并发执行: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#asynchronous-concurrent-execution

[18]
多设备系统: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#multi-device-system

[19]
错误检查: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#error-checking

[20]
调用堆栈: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#call-stack

[21]
纹理和表面内存: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#texture-and-surface-memory

[22]
图形互操作性: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#graphics-interoperability

[23]
上下文: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#context

[24]
即时编译: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#just-in-time-compilation

[25]
运行时和驱动API之间的互操作性: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#interoperability-between-runtime-and-driver-apis

[26]
设备选择: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#device-selection

[27]
异构编程: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#heterogeneous-programming

[28]
Texture and Surface Memory: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#texture-and-surface-memory

[29]
SIMT架构: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#simt-architecture

[30]
硬件多线程: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#hardware-multithreading

[31]
SIMT架构: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#simt-architecture

[32]
硬件多线程: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#hardware-multithreading

[33]
计算能力5.x: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capability-5-x

[34]
计算能力6.x: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capability-6-x

[35]
计算能力7.x: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capability-7-x

[36]
线程层次结构: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#thread-hierarchy

[37]
warp同步性: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#fn2

[38]
Compute Capability 7.x: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capability-7-x

[39]
活动线程: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#simt-architecture-notes

[40]
计算能力: https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#compute-capabilities

公众号后台回复“数据集”获取100+深度学习各方向资源整理