CUDA C++ 编程指南系列教程 —— 编程模型

Programming Model
本章通过概述 CUDA 编程模型背后的主要概念如何在 C++ 中公开来介绍它们。
编程接口中给出了 CUDA C++ 的详细描述。
本章和下一章中使用的向量加法示例的完整代码可以在 vectorAdd CUDA 示例中找到。

1. Kernels

CUDA C++ 通过允许程序员定义称为内核的 C++ 函数来扩展 C++，这些函数在调用时由 N 个不同的 CUDA 线程并行执行 N 次，而不是像常规 C++ 函数那样只执行一次。
内核是使用 global 声明说明符定义的，并且使用新的 <<<…>>> 执行配置语法指定针对给定内核调用执行该内核的 CUDA 线程数（请参阅 C++ 语言扩展）。每个执行内核的线程都会被赋予一个唯一的线程 ID，该 ID 可以在内核中通过内置变量进行访问。
作为说明，以下示例代码使用内置变量 threadIdx 将两个大小为 N 的向量 A 和 B 相加，并将结果存储到向量 C 中：

// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}

int main()
{
    ...
    // Kernel invocation with N threads
    VecAdd<<<1, N>>>(A, B, C);
    ...
}

这里，执行 VecAdd() 的 N 个线程中的每一个都执行一次成对加法。

2. Thread Hierarchy-线程层次结构

为了方便起见，threadIdx 是一个3分量向量，因此可以使用一维、二维或三维线程索引来标识线程，形成一维、二维或三维线程块线程，称为线程块。这提供了一种自然的方式来调用域中元素（例如向量、矩阵或体积）的计算。
线程的索引和线程 ID 之间的关系非常简单：对于一维块，它们是相同的；对于一维块，它们是相同的；对于一维块，它们是相同的。对于大小为(Dx, Dy)的二维块，索引为(x, y)的线程的线程ID为(x + y Dx)；对于大小为 (Dx, Dy, Dz) 的三维块，索引为 (x, y, z) 的线程的线程 ID 为 (x + y Dx + z Dx Dy)。
例如，以下代码将两个大小为 $N$ x $N$ 的矩阵 A 和 B 相加，并将结果存储到矩阵 C 中：

