介绍CUDA编程模型及CUDA线程体系

可伸缩的编程模型

CUDA 编程模型主要有三个关键抽象：层级的线程组，共享内存和栅同步（barrier synchronization）。

这些抽象提供了细粒度的数据并行和线程并行，可以以嵌套在粗粒的数据并行和任务并行中。它们鼓励将问题分解为子问题。每个子问题可以独立的在block threads中并行解决。同时每个子问题分成更细的部分，可以由块中的所有线程并行地合作解决。

这种分解通过允许线程在解决每个子问题时进行协作来保留语言的表达性，同时支持自动可伸缩性。实际上，每个线程块都可以在GPU中任何可用的多处理器上调度，以任何顺序、并发或顺序，因此编译的CUDA程序可以在任意数量的多处理器上执行，如图所示，而且只有运行时系统需要知道物理多处理器的数量。

图1 Automatic Scalability

Note: A GPU is built around an array of Streaming Multiprocessors (SMs) (see Hardware Implementation for more details). A multithreaded program is partitioned into blocks of threads that execute independently from each other, so that a GPU with more multiprocessors will automatically execute the program in less time than a GPU with fewer multiprocessors.

Kernels

CUDA c++ 通过允许程序员定义 c++ 函数( 称为kernel )来扩展 c++，当调用这些函数时，由 N 个不同的 CUDA 线程并行执行 N 次，而不是像常规 c++ 函数那样只执行一次。

使用_ global _ 声明说明符定义内核，并使用新的<<<…>>>执行配置语法(参见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个线程中的每一个都执行一次成对的相加。

Thread Hierarchy

为了方便起见，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中:

// 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个线程 。

但是，一个内核可以由多个形状相同的线程块执行，这样 线程总数就等于每个块的线程数乘以块的数量 。

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

图2 Grid of Thread Blocks

每个块的线程数和每个网格的块数在<<<…>>>语法的类型可以是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个线程)，虽然在本例中是任意的，但却是常见的选择 。网格是用足够的块创建的，每个矩阵元素都有一个线程。

为了简单起见，本示例假设每个维度中每个网格的线程数可以被该维度中每个块的线程数整除，尽管事实并非如此。

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

块中的线程可以通过共享内存共享数据，并通过同步它们的执行来协调内存访问，从而进行协作 。更精确地说，可以通过调用__syncthreads()内部函数来指定内核中的同步点;__syncthreads()充当一个屏障，在允许任何线程继续之前，块中的所有线程都必须等待。除了__syncthreads()之外，Cooperative Groups API还提供了一组丰富的线程同步原语。

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

Thread Block Clusters

随着NVIDIA Compute Capability 9.0的引入 ，CUDA编程模型引入了一个可选的层次结构级别，称为 线程块集群，它由线程块组成 。 与线程块中的线程被保证在流多处理器上同步调度类似，集群中的线程块也被保证在GPU中的GPU处理集群(GPC)上同步调度 。

与线程块类似，集群也被组织成一维、二维或三维，如图3所示。一个集群中的线程块数量可以由用户定义， CUDA支持一个集群中最多8个线程块作为可移植的集群大小 。线程块集群大小是否超过8取决于体系结构，可以使用cudaoccuancymaxpotentialclustersize API进行查询。

图3 Grid of Thread Block Clusters

Note: In a kernel launched using cluster support, the gridDim variable still denotes the size in terms of number of thread blocks , for compatibility purposes. The rank of a block in a cluster can be found using the 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<<

• 运行期指定

线程块集群大小也可以在运行时设置，并且可以使用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);
    cluster_kernel<<

在具有9.0计算能力的GPU中，集群中的所有线程块都被保证在单个GPU处理集群(GPC)上共同调度，并允许集群中的线程块使用cluster Group API cluster.sync()执行硬件支持的同步 。集群组还提供了成员函数，分别使用num_threads()和num_blocks() API根据线程数或块数查询集群组的大小。可以分别通过dim_threads()和dim_blocks() API查询集群组中线程或块的级别。

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