可扩展编程模型

多核 CPU 和多核 GPU 的出现意味着主流处理器芯片已成为并行系统。我们面临的挑战是如何开发应用软件，以透明的方式扩展其并行性，从而充分利用不断增加的处理器内核，就像三维图形应用软件以透明的方式扩展其并行性，以适应内核数量千差万别的多核 GPU 一样。CUDA 并行编程模型旨在克服这一挑战，同时为熟悉 C 等标准编程语言的程序员提供较低的学习曲线。

CUDA并行编程的核心是三个关键的抽象概念:

  • 线程组的层次结构 A hierarchy of thread groups
  • 共享内存 Shared memories
  • 障碍同步 Barrier synchronization

这些抽象概念作为一组最小的语言扩展简单地暴露给程序员。

这种可扩展的编程模型使 GPU 架构能够跨越广泛的市场范围，只需简单地调整多处理器和内存分区的数量即可。

核 Kernels

CUDA C++ 对 C++ 语言进行了扩展，允许用户定义被称作 核(Kernel) 的 C++ 函数。当 kernel 被调用时，它会被 $N$ 个 CUDA 线程并行地执行 $N$ 次，而在普通的 C++ 函数 (function) 中只会被执行一次。

Kernel 由 __global__ 声明符定义。对于每个 kernel 的调用，执行该 kernel 的 CUDA 线程数使用语法 <<<...>>> 给出。执行 kernel 的每个线程都有一个唯一的线程 ID，可通过内置变量在 kernel 中访问。举个例子：

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

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

线程层次结构 Thread Hierarchy

对于CUDA的软件架构我们在逻辑上分为三个层次结构。从小到大依次是线程（thread）、线程块（thread block）、网格（grid）。根据 NVidia 提供的 CUDA 编程指导，三者间的关系为：

  • threadIdx 是一个 3 分量向量，因此可以使用一维、二维或三维线程索引来识别线程，形成一个一维、二维或三维线程组成的块，称为线程块 (thread block)。这为调用向量、矩阵或体积等域中的元素进行计算提供了一种自然的方式。线程索引和线程 ID 之间的关系很简单： 对于一维程序块，它们是相同的；对于大小为 $(Dx, Dy)$ 的二维程序块，索引为 $(x, y)$ 的线程的线程 ID 是 $(x + y\cdot Dx)$；对于大小为 $(Dx, Dy, Dz)$ 的三维程序块，索引为 $(x, y, z)$ 的线程的线程 ID 是 $(x + y\cdot Dx + z\cdot Dx\cdot Dy)$。
  • 线程块被组织成一维、二维或三维的线程块网格 (grid)。网格中线程块的数量通常由处理数据的大小决定，而处理数据的大小通常超过系统中处理器的数量。
层次关系


下例是在多个线程块中进行矩阵加法的一段代码：

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

«<…»> 语法中指定的每个网格的块数和每个块的线程数可以是 int 或 dim3 类型。二维区块或网格可按上例指定。

blockIdx是块（block）在整个网格（grid）中的位置，blockDim.x和blockDim.y是块在水平和竖直方向上的线程数，在这个例子中，它们的大小都是16。

线程块簇 Thread Block Clusters

CUDA 编程还引入了一种可选择的层次结构，称为线程块簇（threa block clusters），由线程块组成。集群中的线程块也保证在 GPU 处理集群（GPU Processing Cluster, GPC）上共同调度。

（最后更新于2024/08/14，待更新）