(二)CUDA C++ 编程指南-编程模型

参考
CUDA C++ Programming Guide

这一章介绍了CUDA编程模型背后的主要概念，并概述了这些概念如何在C++中得以体现。
CUDA C++的详细描述可以在编程接口中找到。
本章和下一章所使用的向量加法示例的完整代码可以在 vectorAdd CUDA sample示例中找到。

2.1. 核函数（CUDA C++函数）

CUDA C++ 扩展了 C++，它允许程序员定义 C++ 函数，这些函数被称为 kernels（核函数）。与常规的 C++ 函数只执行一次不同，这些核函数在被调用时，会由 N 个不同的 CUDA 线程 并行执行 N 次。

核函数是使用 __ global__ 声明说明符定义的，并且该核函数的 CUDA 线程数是使用<<<…>>>执行配置语法（请参阅C++语言扩展）。每个执行内核函数的线程都有一个唯一的线程 ID，可通过内置变量在内核函数中访问该 ID。

内核使用 global 声明修饰符进行定义，而特定内核调用执行该内核的 CUDA 线程数则通过新的 <<<…>>> 执行配置 语法进行指定的（参见 C++ 语言扩展）。每个执行核函数的线程都会被赋予一个唯一的 线程 ID，这个 ID 可以通过核函数内的内置变量进行访问。

例如，以下代码使用内置变量 threadIdx，将两个大小为 N 的向量 A 和 B 相加，并将结果存储到向量 C 中：

// 核函数定义
__global__ void VecAdd(float* A, float* B, float* C){
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}
int main(){
    ...
    // N 个线程的核函数调用  
	VecAdd<<<1, N>>>(A, B, C);
    ...
}  

在这里，执行 VecAdd() 的每个线程都执行一对元素的加法操作。

2.2. 线程层次结构

为了方便起见，threadIdx 是一个三组件向量，因此可以使用一维、二维或三维的线程索引来标识线程，这些线程形成一个一维、二维或三维的线程块，称为线程块。这为在向量、矩阵或体积等域中的元素上调用计算提供了一种自然的方式。

线程索引和它的线程ID 之间有着直接的关系：对于一维线程块，它们是相同的；对于大小为 (Dx, Dy) 的二维线程块，索引为 (x, y) 的线程的线程 ID 是 (x + yDx)；对于大小为 (Dx, Dy, Dz) 的三维线程块，索引为 (x, y, z) 的线程的线程 ID 是 (x + yDx + zDxDy)。

例如，以下代码将两个大小为 NxN 的矩阵 A 和 B 相加，并将结果存储到矩阵 C 中：