// Kernel definition
__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()
{
    ...
    // Kernel invocation with one block of N * N * 1 threads
    int numBlocks = 1;
    dim3 threadsPerBlock(N, N);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

每个块的线程数量是有限的，因为块中的所有线程都应驻留在同一个流式多处理器核心上，并且必须共享该核心的有限内存资源。在当前的 GPU 上，一个线程块最多可以包含 1024 个线程。
然而，一个内核可以由多个形状相同的线程块来执行，因此线程总数等于每个块的线程数乘以块数。
块被组织成一维、二维或三维线程块网格，如图 4 所示。网格中线程块的数量通常由正在处理的数据大小决定，该大小通常超过系统中处理器的数量。
图 4 线程块网格
<<<...>>> 语法中指定的每个块的线程数和每个网格的块数可以是 int 或 dim3 类型。可以如上例所示指定二维块或网格。
网格内的每个块都可以通过一维、二维或三维唯一索引来标识，该索引可通过内置的 blockIdx 变量在内核中访问。线程块的尺寸可以通过内置的 blockDim 变量在内核中访问。
扩展前面的 MatAdd() 示例以处理多个块，代码如下。

// Kernel definition
__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()
{
    ...
    // Kernel invocation
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
    ...
}

16x16（256 个线程）的线程块大小虽然在本例中是任意的，但却是常见的选择。网格是用足够的块创建的，以便像以前一样每个矩阵元素有一个线程。为简单起见，此示例假设每个维度中每个网格的线程数可被该维度中每个块的线程数整除，但情况不一定如此。
线程块需要独立执行：必须能够以任何顺序（并行或串行）执行它们。这种独立性要求允许在任意数量的内核上以任意顺序调度线程块，如图 3 所示，从而使程序员能够编写随内核数量扩展的代码。
块内的线程可以通过某些共享内存共享数据并同步其执行来协调内存访问来进行协作。更准确地说，可以通过调用 __syncthreads() 内部函数来指定内核中的同步点； __syncthreads() 充当屏障，块中的所有线程都必须等待，然后才允许任何线程继续进行。共享内存给出了使用共享内存的示例。除了 __syncthreads() 之外，协作组 API 还提供了一组丰富的线程同步原语。
为了高效合作，共享内存应该是每个处理器核心附近的低延迟内存（很像 L1 缓存），并且 __syncthreads() 应该是轻量级的。

2.1. Thread Block Clusters

随着 NVIDIA 计算能力 9.0 的推出，CUDA 编程模型引入了一个可选的层次结构级别，称为由线程块组成的线程块集群。与如何保证线程块中的线程在流式多处理器上共同调度类似，集群中的线程块也保证在 GPU 中的 GPU 处理集群 (GPC) 上共同调度。
与线程块类似，簇也被组织成一维、二维或三维，如图5所示。簇中线程块的数量可以由用户定义，最多8个线程块在 CUDA 中支持作为可移植集群大小。请注意，在 GPU 硬件或 MIG 配置上太小而无法支持 8 个多处理器时，最大集群大小将相应减小。这些较小配置以及支持超过 8 的线程块簇大小的较大配置的标识是特定于体系结构的，并且可以使用 cudaOccupancyMaxPotentialClusterSize API 进行查询。
图 5 线程块簇网格
Note

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

可以使用 cluster_dims(X,Y,Z) 的编译器时内核属性或使用 CUDA 内核启动 API cudaLaunchKernelEx 在内核中启用线程块集群。下面的示例展示了如何使用编译器时内核属性启动集群。使用内核属性的簇大小在编译时固定，然后可以使用经典的 <<< 、 >>> 启动内核。如果内核使用编译时簇大小，则在启动内核时无法修改簇大小。

// Kernel definition
// 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;
    // Kernel invocation with compile time cluster size
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);

    // The grid dimension is not affected by cluster launch, and is still enumerated
    // using number of blocks.
    // The grid dimension must be a multiple of cluster size.
    cluster_kernel<<<numBlocks, threadsPerBlock>>>(input, output);
}

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

