CUDA C++编程指南(2)——编程模型

AI-安全-功耗 CUBE 博客目录导读

目录

2.1. 内核

2.2. 线程层次结构

2.2.1. 线程块集群

2.3. 内存层次结构

2.4. 异构编程

2.5. 异步SIMT编程模型

2.5.1. 异步操作

2.6. 计算能力

---

本博客通过概述CUDA编程模型在C++中的实现方式，介绍其背后的主要概念。

关于CUDA C++的详细描述，请参阅Programming Interface。

使用的向量加法示例完整代码可在vectorAdd CUDA sample中找到。

2.1. 内核

CUDA C++ 扩展了 C++，允许程序员定义称为内核(kernel)的 C++ 函数。与常规 C++ 函数仅执行一次不同，当调用这些内核时，它们会被 N 个不同的CUDA 线程并行执行 N 次。

内核使用__global__声明说明符定义，针对给定内核调用执行该内核的CUDA线程数量通过新的<<<...>>>执行配置语法指定（参见Execution Configuration）。每个执行内核的线程都会被分配一个唯一的线程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 threadsVecAdd<<<1, N>>>(A, B, C);...
}

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

2.2. 线程层次结构

为了方便起见，threadIdx是一个三维向量，因此可以使用一维、二维或三维线程索引(threadIdx)来标识线程，形成一维、二维或三维的线程组，称为线程块(thread block)。这为跨域元素（如向量、矩阵或体积）调用计算提供了一种自然的方式。

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

例如，以下代码将两个大小为NxN的矩阵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 threadsint numBlocks = 1;dim3 threadsPerBlock(N, N);MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);...
}

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

然而，一个内核可以由多个形状相同的线程块执行，因此线程总数等于每个块的线程数乘以块的数量。

线程块被组织成一维、二维或三维的网格，如下图所示。网格中的线程块数量通常由待处理数据的大小决定，该数量一般会超过系统中处理器的数量。

                                                                    图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 invocationdim3 threadsPerBlock(16, 16);dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);...
}

线程块大小为16x16（256个线程），虽然在此例中是任意选择的，但这是一个常见的选择。网格被创建为包含足够多的块，以便像之前一样每个矩阵元素对应一个线程。为简单起见，此示例假设每个维度中每个网格的线程数能被该维度中每个块的线程数整除，尽管实际情况并非必须如此。

线程块需要独立执行。必须能够以任意顺序、并行或串行执行这些块。这种独立性要求允许线程块以任何顺序调度并在任意数量的核心上运行，使程序员能够编写可随核心数量扩展的代码。

块内的线程可以通过共享内存共享数据，并通过同步执行来协调内存访问，从而实现协作。更准确地说，可以通过调用__syncthreads()内置函数在内核中指定同步点；__syncthreads()充当一个屏障，块中的所有线程都必须在此等待，然后才能继续执行。Shared Memory提供了一个使用共享内存的示例。除了__syncthreads()之外，Cooperative Groups API还提供了一组丰富的线程同步原语。

为了实现高效协作，共享内存应设计为靠近每个处理器核心的低延迟内存（类似于L1缓存），而__syncthreads()则需保持轻量级特性。

2.2.1. 线程块集群

随着NVIDIA Compute Capability 9.0的推出，CUDA编程模型引入了一个称为线程块集群(Thread Block Clusters)的可选层次结构，该结构由线程块组成。类似于线程块中的线程保证在流式多处理器上协同调度，集群中的线程块也保证在GPU处理集群(GPC)上协同调度。

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

                                                                图5 线程块集群网格

【注意】：在使用集群支持启动的内核中，出于兼容性考虑，gridDim变量仍表示线程块数量的维度。可以通过Cluster Group API来查找块在集群中的层级。

可以通过两种方式在内核中启用线程块集群：一种是使用编译时内核属性__cluster_dims__(X,Y,Z)，另一种是使用CUDA内核启动APIcudaLaunchKernelEx。以下示例展示了如何使用编译时内核属性启动集群。使用内核属性设置的集群大小在编译时固定，之后可以使用传统的<<< , >>>语法启动内核。如果内核使用了编译时集群大小，那么在启动内核时无法修改集群大小。

// 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 sizedim3 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-dimensionattribute[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()进行硬件支持的同步。集群组还提供成员函数，分别通过num_threads()和num_blocks() API查询以线程数或块数表示的集群组大小。线程或块在集群组中的索引可以分别使用dim_threads()和dim_blocks() API查询。

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

2.3. 内存层次结构

CUDA线程在执行过程中可以从多个内存空间访问数据，如下图所示。每个线程都有私有的本地内存。每个线程块都有共享内存，该内存对块内的所有线程可见且生命周期与线程块相同。线程块集群中的线程块可以相互对共享内存执行读取、写入和原子操作。所有线程都可以访问相同的全局内存。

所有线程还可以访问两个额外的只读内存空间：常量内存空间和纹理内存空间。全局内存、常量内存和纹理内存空间针对不同的内存使用场景进行了优化（参见Device Memory Accesses）。纹理内存还为某些特定数据格式提供了不同的寻址模式和数据过滤功能（参见Texture and Surface Memory）。

全局内存、常量内存和纹理内存空间在同一应用程序的内核启动过程中是持久存在的。

                                                                图6 内存层次结构

2.4. 异构编程

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

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

统一内存(Unified Memory)提供managed memory来桥接主机和设备内存空间。managed memory可作为具有统一地址空间的单一连贯内存image，供系统中所有CPU和GPU访问。这一功能支持设备内存的超额分配，并通过消除在主机和设备间显式镜像(mirror)数据的需求，能极大简化应用程序移植工作。有关统一内存的介绍，请参阅Unified Memory Programming。

                                                                        图7 异构编程

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

2.5. 异步SIMT编程模型

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

异步编程模型定义了Asynchronous Barrier的行为，用于CUDA线程之间的同步。该模型还解释并定义了如何通过cuda::memcpy_async在GPU进行计算的同时，异步地将数据从全局内存中移动。

2.5.1. 异步操作

异步操作定义为由CUDA线程发起、并由另一个线程异步执行的操作。在编写良好的程序中，一个或多个CUDA线程会与该异步操作进行同步。发起异步操作的CUDA线程不必包含在这些同步线程中。

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

同步对象可以是cuda::barrier或cuda::pipeline。这些对象在Asynchronous Barrier和Asynchronous Data Copies using 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线程将同步。

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

2.6. 计算能力

设备的计算能力由一个版本号表示，有时也称为"SM版本"。这个版本号标识了GPU硬件支持的功能，应用程序在运行时通过它来确定当前GPU可用的硬件特性和指令。

计算能力由主版本号X和次版本号Y组成，表示为X.Y。

具有相同主版本号的设备属于相同的核心架构。主版本号9对应基于NVIDIA Hopper GPU架构的设备，8对应基于NVIDIA Ampere GPU架构的设备，7对应基于Volta架构的设备，6对应基于Pascal架构的设备，5对应基于Maxwell架构的设备，3对应基于Kepler架构的设备。

次版本号对应核心架构的增量改进，可能包括新功能。

Turing是计算能力7.5设备的架构，是基于Volta架构的增量更新版本。

CUDA-Enabled GPUs列出了所有支持CUDA的设备及其计算能力。Compute Capabilities提供了每个计算能力的技术规格。

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

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