// 核函数定义
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
{
    int i = threadIdx.x;
    int j = threadIdx.y;
    C[i][j] = A[i][j] + B[i][j];
}
int main()
{
    ...
    // 使用一个 N * N * 1 线程块的核函数调用（所有线程都在一个块中）   
	int numBlocks = 1;
    dim3 threadsPerBlock(N, N);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

每个线程块的线程数是有限制的，因为块的所有线程都应驻留在同一个流式多处理器内核上，并且必须共享该内核的有限内存资源。在当前的 GPU 上，一个线程块最多可以包含 1024 个线程。

但是，一个核函数可以由多个形状相等的线程块执行，因此线程总数=每个块的线程数X块数。

线程块被组织成一维、二维或三维的线程块的网格，如图 4 所示。网格中的线程块个数通常由正在处理的数据的大小决定，这通常超过系统中的处理器数。

图 4 线程块中的网格

语法 <<<...>>> 中指定的每个块的线程数和每个网格的块数可以是 int类型 或 dim3 类型。二维块或网格可以按照上面的示例进行指定。

网格中的每个线程块都可以通过一维、二维或三维的唯一索引进行识别，这些索引在核函数中可以通过内置变量 blockIdx 访问。线程块的维度在核函数中可以通过内置变量 blockDim 访问。

扩展之前的 MatAdd() 示例以处理多个块,代码如下所示。

// 核函数定义
__global__ void MatAdd(float A[N][N], float B[N][N], float C[N][N]){
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        C[i][j] = A[i][j] + B[i][j];
}

int main(){
    ...
    // 核函数调用（线程分在多个块中）  
	dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N/threadsPerBlock.x, N/threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

线程块大小为16x16（256个线程），尽管在这个示例中是任意的，但这是一个常见的选择。网格创建时，包含足够多的块，以便每个矩阵元素都有一个线程，就像之前一样。为了简化，本例假设每个维度中每个网格的线程数能够被该维度中每个块的线程数整除，尽管实际情况并不一定如此。

线程块需要独立执行：它们可以以任何顺序、并行或串行的方式执行。这种独立性要求允许线程块在任何数量的内核上以任何顺序进行调度，如下图所示，这使得程序员能够编写与内核数量成比例扩展的代码。

一个块内的线程可以通过共享一些共享内存来合作，并通过同步它们的执行来协调内存访问。更确切地说，程序员可以通过调用内置函数__syncthreads()在内核中指定同步点；__syncthreads()作为一个屏障，块中的所有线程在继续执行之前都必须等待。共享内存提供了一个使用共享内存的例子。除了__syncthreads()外，合作组API还提供了一套丰富的线程同步原语。

为了有效合作，期望共享内存是靠近每个处理器核心的低延迟内存（类似于L1缓存），并且期望__syncthreads()是一个轻量级的操作。

2.2.1 线程块集群

随着NVIDIA计算能力9.0的引入，CUDA编程模型引入了一个可选的层次结构，称为线程块集群，它由线程块组成。类似于线程块中的线程被保证在同一个流多处理器上协同调度，集群中的线程块也被保证在GPU中的GPU处理集群（GPC）上协同调度。

与线程块类似，集群也组织成一维、二维或三维，如下图所示。集群中的线程块数量可以由用户定义，CUDA支持集群中最多有8个线程块作为可移植的集群大小。请注意，在GPU硬件或MIG配置太小而无法支持8个流多处理器的情况下，最大集群大小会相应减少。这些较小配置以及支持超过8个线程块集群大小的较大配置的识别是特定于架构的，可以使用cudaOccupancyMaxPotentialClusterSize API进行查询。

图 线程块集群网格

！注意

在使用集群支持启动的内核中，出于兼容性的目的，gridDim变量仍然表示线程块的数量。块在集群中的排名可以通过Cluster Group API找到。

简单来说，当你在使用CUDA编程时启用了集群支持，gridDim变量还是用来表示线程块的数量，这是为了保持兼容性。而如果你想知道一个线程块在集群中的位置或排名，你可以使用Cluster Group API来查询。

线程块集群可以在内核中通过编译时内核属性 __cluster_dims__(X,Y,Z) 或者使用 CUDA 内核启动 API cudaLaunchKernelEx 来启用。下面的例子展示了如何使用编译时内核属性来启动一个集群。使用内核属性设置的集群大小在编译时是固定的，然后可以使用经典的 <<< , >>> 符号来启动内核。如果一个内核使用了编译时的集群大小，那么在启动内核时，集群大小就无法修改了。

// 核函数定义
// Compile time cluster size 2 in X-dimension and 1 in Y and Z dimension
__global__ void __cluster_dims__(2, 1, 1) cluster_kernel(float *input, float* output)
{

}
int main()
{
    float *input, *output;
    // 使用编译时集群大小的核函数调用
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N/threadsPerBlock.x, N/threadsPerBlock.y);
    // 集群启动不会影响网格维度，网格维度仍然是通过线程块的数量来计算的。
    // 网格维度必须是集群大小的倍数。
    cluster_kernel<<<numBlocks, threadsPerBlock>>>(input, output);
}

线程块集群大小也可以在运行时设置，并且可以使用 CUDA 核函数启动 API cudaLaunchKernelEx 来启动内核。下面的代码示例展示了如何使用可扩展 API 来启动集群内核。

// 核函数定义
// 没有将编译时属性附加到核函数上
__global__ void cluster_kernel(float *input, float* output)
{

}
int main()
{
    float *input, *output;
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N/threadsPerBlock.x, N/threadsPerBlock.y);

    // 使用运行时集群大小的核函数调用
    {
        cudaLaunchConfig_t config = {0};
        // 集群的启动不会影响网格维度，网格维度仍然是通过线程块的数量来计算的。
		// 网格维度应该是集群大小的倍数。
        config.gridDim = numBlocks;
        config.blockDim = threadsPerBlock;

        cudaLaunchAttribute attribute[1];
        attribute[0].id = cudaLaunchAttributeClusterDimension;
        attribute[0].val.clusterDim.x = 2; // X维度上集群大小
        attribute[0].val.clusterDim.y = 1;
        attribute[0].val.clusterDim.z = 1;
        config.attrs = attribute;
        config.numAttrs = 1;

        cudaLaunchKernelEx(&config, cluster_kernel, input, output);
    }
}

