2.3 CUDA 编程模型

上一篇中，我们启动了一个 Kernel，并由 GPU 输出了一次“Hello World!”。但 GPU 的优势并不在于只执行一次代码，而在于让大量线程执行相同的 Kernel，并分别处理不同的数据。那么，CUDA 是如何组织这些线程，又怎样让每个线程找到自己负责的数据呢？

CUDA 编程模型

之前我们有所提及，Kernel 中的代码实际由 GPU 线程，也就是 Thread 执行。Thread 是 CUDA 编程模型中最基本的逻辑执行单位，启动 Kernel 后，每个线程都会独立执行一遍核函数中的代码。

多个线程虽然执行相同的 Kernel 代码，但每个线程都拥有自己的线程编号和运行状态，因此可以根据编号处理不同的数据。（这一点很重要！本篇可能涉及不多，但后面的 CUDA 计算基本上都是以此为基础。）

CUDA 程序通常会启动大量线程。为了方便组织和管理，CUDA 不会把所有线程直接堆放在一起，而是先将一组线程组织成一个 Block，也就是线程块。

可以暂时把 Block 理解为一个线程小组，每个小组内的线程重新编号，比如：

Block 0：Thread 0、Thread 1、Thread 2

Block 1：Thread 0、Thread 1、Thread 2

多个 Block 组成一个 Grid（网格），每个 Block 也都有自己的编号。一次 Kernel 会创建一个 Grid，所以可以粗略地理解为它是一次 Kernel 启动所对应的线程组织范围。

那么整个结构就是这样的：一次 Kernel 对应一个 Grid，一个 Grid 包含若干个 Block，一个 Block 又包含若干个线程，它们的结构如图所示：

图中只需了解三者的结构即可，具体的下标、索引这些后面会讲的。

一维线程索引

前面说可以根据线程编号处理不同的数据，但是如果用多个 Block 的话，每个块内线程编号都会重新编号，不方便利用线程编号了（因为不同 Block 内线程号重复），所以有必要建立一个全局的索引，这个需要结合块编号和块内线程编号来确定。

接下来我们全程使用一维的 Grid 和 Block 来完成一次线程索引的计算和打印输出。

以 2 个 Block 和 3 个 Thread 为例，可以很直观看到线程的全局索引，下面结合代码来进一步实现全局索引的打印：

#include <cuda_runtime.h>
#include <stdio.h>

__global__ void checkIndex()
{
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    printf("块编号：%d | 块内线程号：%d | 全局索引号：%d\n", blockIdx.x, threadIdx.x, index);
}

int main()
{
    dim3 block(3); // 每个块 3 个线程
    dim3 grid(2);  // 一共 2 个块

    checkIndex<<<grid, block>>>(); // 启动 GPU 核函数

    cudaDeviceSynchronize();

    return 0;
}

解释一下代码中的一些新语法和全局索引计算方法：

dim3：CUDA 提供的一种数据类型，内部包含 x、y 和 z 三个成员，可以用来描述一维、二维或三维的线程组织结构（可以把 dim3 当成一个装数字的容器，专门用来存 GPU 线程配置的两个关键数值）。这里只使用一维结构，因此只需要关注 x 方向：dim3 block(3) 表示每个块 3 个线程，dim3 grid(2) 表示一共启动 2 个块。

定义完 Grid 和 Block 后，通过上一篇提到的语句启动核函数：Kernel名称<<<Grid大小, Block大小>>>(参数);。尖括号中的第一个参数表示 Grid 的大小，也就是需要启动多少个 Block；第二个参数表示每个 Block 的大小，也就是每个 Block 中包含多少个 Thread。放到我们的示例中就是 checkIndex<<<grid, block>>>();，这样 6 个线程都会独立执行一遍 checkIndex 核函数中的代码。

接下来看核函数，里面使用了三个 CUDA 自动提供的变量：blockIdx.x、blockDim.x、threadIdx.x。它们不需要作为参数传入核函数，每个线程执行时都可以直接使用。

