CUDA编程模型

有梦想的雨

已于 2023-10-14 16:45:10 修改

阅读量167

点赞数

分类专栏： CUDA手册文章标签：开发语言人工智能 cuda

于 2023-10-14 16:44:38 首次发布

CUDA手册专栏收录该内容

2 篇文章 0 订阅

订阅专栏

编程模型

本章介绍了CUDA编程模型背后的主要概念，概述了它们在C++中的公开方式。

在编程接口中对CUDA C++进行了详尽的描述。

本章和下一章中使用的矢量加法示例的完整代码可以在矢量加法CUDA示例中找到。

2.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.2 线程层次结构（Thread Hierarchy）

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

线程的索引及其线程ID以一种简单的方式相互关联：对于一维块，它们是相同的；对于大小为 $（ D x ， Dy ）$ 的二维块，索引为 $（ x ， y ）$ 的线程的线程 ID 为 $（ x + yD x ）$ ；对于尺寸为 $（ D x ， Dy ， Dz ）$ 的三维块，索引为 $（ x ， y ， z ）$ 的线程的线程ID为 $（ x + yD x + zD x - Dy ）$ 。

例如，以下代码将大小为 $N\times 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个线程。

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

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

线程块网格
<<<…>>>中指定的每个块的线程数和每个网格的块数语法可以是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个线程）的线程块大小虽然在这种情况下是任意的，但却是常见的选择。网格是用足够的块创建的，每个矩阵元素像以前一样有一个线程。为了简单起见，本例假设每个维度中每个网格的线程数可以被该维度中每个块的线程数整除，但事实并非如此。

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

自动可扩展性
块内的线程可以通过一些共享内存共享数据，并通过同步它们的执行来协调内存访问来进行协作。更确切地说，可以通过调用__syncthreads()内部函数来指定内核中的同步点: __syncthreads()充当了一个屏障，块中的所有线程都必须在该屏障处等待，然后才允许任何线程继续。共享内存提供了一个使用共享内存的示例。除了__syncthreads()之外，CooperativeGroupsneneneba API还提供了一组丰富的线程同步原语。

为了高效协作，共享内存应该是每个处理器核心附近的低延迟内存（很像一级缓存），__syncthreads()应该是轻量级的。

2.2.1线程块簇（Thread Block Clusters）

随着NVIDIA Compute Capability 9.0的推出，CUDA编程模型引入了一种可选的层次结构级别，称为线程块簇（Thread Block Clusters），由线程块组成。类似于线程块中的线程被保证在流式多处理器上被共同调度的方式，簇中的线程块也被保证在GPU中的GPU处理簇（GPU Processing Cluster——GPC）上被共同调度。

与线程块类似，集群也被组织为一维、二维或三维，如下图所示。集群中线程块的数量可以由用户定义，CUDA中支持集群中最多8个线程块作为可移植集群大小。请注意，对于太小而无法支持8个多处理器的GPU硬件或MIG配置，最大集群大小将相应减小。这些较小配置以及支持线程块集群大小超过8的较大配置的标识是特定于体系结构的，可以使用cudaOccupancyMaxPotentialClusterSize 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）上共同调度，并允许集群中的线程块使用集群组API cluster.sync()执行硬件支持的同步。集群组还提供成员根据线程数或块数查询集群组大小：使用num_threads()和num_blocks()。集群组中线程或块的等级可以分别使用dim_threads()和dim_blocks() API查询。

属于集群的线程块可以访问分布式共享内存。集群中的线程块能够对分布式共享内存中的任何地址进行读取、写入和执行原子操作。

2.3 内存层次结构（Memory Hierarchy）

CUDA线程在执行过程中可以访问来自多个内存空间的数据，如下图所示。每个线程都有专用的本地内存。每个线程块的共享内存都对该块的所有线程可见，并且与该块具有相同的生存期。线程块集群中的线程块可以对彼此的共享内存执行读、写和原子操作。所有线程都可以访问相同的全局内存。

还有两个额外的只读内存空间可供所有线程访问：常量内存空间和纹理内存空间（texture memory spaces）。全局、常量和纹理内存空间针对不同的内存使用进行了优化。纹理内存还为某些特定的数据格式提供不同的寻址模式以及数据过滤。

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

内存层次结构

2.4 异构编程（Heterogeneous Programming）

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

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

统一内存提供托管内存以桥接主机和设备内存空间。托管内存可以从系统中的所有CPU和GPU访问，作为具有公共地址空间的单个连贯内存映像。这种功能可以实现设备内存的超额预订，并且可以通过消除在主机和设备上显式镜像数据的需要，大大简化移植应用程序的任务。
异构编程

2.5 异步SIMT编程模型（Asynchronous SIMT Programming Model）

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

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

2.5.1 异步操作（Asynchronous Operations）

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

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

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

Thread Scope(线程作用域)	描述
`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 计算能力（Compute Capability）

设备的计算能力由版本号表示，有时也称为“SM版本”。该版本号标识GPU硬件所支持的功能，并且由应用程序在运行时使用该版本号来确定哪些硬件功能和/或指令在当前GPU上可用。

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

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

次要修订号对应于对核心体系结构的增量改进，可能包括新功能。

图灵是计算能力为7.5的设备的架构，是基于Volta架构的增量更新。

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