在计算能力为 9.0 的 GPU 中，集群中的所有线程块都保证会被共同调度到一个单一的 GPU 处理集群（GPC）上，并且允许集群中的线程块使用 Cluster Group API 中的 cluster.sync() 函数来执行硬件支持的同步操作。Cluster Group 还提供了成员函数，可以使用 num_threads() 和 num_blocks() API 分别查询集群组的大小（以线程数量或线程块数量为单位）。集群组中线程或线程块的排名可以通过 dim_threads() 和 dim_blocks() API 分别进行查询。

属于集群的线程块可以访问分布式共享内存。集群中的线程块能够读取、写入分布式共享内存中的任何地址，并执行原子操作。分布式共享内存提供了一个在分布式共享内存中执行直方图计算的示例。

2.3 内存等级

CUDA线程在执行过程中可以访问多个内存空间，这一点可以通过图6来直观展示。每个线程都有私有的本地内存。每个线程块都有共享内存，该内存对所有属于该线程块的线程都是可见的，并且其生命周期与线程块相同。位于线程块集群中的线程块可以对彼此的共享内存执行读取、写入和原子操作。所有线程都可以访问相同的全局内存。

还有两个可供所有线程访问的附加只读内存空间：常量内存空间和纹理内存空间。
全局内存空间、常量内存空间和纹理内存空间针对不同的内存使用情况进行了优化（请参阅设备内存访问）。纹理内存还为某些特定数据格式提供不同的寻址模式以及数据过滤（请参阅纹理和表面内存）。
全局内存、常量内存和纹理内存空间在同一应用程序启动内核时是持久的。

图6 内存等级

2.4. 异构编程

如图7所示，CUDA编程模型假设CUDA线程在物理上独立的设备上执行，该设备作为运行C++程序的主机的协处理器。例如，当内核在GPU上执行，而C++程序的其余部分在CPU上执行时，就会出现这种情况。

CUDA 编程模型还假设主机和设备在 DRAM 中维护自己独立的内存空间，分别称为主机内存和设备内存。因此，程序通过调用 CUDA 运行时（在 编程接口中描述）来管理内核可见的全局、常量和纹理内存空间。这包括设备内存分配和释放，以及主机和设备内存之间的数据传输。

统一内存提供托管内存以桥接主机和设备内存空间。托管内存可从系统中的所有 CPU 和 GPU 作为具有公共地址空间的单个连贯内存映像进行访问。此功能支持设备内存的超额订阅，并且无需在主机和设备上显式镜像数据，从而大大简化移植应用程序的任务。有关统一内存的介绍，请参阅 统一内存编程。

图7 异构编程

注意
串行代码在主机上执行，而并行代码在设备上执行。

2.5.异步SIMT编程模型