blockIdx.x 表示当前线程所在 Block 在 Grid 中的编号。我们的代码中用了两个 Block，因此其编号分别为 0、1；threadIdx.x 表示当前线程在所属 Block 内部的编号。每个 Block 有 3 个线程，因此每个 Block 内部的线程编号都是 0、1、2；blockDim.x 表示每个 Block 在 x 方向包含多少个线程，这里 blockDim.x = 3。（注意哦，前俩分别是线程和块的编号，最后一个是块大小。）

所以，通过 int index = blockIdx.x * blockDim.x + threadIdx.x; 计算线程全局索引。（结合前面的示意图，这个不难理解。）

代码运行结果如下：

输出正确，全局索引和块编号及块内线程号一一对应。但是为什么会先输出第 1 块，后输出第 0 块？如果你把参数调得更大，会发现输出更加“乱七八糟”，这并不奇怪，CUDA 不保证不同 Block 的执行顺序，它们可以按任意顺序执行，并行串行均有可能。

Block 和 Grid 的大小

前面的示例直接指定了块和线程的数量，这是为了方便我们观察而任意指定的，但是实际应用中，我们的数据量要远大于此。（不然也不会专门用 GPU 来计算吧~）

当前常见的 NVIDIA GPU 上，一个 Block 通常最多只能开 1024 个线程，一般选 128、256 或者 512，并且最好是 32 的倍数。（这个后面讲到 warp 就知道了。）

Grid 的大小是由所需线程数量和块大小决定，例如我们需要处理 1024 个数据，块大小设置为 256，就需要 1024 / 256 = 4 块，Grid 的大小也就是 4；如果块大小设置为 512，就需要 1024 / 512 = 2 块，Grid 大小就是 2。

不过这也只是我们理想情况下的示例，通常需要处理的元素不一定刚刚好是块大小的倍数，比如需要处理 1000 个数据，每个 Block 包含 256 个线程，那么 3 个 Block 只能提供 768 个线程，无法覆盖全部数据，因此需要向上取整，启动 4 个 Block。

因此，Grid 大小必须进行向上取整。CUDA 程序中经常使用下面的写法：grid.x = (num + block.x - 1) / block.x;，这种写法比较简洁和通用，当然也可以取余后进行判断，看个人习惯。

下面通过一个简单示例观察处理 3000 个数据，不同的块大小所需的数量：

#include <cuda_runtime.h>
#include <iostream>

using namespace std;

// 这里只是计算 Grid 大小，在主函数计算即可，不用核函数。
int main()
{
    int num = 3000; // 总数据量

    dim3 block;
    dim3 grid;

    // 情况 1：每个块开 1024 个线程
    block.x = 1024;
    grid.x = (num + block.x - 1) / block.x;
    cout << "每个块的线程数：" << block.x << " | 总块数：" << grid.x << endl;

    // ================== 情况 2：每个块开 512 个线程 ==================
    block.x = 512;
    grid.x = (num + block.x - 1) / block.x;
    cout << "每个块的线程数：" << block.x << " | 总块数：" << grid.x << endl;

    // ================== 情况 3：每个块开 256 个线程 ==================
    block.x = 256;
    grid.x = (num + block.x - 1) / block.x;
    cout << "每个块的线程数：" << block.x << " | 总块数：" << grid.x << endl;

    // ================== 情况 4：每个块开 128 个线程 ==================
    block.x = 128;
    grid.x = (num + block.x - 1) / block.x;
    cout << "每个块的线程数：" << block.x << " | 总块数：" << grid.x << endl;

    return 0;
}

至此，我们已经了解了 CUDA 中 Grid、Block 和 Thread 的三层组织结构，并完成了线程全局索引的计算，但目前我们只是把这些编号打印了出来。下一篇将进一步完成一个完整的向量加法程序，让每个线程根据自己的全局索引找到对应的数组元素，并真正参与并行计算。