// Kernel definition
// No compile time attribute attached to the kernel
__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);

    // Kernel invocation with runtime cluster size
    {
        cudaLaunchConfig_t config = {0};
        // The grid dimension is not affected by cluster launch, and is still enumerated
        // using number of blocks.
        // The grid dimension should be a multiple of cluster size.
        config.gridDim = numBlocks;
        config.blockDim = threadsPerBlock;

        cudaLaunchAttribute attribute[1];
        attribute[0].id = cudaLaunchAttributeClusterDimension;
        attribute[0].val.clusterDim.x = 2; // Cluster size in X-dimension
        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集群执行硬件支持的同步。同步（）。集群组还提供成员函数，分别使用 num_threads() 和 num_blocks() API 以线程数或块数查询集群组大小。可以分别使用dim_threads()和dim_blocks() API查询集群组中线程或块的等级。
属于集群的线程块可以访问分布式共享内存。集群中的线程块能够对分布式共享内存中的任何地址进行读取、写入和执行原子操作。分布式共享内存给出了在分布式共享内存中执行直方图的示例。

3. Memory Hierarchy-内存层次结构

CUDA 线程在执行期间可以访问多个内存空间中的数据，如图 6 所示。每个线程都有私有本地内存。每个线程块都有对该块的所有线程可见的共享内存，并且与该块具有相同的生命周期。线程块簇中的线程块可以对彼此的共享内存执行读、写和原子操作。所有线程都可以访问相同的全局内存。
还有两个可供所有线程访问的附加只读内存空间：常量内存空间和纹理内存空间。全局、常量和纹理内存空间针对不同的内存使用情况进行了优化（请参阅设备内存访问）。纹理内存还为某些特定的数据格式提供不同的寻址模式以及数据过滤（请参阅纹理和表面内存）。
全局、常量和纹理内存空间在同一应用程序的内核启动过程中是持久的。
图 6 内存层次结构

4. Heterogeneous Programming-异构编程

如图 7 所示，CUDA 编程模型假设 CUDA 线程在物理上独立的设备上执行，该设备作为运行 C++ 程序的主机的协处理器运行。例如，当内核在 GPU 上执行而 C++ 程序的其余部分在 CPU 上执行时，就会出现这种情况。
CUDA 编程模型还假设主机和设备都在 DRAM 中维护自己独立的内存空间，分别称为主机内存和设备内存。因此，程序通过调用 CUDA 运行时（在编程接口中描述）来管理内核可见的全局、常量和纹理内存空间。这包括设备内存分配和释放以及主机和设备内存之间的数据传输。
统一内存提供托管内存来桥接主机和设备内存空间。托管内存可作为具有公共地址空间的单个一致内存映像从系统中的所有 CPU 和 GPU 进行访问。此功能可实现设备内存的超额订阅，并且无需在主机和设备上显式镜像数据，从而大大简化移植应用程序的任务。有关统一内存的介绍，请参阅统一内存编程。

图 7 异构编程
Note

串行代码在主机(host)上执行，而并行代码在设备(device)上执行。

5. Asynchronous SIMT Programming Model-异步SIMT编程模型

在 CUDA 编程模型中，线程是执行计算或内存操作的最低抽象级别。从基于 NVIDIA Ampere GPU 架构的设备开始，CUDA 编程模型通过异步编程模型提供内存操作加速。异步编程模型定义了与 CUDA 线程相关的异步操作的行为。
异步编程模型定义了用于 CUDA 线程之间同步的异步屏障的行为。该模型还解释并定义了如何使用 cuda::memcpy_async 在 GPU 中计算时从全局内存异步移动数据。

5.1. Asynchronous Operations

异步操作被定义为由 CUDA 线程发起并像由另一个线程一样异步执行的操作。在格式良好的程序中，一个或多个 CUDA 线程与异步操作同步。启动异步操作的 CUDA 线程不需要位于同步线程中。
这样的异步线程（as-if 线程）始终与启动异步操作的 CUDA 线程相关联。异步操作使用同步对象来同步操作的完成。这样的同步对象可以由用户显式管理（例如，cuda::memcpy_async）或在库内隐​​式管理（例如，cooperative_groups::memcpy_async）。
同步对象可以是 cuda::barrier 或 cuda::pipeline。这些对象在使用 cuda::pipeline 的异步屏障和异步数据副本中详细解释。这些同步对象可以在不同的线程范围内使用。范围定义了可以使用同步对象来与异步操作同步的线程集。下表定义了 CUDA C++ 中可用的线程范围以及可以与每个线程同步的线程。

Thread Scope	Description
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++ 的扩展来实现。

6. Compute Capability

设备的计算能力由版本号表示，有时也称为“SM 版本”。该版本号标识 GPU 硬件支持的功能，并由应用程序在运行时使用来确定当前 GPU 上可用的硬件功能和/或指令。
计算能力包括主要修订号X和次要修订号Y并且用X.Y表示。
具有相同主修订号的设备具有相同的核心架构。对于基于 NVIDIA Hopper GPU 架构的设备，主要修订号为 9；对于基于 NVIDIA Ampere GPU 架构的设备，主要修订号为 8；对于基于 Volta 架构的设备，主要修订号为 7；对于基于 Pascal 架构的设备，主要修订号为 6；对于基于 NVIDIA GPU 架构的设备，主要修订号为 5。 Maxwell 架构，3 表示基于开普勒架构的设备。
次要修订号对应于核心架构的增量改进，可能包括新功能。
Turing是计算能力7.5设备的架构，是基于Volta架构的增量更新。
支持 CUDA 的 GPU 列出了所有支持 CUDA 的设备及其计算能力。计算能力给出了每种计算能力的技术规格。
Note

特定 GPU 的计算能力版本不应与 CUDA 版本（例如 CUDA 7.5、CUDA 8、CUDA 9）混淆，后者是 CUDA 软件平台的版本。应用程序开发人员使用 CUDA 平台来创建在多代 GPU 架构上运行的应用程序，包括尚未发明的未来 GPU 架构。虽然 CUDA 平台的新版本通常通过支持新 GPU 架构的计算能力版本来添加对该架构的本机支持，但 CUDA 平台的新版本通常还包括独立于硬件生成的软件功能。

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