CUDA 12.4文档2 内核线程架构

本博客参考官方文档进行介绍，全网仅此一家进行中文翻译，走过路过不要错过。

官方网址：https://docs.nvidia.com/cuda/cuda-c-programming-guide/

本文档分成多个博客进行介绍，在本人专栏中含有所有内容：

https://blog.csdn.net/qq_33345365/category_12610860.html

CUDA 12.4为2024年3月2日发表，本专栏开始书写日期2024/4/8，当时最新版本4.1

本人会维护一个总版本，一个小章节的版本，总版本会持续更新，小版本会及时的调整错误和不合理的翻译，内容大部分使用chatGPT 4翻译，部分内容人工调整

开始编辑时间：2024/4/8；最后编辑时间：2024/4/10

5.1 内核Kernels

CUDA C++通过允许程序员定义C++函数，称为内核，当被调用时，这些函数由N个不同的CUDA线程并行执行N次，而不是像常规的C++函数那样只执行一次。

内核使用 global 声明说明符定义，并且给定内核调用的执行该内核的CUDA线程数使用新的 <<<...>>>执行配置语法指定。执行内核的每个线程都被赋予一个在内核内通过内置变量可以访问的唯一线程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()
{
    ...
    ∕∕ Kernel invocation with N threads
    VecAdd<<<1, N>>>(A, B, C);
    ...
}

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

5.2 线程架构

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

线程的索引和线程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中：

__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()示例扩展为处理多个块的情况：

// 内核定义
__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()预计是轻量级的。

5.2.1 线程块集群 Thread Block Clusters

设计线程数：共享内存上同步 < 线程块集群内同步 < 全局同步

随着NVIDIA计算能力9.0的引入，CUDA编程模型引入了一个叫做线程块群集的可选等级层次，这些都是由线程块构成的。与线程块中的线程保证在流式多处理器上被并行调度类似，线程块群集中的线程块也保证在GPU处理集群（GPC）上进行并行调度。

与线程块类似，群集也以一维、二维或三维的方式组织，如图5所示。群集中的线程块数量可以由用户定义，CUDA中支持以8个线程块为单位的群集大小作为最大限制。注意，在GPU硬件或MIG配置中，如果太小以致不能支持8个多处理器，那么最大群集大小将相应减小。识别这些较小的配置，以及支持线程块群集大小超过8的较大配置，是架构特定的，并可以使用cudaOccupancyMaxPotentialClusterSize API进行查询。

在这里插入图片描述

图5：线程块集群的网格

在使用群集支持启动的内核中，为了兼容性，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;
    ∕∕ 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启动一个群集内核。

∕∕ 内核定义
∕∕ 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 cluster.sync()执行硬件支持的同步。群集组还提供成员函数，使用num_threads()和num_blocks()API分别查询群集组大小，以线程数量或块数量表示。群集组中的线程或块的排名可以使用dim_threads()和dim_blocks()API分别进行查询。

属于群集的线程块可以访问分布式共享内存。群集中的线程块具有读取、写入和执行分布式共享内存中任何地址的原子操作的能力。分布式共享内存章节给出了在分布式共享内存中执行直方图的示例。