(注:本文大量参考了参考资料的文章和书本,许多话直接就是原话,本文的目的主要是通过整理来理清思路。)

GPU 并行计算

我们说的「GPU 并行计算」实际上指的是 CPU+GPU 的异构计算架构,在这个架构中,CPU 与 GPU 通过 PCIe 总线连接,协同工作,这里 CPU 所在位置被称为主机端(host),GPU 所在位置则被称为设备端(device)。

CPU+GPU 异构计算

由于 GPU 并行计算是一个异构计算,我们需要区分 CPU 和 GPU 的代码。在 CUDA 编程中,我们主要通过函数类型限定词来区别,主要的三个限定词如下:

  1. __global__:核函数(kernel),即真正在 device 的线程中并行执行的函数,一般从 host 异步调用,在 device 并行执行,返回类型为 void。调用时需要使用 <<<grid, block>>> 来指定线程数量和结构(见下文)。在 CUDA 中,每个线程都会分配到一个唯一的 thread ID(线程号),这个 ID 可以通过核函数的内置变量 threadIdx 来获得。
  2. __device__:在 device 调用,在 device 执行。
  3. __host__:在 host 调用,在 host 执行,一般省略不写。

__device__ 与 __host__ 可以一起使用(即该函数在 host 与 device 都会编译),但 __global__ 只能单独使用。

典型的 CUDA 程序执行流程如下:

  1. 分配 host 内存,进行数据初始化;
  2. 分配 device 内存,将数据从 host 拷贝到 device 上;
  3. 调用 CUDA 核函数在 device 上完成指定运算;
  4. 将 device 上的运算结果拷贝到 host 上;
  5. 释放 device 和 host 上分配的内存。

GPU 逻辑架构

层次逻辑架构

GPU(device)的层次逻辑结构(可以对照下面两张图看):

  • 每个 device 包含多个 grid(网格),同一个 grid 上的所有线程共享相同的内存空间。
  • 每个 grid 包含多个 block(线程块),还有 global memory(全局内存)、constant memory(常量内存)、texture memory(纹理内存)。
  • 每个 block 包含很多 thread(线程),还有 shared memory(共享内存)。Shared memory 的「共享」指的是可以被该 block 内的所有线程共享,其生命周期与 block 一致。同一个 block 内部的 thread 可以同步,也可以通过 shared memory 进行通信。
  • 每个 thread 都有自己的 local memory(本地内存),也可以访问 shared memory,还可以访问 global memory、constant memory、texture memory 等。

线程模型

GPU 线程模型

如上图所示,这里展示了一个「Device -> Grid -> Block -> Thread」的示例,其中每一个层次都只使用了 2 维结构(实际可以是 1、2、3 维)。CUDA 使用 dim3 类型来表示定位变量,它可以看成是包含 3 个无符号整数的结构体变量,分别表示 3 个维度,默认为 1。对于图中的结构,其定义代码如下所示,在调用 kernel 时也必须通过执行配置 <<<grid, block>>> 来指定 kernel 所使用的线程数及结构:

1
2
3
dim3 grid(3, 2);
dim3 block(5, 3);
kernel_fun<<<grid, block>>>(prams...);

如果我们要定位一个具体的线程,那么我们需要用到两个内置的坐标变量 blockIdxthreadIdx 来唯一标识,这两个变量都是 dim3 类型,分别指明了线程在 grid 中的位置和在 block 中的位置。举例来说,图中的 Thread $(1,1)$ 满足:

1
2
3
4
threadIdx.x = 1
threadIdx.y = 1
blockIdx.x = 1
blockIdx.y = 1

除此之外,我们还需要知道 grid、block 的组织结构(即多少维、每维多大),从而获取该线程的全局 ID。block 的组织结构同样是通过线程的内置变量 blockDim 来获得的,对于一个 2 维的 block $(D_x, D_y)$,线程 $(x, y)$ 的 ID 为 $(x + y \cdot D_x)$,若是 3 维 block $(D_x, D_y, D_z)$,线程 $(x, y, z)$ 的 ID 则为 $(x + y \cdot D_x + z \cdot D_x \cdot D_y)$。与 blockDim 类似,线程有另一个内置变量 gridDim 来获取 grid 的组织结构。

Kernel 的这种线程组织结构天然适合 vector、matrix 等运算。还是以上图为例,假设我们将使用该组织结构实现两个矩阵的加法,每个线程负责处理每个位置的两个元素相加,线程块大小为 $(16, 16)$,然后将 $N \times N$ 的矩阵均分为不同的线程块来执行加法运算,代码如下所示:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
__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()
{
    // ...

    // Block: (16, 16) threads
    dim3 threadsPerBlock(16, 16);

    // Grid: (N / 16, N / 16) blocks
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y)

    // Call kernel
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);

    // ...
}

内存模型

GPU 内存模型

GPU 物理架构

GPU 的层次物理结构

  • 每个 GPU 包含多个 SM(streaming multiprocessor,流式多处理器),每个 GPU 的 SM 数量根据 GPU 的高中低端来决定,例如 Maxwell 架构的 GTX 1070 和 GTX 1080 分别有 15 个 SM(1920 个 SP)、20 个 SM(2560 个 SP)。
  • 每个 SM 包括多个 SP(streaming processor,流处理器),还包括 shared memory(共享内存)、register(寄存器)、warp scheduler(束调度器)等。SM 也被称为「大核」。每个 SM 中的 SP 数量依据 GPU 架构而不同,例如 Kepler 架构和 Maxwell 架构分别有 192、128 个。
  • SP 是 GPU 最基本的计算单元,也被称为 GPU 核(core),一个 SP 可以执行一个 thread,GPU 并行计算中具体的指令和任务都是在 SP 上进行处理的。

纹理贴图单元(texture mapping units,TMUs

光栅化处理单元(raster operations units,ROPs

逻辑架构与硬件架构的对应

GPU 逻辑架构与硬件架构的对应

GPU 包含多个 SM,中每个 SM 都包含多个 SP,所以一般可达数千个 SP,而每个 SP 都能执行一个线程,因此 GPU 支持这么多个线程并行进行。当一个 kernel 启动时,block(逻辑层)会被自动分配到各个 SM(物理层),而 SM 可以调度多个 block,并采用 SIMT(single instruction multiple thread,单指令多线程)架构,其基本的执行单元为 warp(即 GPU 调度和运行的基本单元),一个 warp 包括多个线程(例如 32),这些线程必须执行相同的指令,不过每个线程都有它自己的程序计数器、状态寄存器、执行路径,虽然 warp 中的所有线程同时从同一程序地址执行,但是比如遇到分支结构,不同的线程可能进入不同的分支结构,而在某些线程执行某个分支的指令时,另外一些不需要执行这些指令的线程则会死等,从而降低了性能。由于资源的限制,SM 通过 warp scheduler 来调度多个 warps,为每个 block 分配 shared memory,也要为每个 warp 中的线程分配独立的寄存器。因为一般 GPU 的 warp 大小为 32(Nvidia 可能会改),所以一般 block 的大小会设置为 32 的倍数

参考