在 CUDA 编程模型中，线程是执行计算或内存操作的最低抽象级别。从基于 NVIDIA Ampere GPU 架构的设备开始，CUDA 编程模型通过异步编程模型提供内存操作加速。异步编程模型定义了与 CUDA 线程相关的异步操作的行为。

异步编程模型定义了CUDA线程之间同步的异步屏障的行为。该模型还解释和定义了如何使用cuda::memcpy_async在GPU进行计算的同时，异步地从全局内存移动数据。

2.5.1. 异步操作

异步操作定义为由 CUDA 线程启动并由另一个线程异步执行的操作。在格式良好的程序中，一个或多个 CUDA 线程与异步操作同步。启动异步操作的 CUDA 线程不需要位于同步线程中。

这样的异步线程（as-if 线程）总是与发起异步操作的CUDA线程相关联。异步操作使用同步对象来同步操作的完成。这样的同步对象可以由用户显式管理（例如，cuda::memcpy_async），也可以在一个库内部隐式管理（例如，cooperative_groups::memcpy_async）。

同步对象可以是cuda::barrier或cuda::pipeline。这些对象在异步屏障和使用cuda::pipeline的异步数据复制中有详细的解释。这些同步对象可以在不同的线程作用域中使用。作用域定义了可能使用同步对象来与异步操作同步的线程集合。下面的表格定义了CUDA C++中可用的线程作用域以及可以与每个作用域同步的线程。

线程作用域	描述
cuda::thread_scope::thread_scope_thread	只有启动异步操作的 CUDA 线程才会同步。
cuda::thread_scope::thread_scope_block	与发起线程位于同一线程块内的所有或任何CUDA线程都会同步。
cuda::thread_scope::thread_scope_device	与发起线程位于同一GPU设备内的所有或任何CUDA线程都会同步。
cuda::thread_scope::thread_scope_system	与发起线程位于同一系统中的所有或任何CUDA或CPU线程都会同步。

这些线程作用域作为CUDA 标准 C++库中标准 C++ 的扩展实现。

2.6. 计算能力

设备的计算能力由一个版本号表示，有时也被称为其“SM版本”。该版本号标识了GPU硬件所支持的特性，并在应用程序运行时用于确定当前GPU上可用的硬件特性和/或指令。

计算能力包括主要修订号 X 和次要修订号 Y，并用 X.Y 表示。

具有相同主要修订号的设备具有相同的内核体系结构。基于 NVIDIA Hopper GPU 架构的设备的主要修订号为 9，基于 NVIDIA Ampere GPU 架构的设备为 8，基于 Volta 架构的设备为 7，基于 Pascal 架构的设备为 6，基于 Maxwell 架构的设备为 5，基于 Kepler 架构的设备为 3。

具有相同主要修订号的设备具有相同的核心架构。
基于NVIDIA Hopper GPU架构的设备主要修订号为9，
基于NVIDIA Ampere GPU架构的设备为8，
基于Volta架构的设备为7，
基于Pascal架构的设备为6，
基于Maxwell架构的设备为5，
而基于Kepler架构的设备为3。

次要修订号对应于对核心架构的增量改进，可能包括新功能。
Turing 是计算能力为 7.5 设备的架构，是基于 Volta 架构的增量更新。

支持CUDA的GPU列出了所有支持CUDA的设备及其计算能力。计算能力给出了每种计算能力的技术规格。

注意
特定GPU的计算能力版本不应与CUDA版本（例如CUDA 7.5、CUDA 8、CUDA 9）混淆，CUDA版本是CUDA软件平台的版本。CUDA平台被应用程序开发人员用来创建可以在多代GPU架构上运行的应用程序，包括未来尚未发明的GPU架构。虽然CUDA平台的新版本通常通过支持该架构的计算能力版本来添加对新GPU架构的原生支持，但CUDA平台的新版本通常也包括与硬件无关的软件功能。

分别从 CUDA 7.0 和 CUDA 9.0 开始不再支持 Tesla 和 Fermi 架构。