(0) 概述和基本入门

(0.0) GPU 硬件简介

下图展示了 CPU 和 Nvidia GPU 的一般结构：

左图表示 CPU 结构，右图表示 GPU 结构。

图中紫色部分表示 L1 缓存，黄色部分表示控制器，绿色部分表示算术逻辑单元，蓝色部分表示 L2 及 L2 以上级别的缓存，橙色部分表示数据存储器 DRAM。

一般而言，从 CPU 的 DRAM 还有一条 PCIe 总线和 GPU 的 DRAM 相连。

下图可以看出 CPU 和 GPU 在结构上的巨大差异，前者将大量的晶体管单元用于控制器和缓存器的构造，而后者将大部分晶体管单元用于算术逻辑单元。这种结构差异导致，我们一般说的 GPU 计算并不是纯粹由 GPU 完成的，而是由 CPU 调度、GPU 计算共同完成的异构计算任务。

CPU 与 GPU 硬件结构

在异构计算任务中，CPU 被称为主机（Host），GPU 被称为设备（device）。

通过 nvidia-smi 命令可以轻松地获取当前设备下 GPU 的基本信息和状态参数，关于该命令的更多具体用法可以参考官方文档。

(0.1) CUDA 程序结构

CUDA 程序一般使用 .cu 后缀，编译 CUDA 程序则使用 nvcc 编译器。一般而言，一个 CUDA 程序的结构如下：

int main() {
    主机代码;
    核函数调用;
    主机代码;
    核函数调用;
    ......
    return 0;  
}

__global__ void 核函数1(parameters) {
    ......
}

__global__ void 核函数2(parameters) {
    ......
}

......

可以看到，在 main 函数中，我们穿插地写主机代码和核函数调用代码。主机代码主要负责 CPU 和 GPU 的内存管理、计算任务的分派，而核函数调用主要负责完成主要的计算工作。

前缀 __global__ 用来定义一个核函数，在 CUDA 中，核函数只能返回 void 类型，这意味着当我们需要写计算结果时，应该在参数列表中传入一个用来存放计算结果的指针，然后将计算结果写回到这个指针指向的存储空间中。

除了 __global__ 以外，CUDA 程序中的函数还可以使用 __device__ 和 __host__ 来修饰函数

__divice__ 修饰的函数称为设备函数，只能被核函数或是其它设备函数调用，只能在设备中执行。

__host__ 修饰的函数称为主机函数，它就是主机端的普通 C++ 函数，在主机（CPU）中调用和执行，可以忽略。

调用核函数进行运算涉及到 CUDA 的线程组织方式，见下一节。

(0.2) 主机和设备间数据传输

主机和设备间的数据传输是开启 CUDA 并行前的准备，这是因为CUDA 核函数传入的参数必须是指向设备内存的函数。因此，我们必须预先在主机代码中分配设备内存并初始化。分配设备内存可以使用 cudaMalloc 函数，初始化设备内存则可以将一段已经初始化好的主机内存拷贝到刚分配好的设备内存中，这可以使用 cudaMemcpy 函数实现，这两个函数的函数原型如下：

void cudaMalloc(void** d_ptr, unsigned int size);
void cudaMemcpy(void* d_ptr, void* h_ptr, unsigned int size, enum cudaMemcpyKind)

cudaMalloc 函数中需要注意的是，第一个参数是一个指针的指针，即二级指针，因为这个函数需要在 GPU 显存上分配出一片空间，并且让 d_ptr 指向这个空间，由于 cudaMalloc 函数没有返回值，因此需要传入一个二级指针，然后由 cudaMalloc 函数负责改变这个指针，使它指向刚分配出的那片设备内存。 一般而言，我们将指向主机内存的指针加上 h_ 前缀，而指向设备内存的指针加上 d_ 前缀。

cudaMemcpy 函数中，第一个参数是指向设备内存的指针，第二个参数是指向主机内存的指针，最后一个参数是 enum 类型的变量，用于指出数据传输的方向，它有五种取值，根据变量名就很容易看出数据传输的方向，比较常用的是前面两种。

cudaMemcpyHostToDevice

cudaMemcpyDeviceToHost

cudaMemcpyHostToHost

cudaMemcpyDeviceToDevice

cudaMemcpyDefault

下面是一个分配设备内存的例子，该程序希望完成的功能是将 h_x 和 h_y 数组加和的结果存放到 h_z 数组中，我们使用 CUDA 来解决这个问题时，需要先将 h_x 数组和 h_y 数组拷贝到位于设备内存的 d_x 和 d_y 数组中，下面这个程序片段即用来完成设备内存分配和初始化的功能。

double *h_x = (double*) malloc(M);
double *h_y = (double*) malloc(M);
double *h_z = (double*) malloc(M);

for (int n = 0; n < N; ++n) {
    h_x[n] = a;
    h_y[n] = b;
}

double *d_x, *d_y, *d_z;
cudaMalloc((void **)&d_x, sizeof(double)*N);
cudaMalloc((void **)&d_y, sizeof(double)*N);
cudaMalloc((void **)&d_z, sizeof(double)*N);
cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);

(0.3) CUDA 线程组织

核函数的调用语法如下所示：

kernel_function<<<grid_size, block_size>>>(parameters)

grid_size 和 block_size：grid_size 和 block_size都可以是一个 dim3类型的结构体，或是一个 unsigned 类型的无符号整数变量。前者表示网格大小，后者表示线程块大小。在 CUDA 的线程组织模型中，线程是最基本的单位，线程块由线程组成，而网格由线程块组成。它们的存在都是为了将线程组织成逻辑上更容易理解的线程单元。

SIMD：核函数的调用逻辑遵循 Single Instruction Multiple Data*（SIMD），即核函数的代码（Single Instruction）是对多份数据（Multiple Data）执行相同操作的一段指令。在调用核函数后，每个线程都会执行核函数中的代码。在下一节（CUDA 核函数设计）将会指出，如何通过设计特定的分支结构，来*使不同线程对不同数据执行核函数中指定的操作（这需要指定一种从线程到对应位置数据的索引方式）。

当 grid_size 和 block_size 均为无符号整数时。对应的是最简单的情形：

gridDim 和 blockDim 为一维变量

而当 grid_size 和 block_size 均为 dim3 结构体类型的变量时，情况会更加复杂，一个 dim3 结构体可以定义为二维或者三维的变量，如下：

dim3 grid_size(2, 3);    // 定义为二维变量
dim3 grid_size(2, 2, 2); // 定义为三维变量

当定义为二维变量时，可以想象，每个网格是由二维排列的块构成的，每个块内部由二维排列的线程构成。如下图：

gridDim 和 blockDim 为二维变量

当定义为三维变量时，可以想象，每个网格是由三维排列的块构成的，每个块内部由三维排列的线程构成。如下图：

gridDim 和 blockDim 为三维变量

显然，参与该核函数计算的线程个数，可以用如下公式表示：

num_threads = grid_size.x * grid_size.y * grid_size.z * block_size.x * block_size.y * block_size.z

维度为 1 或 2 的情形，只要将 y 、 z 分量视为 1 即可。

(0.4) CUDA 核函数设计

(0.4.0) 核函数设计范式

CUDA 的核函数设计一般遵循如下范式：

__global__ void kernel_function(data1, data2, ..., result) {
   index1, index2, ... = get_index(thread_info)   
   result = some_operations(data1[index1], data2[index2], ...)
}

data1，data2 ... 表示需要处理的数据指针，index1 和 index2 ... 用来定位需要计算的数据的位置，some_operation 对这些数据进行指定的计算操作，然后写回到参数列表中传入的用于记录结果的 result 指针中。

总结下来就是两部曲：

确定线程和数据的对应

对需要处理的数据执行操作

(0.4.1) 使用 CUDA 内建变量实现线程和数据的对应

上述代码中的 thread_info 就表示执行这个核函数的线程对应的信息，这些信息存储在 CUDA 的内建变量中：包括 gridIdx 和 blockIdx，以及一些全局的信息比如 gridDim 和 blockDim。

全局信息主要指明线程的个数和组织方式：

gridDim 是一个 dim3 类型的结构体，包含 x, y, z 三个变量，这个变量对应调用核函数时的 grid_size。

blockDim 是一个 dim3 类型的结构体，包含 x, y, z 三个变量，这个变量对应调用核函数时的 block_size。

局部信息主要指明该线程在所有参与执行该核函数的线程中的位置：

blockIdx 指明一个线程所在的线程块在网格中的位置。

blockIdx.x 的范围为 0 到 gridDim.x-1

blockIdx.y 的范围为 0 到 gridDim.y-1

blockIdx.z 的范围为 0 到 gridDim.z-1

threadIdx 指明一个线程在它所在的线程块中的位置。

threadIdx.x 的范围为 0 到 blockDim.x-1

threadIdx.y 的范围为 0 到 blockDim.y-1

threadIdx.z 的范围为 0 到 blockDim.z-1

不论是在 blockIdx 还是 threadIdx 中，x 都是变化最快的分量，其次是 y。

在实际核函数的设计中，我们就是通过 blockIdx 和 threadIdx 两个结构体来获取需要计算的数据的位置，即确定线程和数据的对应。

例如在上文中数组相加的例子中，我们可以将核函数这样设计：

void __global__ add(const double *x, const double *y, double *z, const int N) {
    const int n = blockDim.x * blockIdx.x + threadIdx.x;
    if (n < N) {  
        z[n] = x[n] + y[n];
    }
}

这个核函数中，负责实现线程和数据对应的代码是 const int n = blockDim.x * blockIdx.x + threadIdx.x;，负责对相应位置数据进行计算的代码是 z[n] = x[n] + y[n]; （即完成一个简单的加和操作）

注意这个核函数中的 if 语句，这个 if 语句的作用在于防止出现数组越界，这是一种非常惯用的处理方法：我们通常在确定 gridDim 和 blockDim 的大小时，会让实际执行的线程总数大于需要处理的数据数量（比如在数组加和的例子，数据数量就是数组长度 N），这样，一些线程实际上并不做任何计算，因为所有数据都已经被分配给其它线程处理。

众所周知，高维坐标实际上都可以一维坐标对应，我们用 bid 表示线程所在线程块在网格中的序号，用 tid 表示线程在它所在线程块中的序号，bid 和 tid 可以采用如下方式计算：

unsigned int bid = blockIdx.z * gridDim.x * gridDim.y + blockIdx.y * gridDim.x + blockIdx.x;
unsigned int tid = threadIdx.z * blockDim.x * blockDim.y + threadidx.y * blockDim.x + threadIdx.x;

在不同情况下，线程到数据的确定有不同的方法，这在不同

(0.4.2) 本节完整代码

上述数组求和的完整代码罗列如下，可以在这个链接找到：CUDA-Programming/src/03-basic-framework/add1.cu at master · brucefan1983/CUDA-Programming (github.com)，本文的这段代码参考自书本 《CUDA 编程基础与实践》（清华大学出版社）

#include <math.h>
#include <stdio.h>

const double EPSILON = 1.0e-15;
const double a = 1.23;
const double b = 2.34;
const double c = 3.57;
void __global__ add(const double *x, const double *y, double *z);
void check(const double *z, const int N);

int main(void) {
    const int N = 100000000;
    const int M = sizeof(double) * N;
    double *h_x = (double*) malloc(M);
    double *h_y = (double*) malloc(M);
    double *h_z = (double*) malloc(M);

    for (int n = 0; n < N; ++n) {
        h_x[n] = a;
        h_y[n] = b;
    }

    double *d_x, *d_y, *d_z;
    cudaMalloc((void **)&d_x, M);
    cudaMalloc((void **)&d_y, M);
    cudaMalloc((void **)&d_z, M);
    cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
    cudaMemcpy(d_y, h_y, M, cudaMemcpyHostToDevice);

    const int block_size = 128;
    const int grid_size = N / block_size;
    add<<<grid_size, block_size>>>(d_x, d_y, d_z);

    cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
    check(h_z, N);

    free(h_x);
    free(h_y);
    free(h_z);
    cudaFree(d_x);
    cudaFree(d_y);
    cudaFree(d_z);
    return 0;
}

void __global__ add(const double *x, const double *y, double *z, const int N) {
    const int n = blockDim.x * blockIdx.x + threadIdx.x;
    if (n < N) {
        z[n] = x[n] + y[n];
    }
}

void check(const double *z, const int N) {
    bool has_error = false;
    for (int n = 0; n < N; ++n) {
        if (fabs(z[n] - c) > EPSILON) {
            has_error = true;
        }
    }
    printf("%s\n", has_error ? "Has errors" : "No errors");

(0.5) 编译 CUDA 程序

编译 CUDA 程序需要使用 Nvidia 官方提供的编译器 nvcc。nvcc 会先将所有源代码先分离成主机代码和设备代码，主机代码完整支持 C++ 语法，设备代码只部分支持 C++ 语法。nvcc 先将设备代码编译为 PTX（parallel thread execution）伪汇编代码，再将 PTX 代码编译为二进制的 cubin 目标代码。

CUDA 中核函数也因此不能直接作为类的成员函数，如果希望使用面向对象，我们一般通过包装函数调用核函数，然后将这个包装函数作为成员函数。

在源设备代码编译为 PTX 代码时，需要使用 -arch=compute_XY 来指定一个虚拟架构的计算能力，在将 PTX 代码编译为 cubin 代码时，需要使用选项 -code=sm_ZW 指定一个真实架构的计算能力，真实架构号必须大于等于虚拟架构号。

X, Y, Z, W 实际使用时需要指定具体数字。

读者可以使用 nvidia-smi 查看本机选卡型号，然后在官网根据型号查询架构。

关于虚拟架构计算能力和真实架构计算能力的介绍，可以详见附录。

(1) CUDA 程序加速关键

本节是承前启后的一节，主要介绍影响 CUDA 程序加速性能的因素，包括数据传输比例、核函数内算术强度和访存时间，以及 SM 中驻留的线程数和并行规模。后两者是值得深入讨论的内容，事实上本笔记剩余的内容都可以看作是对后两者的优化，尤其是算术强度和访存时间的优化，它实在包括了太多的内容和技巧。

(1.0) 数据传输比例

在上一章节的介绍中，我们已经提到，进行 CUDA 并行计算前的必要准备是设备内存的分配和初始化，这一过程涉及到将主机内存拷贝到设备内存的操作。事实上，GPU 计算核心和设备内存之间数据传输的带宽要远高于 GPU 和 CPU 之间数据传输的带宽，也就是说，使用 cudaMemcpy 将数据从 CPU 转移到 GPU 的操作是非常费时的。

因此，要获得可观的 GPU 加速，就必须尽量缩减数据传输所花时间的比例。如果实际核函数中计算任务很重，那么 CPU 和 GPU 之间数据传输的时间就可以忽略不计，此时加速效果就会非常明显。

(1.1) 核函数内算术强度和访存时间

一个计算问题的算术强度是指其中算术操作的工作量（时间）与必要的内存操作的工作量（时间）之比。显然，如果一个计算问题的算术强度越高，那么它的加速效果越好。

提高算术强度和降低访存时间是一致的，在计算任务固定的情况下，这可以通过两个手段达到：

尽可能减少访存操作的次数

通过共享内存、寄存器内存等 CUDA 内存结构的合理使用来提高访存操作的次数。这也是本节讨论的主要内容

(1.2) SM 驻留线程数和并行规模

SM 全称为 Stream Multiprosessor，即流多处理器，一个 GPU 中包含众多 SM。要使一个 CUDA 程序达到最优性能，就需要提高每个流处理器能够驻留的线程数量，这和并行规模是息息相关的：并行规模可以定义为一个 GPU 中同时驻留的线程数目，我们可以通过调用核函数时指定 gridDim 和 blockDim 来确定并行规模。可以这样理解：

并行规模是从整个 GPU 处理器的角度看待驻留线程的数量。

SM 中驻留线程的数量则是从更微观角度看待。

影响 SM 中驻留线程数主要有两个因素：

并行规模。当并行规模不够大时，分配每个 SM 的线程数量太少，远远不能达到它能够处理的理论最大线程数，这时我们称 SM 的占有率很低，不能激发 GPU 的最大性能。

资源限制。每个 SM 都具有一定数量的寄存器、共享内存等资源，当我们在核函数中使用这些资源超过一定量时（单个线程使用量超标），SM 就不得不减少其内部驻留的线程数。

SM 总资源 = 单个线程占有资源数 * SM中线程数量

对于并行规模，这是根据具体问题确定的，这也是为什么更大规模的问题更具有利用 GPU 进行并行计算的价值。而对于资源限制，我们将在本章节中详细介绍完 CUDA 内存结构后单开一个话题讨论。

(2) CUDA 内存组织

(2.0) CUDA 内存结构

GPU 中的内存结构也具有高容量高延迟的特性，和 CPU 一样，GPU 中的内存也分为多个层次，不同层次具有不同的访问速度和容量大小。下图即表示了 CUDA 中不同内存的可见范围和访问速度（绿色最快，黄色次之，红色最慢）。

CUDA 内存结构

可以用如下表格表示：

我们可以从这个表格中观察到一些有趣的现象：并非可见范围越小的内存访问速度越快，比如局部内存和寄存器内存都是单个线程可见，但是它们之间的访问速度却差异巨大。在下文具体介绍这些内存时，将会详细指出。

(2.1) 动态全局内存

全局内存是核函数中所有线程都能够访问的内存，即全局内存对整个网格的所有线程都是可见的。而动态全局内存则特指采用 cudaMalloc函数分配的内存。

注意将它和 C++ 中的全局内存相区别，这里的“全局”强调的是一种线程可见范围。它的物理位置和访问速度是和 C++ 中的全局内存截然不同的。

全局内存由于位于 GPU 芯片外，因此具有较高的延迟和较低的访问速度，但其容量是所有设备内存中最大的，基本上就等于显存容量。

任何线程都能够访问全局内存的任何位置，考虑我们之前设计的核函数 add 中：

void __global__ add(const double *x, const double *y, double *z, const int N)) {
    const int n = blockDim.x * blockIdx.x + threadIdx.x;
    if (n < N) {  
    z[n] = x[n] + y[n];
    }
}

我们设计第 n 个线程能够访问位于数组 d_x, d_y, d_z （它们被分配在全局内存上）的第 n 个位置，但事实上，并非每个线程都一定要访问固定的位置，而是可以访问任意位置的数据。

(2.2) 静态全局内存

动态全局内存通过 cudaMalloc 函数分配，其所占空间是在运行时确定，CUDA 中还提供了一种静态全局内存，它所占空间在编译期间就能确定。

静态全局变量必须在所有主机和设备函数的外部定义，并且用 __device__ 关键字修饰，比如：

__device__ float x;
__device__ float y[N];

核函数可以对静态全局变量直接进行访问，访问方法和动态全局变量一致，但主机函数中不能对静态全局变量直接访问，而是需要要用 cudaMemcpyFromSymbol 和 cudaMemcpyToSymbol 来进行读写。这两个函数的原型如下：

cudaError_t cudaMemcpyFromSymbol(
    const void* symbol, // 静态全局内存变量名
    const void* src,    // 主机内存缓冲区指针
    size_t count,   // 复制的字节数
    size_t offset = 0,  // 从 symbol 对应设备地址开始偏移的字节数
    cudaMemcpyKind kind = cudaMemcpyHostToDevice // 可选参数
);

cudaError_t cudaMemcpyToSymbol(
    const void* dst,    // 主机内存缓冲区指针
    const void* symbol, // 静态全局内存变量名
    size_t count,   // 复制的字节数
    size_t offset = 0,  // 从 symbol 对应设备地址开始偏移的字节数
    cudaMemcpyKind kind = cudaMemcpyHostToDevice // 可选参数
);

这两个函数的最后两个参数都是可选参数，一般不需要指定。关于这两个函数的用法，在实践章节中会有更加详细的示例。

(2.3) 常量内存

常量内存是有常量缓存的全局内存，虽然理论上常量内存也是全局内存的一种，物理位置上都位于芯片外，但是由于常量缓存的存在，常量内存的访问速度是高于全局内存的。

我们可以在核函数外使用 __constant__ 标识符来定义分配到常量内存上的变量，然后用上文介绍的 cudaMemcpyToSymbol 函数赋值；

此外，向核函数传入的 const 参数也会被分配在常量内存上。 类似上文中核函数 add 的定义，我们传入的 const int N 参数实际上也会被分配在常量内存上。

对于上述的第二点，有一种技巧是将常量数组封装在结构体内，然后将结构体作为 const 参数传给核函数，此时同样会使用常量内存，在实践章节会对类似的技巧作更详细的介绍。

(2.4) 纹理内存和表面内存

纹理内存和表面内存类似于常量内存，它也是一种具有缓存的全局内存，但纹理内存和表面内存容量更大。

对于计算能力不小于 3.5 的 GPU 设备，可以使用 __ldg() 函数将一些只读全局内存加载到只读数据缓存中（read-only data cache），该函数在实践章节作更具体介绍。

帕斯卡架构及其更高架构默认使用 __ldg() 函数读取全局内存，故不需要显示使用。

(2.5) 寄存器内存

寄存器内存是单个线程可见的、CUDA 中访问速度最快的内存。一些高频访问的变量都应该放到寄存器内存中。比如，内建变量 gridDim, blockDim, blockIdx, threadIdx 实际上都存放在寄存器内存中，而上文中数组加和的例子，const int n = blockDim.x * blockIdx.x + threadIdx.x 定义的变量 n 实际上就位于寄存器内存中。

一般而言，单个线程可以使用的寄存器数上限为 255（每个寄存器可以存放 4 字节数据），一个线程块可以使用的寄存器数上限为 64k。因此我们在核函数中定义数据时需要时刻注意，一旦定义的数据量过大， 溢出的部分会被编译器放到局部内存中，（这会在下文介绍），局部内存的访问延迟是远不如寄存器内存的。

(2.6) 局部内存

局部内存在上文中实际上已经介绍过了，核函数中不加任何限定符的变量有可能位于寄存器内存中，也可能位于局部内存中。寄存器中放不下的变量都会被编译器分配到局部内存中。然而，尽管局部变量的可见范围也是单个线程，但是从硬件上看，局部内存也是全局内存的一部分，因此它的访问延迟和全局内存相当，在核函数中过多地使用局部内存会影响程序的性能。

(2.7) 共享内存

共享内存具有仅次于寄存器内存的访问速度，作为速度的代价，它的数量和可见范围也十分有限。共享内存仅对整个线程块可见，一个线程块中的所有线程都可以访问该线程块的共享内存，但不能访问其它线程块的共享内存。

共享内存分为静态共享内存和动态共享内存两种。作为共享内存的变量，一般都在前面添加 s_ 前缀修饰。

一般而言，静态共享内存和动态共享内存两种方式在执行时间上并无明显差别，后者有时还能够提高程序的可维护性。

使用静态共享内存，用 __shared__ 标识符在核函数中定义即可：

__shared__ float s_data [128];

使用动态共享内存，则相对复杂些，需要分为两步：

在调用核函数时，在 <<< >>> 中写入第三个参数，表示每个线程块中需要使用的动态共享内存字节数，如下：

kernel_function<<<grid_size, block_size, dynamic_shared_size>>>(parameters);

在核函数中，使用 extern 标识符以数组的形式声明动态共享内存，并且不能指定数组的大小：

__global void kernel_function(parameters) { 
   extern __shared__ float ds_data [];         
   ......     
}

利用共享内存进行访存优化是非常重要的 GPU 优化策略，我们在单独的一节讲解。

(2.8) 全局内存和局部内存的 L1 和 L2 Cache

为了缓解全局内存和局部内存（它们在物理上同属于一片内存区域）的访问延迟，CUDA 还在这部分内存上增设了 SM 层次的 L1 Cache 和设备层次的 L2 Cache（如下图所示）。L1 Cache 和共享内存同属于一块物理芯片，因此也具有较高的访问效率，L2 Cache 则次之。

全局内存和局部内存的 L1 和 L2 Cache

L1 Cache 和 L2 Cache 和共享内存不同，后者属于可编程的内存，即 CUDA 工程师可以指定将哪部分数据放到共享内存中进行加速，而 L1 Cache 和 L2 Cache 则没有共享内存那么灵活，将什么数据放到这部分内存上是 GPU 内部的算法决定的，一般而言，近期频繁使用的数据会被放到这上面，因此和 CPU 上编程一样，我们需要尽量提高数据的局部性来充分利用 L1 和 L2 Cache 改善程序的性能。

并且，这里的数据局部性是相对的，比如 L1 Cache 是 SM 层次的，那么我们会希望同一个线程块中的线程读取的全局内存位置具有局部性，这样这个线程块访问数据时将具有加速效果。

下文介绍 SM 时将会提到，一个线程块固定分配给一个 SM 进行执行。

一些之前没有接触过 CPU 程序优化的读者可能不熟悉数据的局部性，可以详见章节《基于全局内存合并访问的加速策略》。

(3) 基于缓解 SM 资源限制的加速策略

(3.0) 再看 GPU 结构

一些读者可能在读到目前部分时，会对 GPU 的结构产生一些误会，认为 GPU 是由一系列线程块组成的，因为我们在调用核函数指派线程数目时，就是将它们组织成一个个线程块。但是 GPU 计算部分的组成单元实际上是 SM（Stream Multiprosessor）。GPU、SM 和线程块的关系实际上如下图所示：

GPU、SM 和线程块的关系

一个 GPU 包括多个 SM。

在执行核函数时，一个 SM 可能会被分配执行若干个线程块。但每个线程块只会被一个 SM 执行。即调度映射

是一个函数。 一个 SM 最多被分配到的线程块数目为 16 或 32（取决于架构，详见官网），最多被分配到的线程数目为 1024 或 2048（取决于架构，详见官网）

SM 在执行时，一般以线程束（warp）为单位产生、管理、调度和执行线程，一个线程束包括 32 个线程。 针对线程束特性的优化将在后文介绍

SM 中应该包含如下结构（如图所示）：

一定数量的寄存器内存

一定数量的共享内存

一定数量的常量缓存

一定数量的纹理和表面缓存

一定数量的全局和局部内存的 L1 Cache

若干线程束调度器（warp scheduler），用于在不同线程的上下文之间迅速地切换

若干执行核心（core）

(3.1) 优化 SM 的占有率

在前面已经提到，一个 SM 中最多驻留的线程数为 2048（图灵架构为 1024，以下讨论均基于 2048），最多驻留的线程块数量为 16（部分架构为 32，以下讨论均基于 16），当一个 SM 能够运行 2048 个线程时，我们称此时 SM 的占有率为 100%。这不仅需要足够大的并行规模，还需要精心限制每个线程中使用的资源，这是因为每个 SM 中的寄存器内存容量和共享内存容量都是有限的。

对于寄存器内存：比如，若希望在 SM 中满负荷地驻留 2048 个线程，考虑到 SM 中最多允许使用的寄存器容量为 64 KB，每个线程最多使用 32 个寄存器，假设一个线程使用了 64 个寄存器，这虽然没超过额定的最大值（一般为 255，详见寄存器内存一节），不至于溢出到局部内存中，但是此时 SM 中能够驻留的线程数则变为 1024，此时占有率为

。

对于共享内存：由于共享内存是线程块共用的，假设此时一个核函数指定线程块大小为 128，那么一个 SM 至少需要同时处理 16 个线程块才能达到满负荷 2048 线程数运行。假设我们的芯片为 Tesla V100，共享内存大小为 96 KB，那么每个线程块使用 6 KB 以内的共享内存才能达到

占有率。

如果线程块使用的共享内存大于 SM 上允许的最大容量，将直接导致核函数无法运行

总之，上述讨论只是给出了分析计算的一般方法，实际架构需要结合具体参数进行计算。SM 的占有率也不一定需要维持

，一个比较可观的数值即可。总的原则就是尽可能不要浪费这些资源，以下援引 Nvidia 官方文档：

the fewer registers a kernel uses, the more threads and thread blocks are likely to reside on a multiprocessor, which can improve performance.

(3.2) SM 运行参数查询和控制

查询：使用编译器选项 --ptxas-options=-v 可以查询每个核函数的寄存器容量使用情况

使用 __launch_bounds__ 修饰符：该修饰符可以指定每个线程块的最大线程数，以及每个 SM 至少应该被分配多少个块。

#define MAX_THREADS_PER_BLOCK 256   
#define MIN_BLOCKS_PER_MP     2
__global void launch_bounds(MAX_THREADS_PER_BLOCK, MIN_BLOCKS_PER_SM)   
fooKernel(int inArr, int outArr) {
    // ... Computation of kernel   
}  

使用 --maxrregcount 编译选项：--maxrregcount 可以指定每个线程能够使用的寄存器数量，多出的部分，会溢出到局部内存中，因此这个选项是可能造成负面影响的：SM 的占有率可能由于这个限制提高了，但延迟较高的局部内存也是我们不愿看到的。

(4) 基于全局内存合并访问的加速策略

虽然标题为“基于全局内存合并访问的加速策略”，但是这些结论对于局部内存也奏效，因为我们知道局部内存和全局内存物理上属于芯片的同一块区域。

在本节，读者需要预先了解以下三个基本事实：

一个 SM 中执行的线程束（warp），每次请求的字节数一般是 32 字节。

使用 cudaMalloc 分配的内存首地址是 256 字节的整数倍。

加入 Cache 后的全局内存访问模式（见下文）。

(4.0) 加入 Cache 后的全局内存访问模式

上文提到过，由于全局内存的高访问延迟，我们加入 SM 层次的 L1 Cache 和设备层次的 L2 Cache 来缓解这一情况。加入这两级 Cache 后的访问模式是这样的：

线程首先访问它所在的 SM 上的 L1 Cache，如果发现需要的数据，则直接取回。

如果需要的数据不在 L1 Cache，即发生 L1 Cache Miss，就访问设备层次的 L2 Cache，如果发现需要的数据，则直接取回。

如果需要的数据不在 L2 Cache，即发生 L1 Cache Miss 和 L2 Cache Miss，就访问 DRAM 获取数据。这种情况的访问延迟最大。

加入 Cache 后的全局内存访问模式

(4.1) 数据局部性带来的访存优势

在本节，我们以数组倍增为例，具体核函数代码如下：

__global__ void double_array(float* data_in, float* data_out) {
     unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
     data_out[tid] = data_in[tid] * 2;
}

假设 data_in 的首地址被分配在 0 处，blockDim.x 为 64，threadDim.x 为 32。假设一开始 L1 Cache 为空，并且假设由 0 号线程先执行（我们用 tid 对线程进行编号），那么访问 data_in 的的第一个元素时就会发生 Cache Miss（此例子中我们只考虑 L1 Cache）。但是相应的，与 0 号相邻的 7 个线程在访问它们对应位置数据时，由于数据已经被加载到 L1 Cache 中，因此它们不会发生 Cache Miss。这样，在一个线程束的执行中，只会发生 4 次 Cache Miss。显然这四次 Cache Miss 已经是“必要”发生的最少情形，我们称这样的访问为“合并访问（coalesced）”。

非合并访问的例子随处可见，考虑矩阵乘法的例子， A*B=C,我们按行读取A矩阵的数据，按列读取B矩阵的数据，假设

B矩阵的每行元素足够多（以至于相邻行的元素地址间隔都在 32 字节以上），那么读取B矩阵列的访问必然属于非合并访问，并且合并度极低（访问每个元素都触发 Cache Miss）。

请不要沮丧，在实践章节中，我们将展示一些技巧来解决这个问题。

(5) 基于共享内存的加速策略

(5.0) 使用共享内存缓解非合并访问

在共享内存一节中已经介绍了如何在核函数中定义共享内存，包括静态共享内存和动态共享内存。而本节将重点介绍共享内存的使用场景。

考虑上文提到的矩阵乘法的例子，我们似乎已经有一个比较沮丧的结论：对矩阵

的合并访问必然伴随着对矩阵

的非合并访问，并且访问矩阵

的合并度是极低的，但是共享内存的出现似乎带来了转机：

在假设共享内存无限大的情况下：我们可以把整个矩阵

装到共享内存中，这样，对矩阵

行的访问可以被 L1 Cache 加速，对矩阵

列的访问可以被共享内存加速。

事实上，共享内存并非无限大的，矩阵

无法整个装入共享内存中，因此我们往往采用矩阵分块的方式，让每个线程块负责一个小方块的计算任务，如下图示意。这些技巧都会在实践章节具体讲解。

矩阵分块计算乘积

以上所述就包含了利用共享内存的一般优化思路：对于一些不得不造成非合并访问的访存操作，我们可以预先将它们加载到共享内存中，然后再进行计算。这一操作思路可以概括为如下编程范式：

__global__ void kernel_function(parameters) {
     // 假设算法的特殊设计导致对变量 data 的访问不得不是非合并的
     // 1. 定义共享内存
     __shared__ float s_data [data_size];
     // 2. 将 data 复制到共享内存
     s_data[copy_index] = data[copy_index];
     // 3. 等待所有线程完成复制操作
     __syncthreads();
     // 4. 进行操作（包含非合并内存访问）
     operations(data);
}

即大致分为四个步骤，其中第三步调用了函数 __syncthreads()，这是因为进行正式的操作前应该等待所有线程将数据拷贝结束，确保结果的正确性。

以上内容都局限于纸上谈兵，但作为初学者教程已经足矣，实践章节将见证共享内存带来的真正性能提升。

(5.1) 规避共享内存 Bank 冲突

在介绍 Bank 冲突前，我们需要先介绍 GPU 共享内存的物理结构。

对硬件结构了解的读者会知道，为了实现并行读取数据，一些内存结构需要提供多端口读取能力。而真正的多端口读写结构需要消耗大量的布线资源。

题外话：

假设我们的共享内存结构如下图所示：一个 warp 的共享内存为 4 KB。

设想一下，我们需要支持 warp 中的 32 个线程都能访问共享内存，那么就需要 32 个访问端口，每个端口都需要能够索引 1024 个位置（四字节对齐，4096/4），需要 10 位二进制数，32 个端口一共需要 320 bit 大小的端口。

因此实际设计时，多端口结构往往被设计成分区式：字节数模 4 相等的地址被分配到同一个分区（bank）。每个分区（bank）都分配一个读写端口。

题外话：

在上述情况下，每个分区包含 32 个 4 字节数据。需要 5 位二进制数寻址，32 个端口一共需要 160 bit 大小的端口。端口总位宽减少了一半。

GPU 共享内存物理结构

在每个计算周期，一个线程可以被分配一个访问端口，不同线程不会被分配同一个访问端口。因此，当一个线程束中的多个线程需要访问同一个分区的不同数据时，会发生 bank 冲突——同一个访问端口无法在同一周期为多个线程服务。这样会导致计算性能下降。

注意：访问同一个分区的相同数据并不会导致 Bank Conflict

在实践中，我们可以通过查阅官方文档查看共享内存结构，改写代码避免出现 Bank Conflict。

(6) 原子函数

本笔记假设读者对引入并行计算可能伴生的一系列问题（如读写冲突、写写冲突）已经有一定了解，这些内容已经是老生常谈，此处并不作重复讨论。我们直接看原子函数的用法和作用：

原子函数的作用在于对某个位置的数据进行一次“不可分割”的“读-改-写”操作，下面展示这些原子函数的用法。

原子加法：T atomicAdd(T *addr, T val);C++ *addr = *addr + val; return *addr;

原子减法：T atomicSub(T *addr, T val);C++ *addr = *addr - val; return *addr;

原子交换：T atomicExch(T *addr, T val);C++ *addr = val; return *addr;

原子最小值：T atomicMin(T *addr, T val);C++ *addr = *addr < val ? *addr : val; return *addr;

原子最大值：T atomicMax(T *addr, T val);C++ *addr = *addr > val ? *addr : val; return *addr;

原子自增：T atomicInc(T *addr, T val);C++ *addr = *addr >= val ? 0 : *addr+1; return *addr;

原子自减：T atomicDec(T *addr, T val);C++ *addr = *addr==0 || *addr>val ? val : *addr-1; return *addr;

原子比较-交换：T atomicCAS(T *addr, T compare, T val);C++ *addr = *addr == compare ? val : *addr; return *addr;

原子按位与：T atomicAnd(T *addr, T val);C++ *addr = *addr & val; return *addr;

原子按位或：T atomicOr(T *addr, T val);C++ *addr = *addr | val; return *addr;

原子按位异或：T atomicXor(T *addr, T val);C++ *addr = *addr ^ val; return *addr;

需要注意的是：上面这些函数都是设备函数，即它们只能在核函数中使用，并且和其它并行工具一样，CUDA 的原子函数也会不可避免地伤害性能，我们应该只在必要时使用。

(7) CUDA 线程束

(7.0) SIMT 特性引发的线程发散

SIMT 是 Single Instruction Multiple Threads 的缩写，即单指令-多线程的执行模式，它指的是：在同一时刻，一个线程束中的线程只能执行一个共同的指令或者闲置。

当核函数出现条件分支时，假设条件分支指令有两个目的地址：if(condition) A else B，一个线程束实际上需要分两个计算周期来完成这件事：

首先，满足分支条件的线程执行分支 A，其它线程闲置。

然后，不满足分支条件的线程执行分支 B，其它线程闲置。

SIMT 的特性是针对一个线程束而言的，也就是说，我们并不必要整个系统的线程、或是一个线程块中的线程尽量走一个同分支，我们只需要让一个线程束中的所有线程尽量走同一个分支即可。

判断线程是否属于同一个线程束只要看 tid 即可，tid / 32 相等的线程会被分入同一个线程束，然后被同一个 SM 执行。

由于线程束的大小为 32，所以我们一般会将线程块的大小指定为 32 的倍数，让它恰好被分入若干个线程中。

(7.1) 线程束同步函数

在上文中，我们已经使用过线程同步函数 __syncthreads()，该函数的功能是阻塞所有已经到达该位置的线程，直到所有线程都已经到达这个位置才会放行：

阻塞所有已经到达该位置的线程，直到所有线程都已经到达这个位置才会放行

另一个常用的线程束同步函数是 __syncwarps()，它和 __syncthreads() 的区别是：

__syncthreads 等待范围更广，包括系统的所有线程

__syncwarps 等待范围更窄，包括同一个线程束中的所有线程，而不需要等待其它线程束中的线程。

syncwarps 函数还有一个 unsigned int 类型的 mask 参数，默认值为 0xFFFFFFFF，这个参数相当于一个 bitmap，对应二进制位为 1 的线程都参与同步，使用这个掩码可以实现更加精细的同步。

(7.2) 线程束表决函数

线程束表决函数实际上也完成了线程束同步函数的功能（具体解释详见附录），也就是说，它们在实现 syncwarps 函数“同步”功能的基础上，还能够额外地完成一系列条件判断功能。

__ballot_sync(unsigned mask, int condition)：该函数在同步功能的基础上，还能获取参与同步且满足指定条件的线程束索引。

mask：指定线程束内的哪些线程参与同步和条件判定。

condition：用于条件判定，每个线程都应该预先计算出一个 condition 值并作为参数传入。

return：返回一个 bitmap（unsigned int），如果一个线程参与了同步，并且它在核函数中调用 __balloc_sync 时传入的 condition 值非零，那么这个线程在 bitmap 中对应的位置就会被置为 1，否则为 0。这个 bitmap 相当于指出所有满足上述条件的线程在线程束中的位置。

__all_sync(unsigned mask, int condition)：该函数在同步功能的基础上，还能判断参与同步的线程是不是全部满足条件。

mask：指定线程束内的哪些线程参与同步和条件判定。

condition：用于条件判定，每个线程都应该预先计算出一个 condition 值并作为参数传入。

return：返回 0 或 1，如果参与同步的所有线程在核函数中调用 __all_sync 时传入的 condition 值均非零，就返回 1，否则返回 0。

__any_sync(unsigned mask, int condition)：该函数在同步功能的基础上，还能判断参与同步的线程是不是存在至少一个满足条件。

mask：指定线程束内的哪些线程参与同步和条件判定。

condition：用于条件判定，每个线程都应该预先计算出一个 condition 值并作为参数传入。

return：返回 0 或 1，如果参与同步的所有线程在核函数中调用 __any_sync 时传入的 condition 值有一个非零，就返回 1，否则返回 0。

(7.3) 线程束洗牌函数

线程束洗牌函数实际上也完成了线程束同步函数的功能（具体解释详见附录），也就是说，它们在实现 syncwarps 函数“同步”功能的基础上，还能够额外地完成在线程间共享数据的功能。

以下四个函数都是基于类似的逻辑，即线程束内部的数据广播机制。每个线程都可以在同步时将一个寄存器的数据 val 送入广播通路，届时其它线程将收到通路广播的内容。而 width 参数用于指明在同步和广播时，以多少个线程为一组，默认就是线程束的大小 32。然后会有一个参数 srcLane 指明哪个线程的数据需要广播给其它线程。下图即展示了 width = 8 （即组宽度为 8）函数 __shfl_sync 的工作逻辑。

线程束内部数据广播机制

__shfl_sync(unsigned mask, int val, int srcLane, int width)：该函数在同步功能的基础上，还能将指定线程中的指定寄存器值共享到其它线程中。

mask：指定线程束内的哪些线程参与同步和条件判定。

val：需要广播的寄存器值。

srcLane：提供广播寄存器值的源线程标号。

width：指明在同步和广播时，以多少个线程为一组。

__shfl_up_sync(unsigned mask, int val, int dist, int width)

mask：指定线程束内的哪些线程参与同步和条件判定。

val：需要广播的寄存器值。

dist

如果调用该函数的线程标号为 tid，将返回标号为 tid-dist 的线程传入的 val 值。

如果调用该函数的线程标号 tid < dist，直接返回该线程自身的 val（因为此时 tid - dist 小于零）。

width：指明在同步和广播时，以多少个线程为一组。

__shfl_down_sync(unsigned mask, int val, int dist, int width)：

mask：指定线程束内的哪些线程参与同步和条件判定。

val：需要广播的寄存器值。

dist

如果调用该函数的线程标号为 tid，将返回标号为 tid+dist 的线程传入的 val 值。

如果调用该函数的线程标号 tid+dist>=width，直接返回该线程自身的 val（因为此时 tid+dist 溢出）。

width：指明在同步和广播时，以多少个线程为一组。

__shfl_xor_sync(unsigned mask, int val, int laneMask, int width)：

mask：指定线程束内的哪些线程参与同步和条件判定。

val：需要广播的寄存器值。

laneMask：调用该函数的线程标号为 tid，将返回标号为 tid ^ laneMask 的线程传入的 val 值。

width：指明在同步和广播时，以多少个线程为一组。

线程束洗牌函数的数据共享更适用于在线程束内实现轻量级的数据共享，它的效率比直接使用共享内存更高，但是共享内存的共享范围是一个线程块，能够共享的数据规模更大，而线程束洗牌函数的共享范围最大也不超过一个线程束的大小（warpWidth = 32），并且传输规模仅为一个寄存器大小。

(7.4) 线程束洗牌函数实例

由于线程束洗牌函数相对而言较为抽象，因此此处补充一个实例来说明此系列函数的效果。

本部分施工中......

(8) CUDA 异步执行

(8.0) CUDA 流

所有和 CUDA 相关的操作都被称为 CUDA 操作，比如核函数的调用、主机设备端数据的传输。由 CUDA 操作组成的序列被称为 CUDA 流。

在没有明确指定一个 CUDA 操作要在哪个流中执行时，CUDA 操作都会被分配到空流（默认流），此前讨论的所有操作都在空流中进行，我们可以手动创建一个自己的流：

cudaStream_t my_stream;
// 创建流
cudaStreamCreate(&my_stream);
// 销毁流
cudaStreamDestroy(my_stream);

首先声明一个 cudaStream_t 类型的变量，然后将它的指针传入 cudaStreamCreate 函数即可创建一个新的流（该函数需要改变 my_stream 的值）。将 my_stream 的值传入 cudaStreamDestroy 即可完成流的销毁。

cudaStreamCreate 和 cudaStreamDestroy 两个函数均会返回一个 cudaError_t 类型的返回值，用于记录在创建/销毁过程中发生的错误，如何解析这个参数将在后续章节介绍。

CUDA 还提供了两个函数来让主机灵活地查询流的状态并决定是否挂起：

cudaError_t cudaStreamSynchronize(cudaStream_t stream); 
cudaError_t cudaStreamQuery(cudaStream_t stream);

第一个函数会阻塞主机，直到参数 stream 指定的流执行完毕；第二个函数不会造成阻塞，它仅仅返回参数 stream 指定的流的执行状态，如果执行完毕则返回 cudaSuccess，否则返回 cudaErrorNotReady。

cudaErrorNotReady 虽然是一个 CUDA 错误，但它却不是真正意义的错误，只是为了

(8.1) 主机代码和核函数的并行

CUDA 流可以理解 CUDA 操作的载体，处于同一个流中的 CUDA 操作需要遵循如下规则：

规则 1：核函数的启动是异步的，即 CPU 在执行到调用核函数的代码时，仅仅是向 GPU 发出一系列命令，使得 GPU 可以开始运算，在 GPU 完成计算前，CPU 并不会阻塞在这行代码，而是会继续执行后面的主机代码。

规则 2：基于 cudaMalloc 的主机和设备间数据的传输是阻塞的，主机发出 cudaMemcpy 命令之后，会等待数据传输的结束，在此过程中，主机会被挂起而并不往下执行。

规则 3：同一个 CUDA 流中的操作必须严格依序执行，如果上一个操作还没执行完，那么该操作将会被阻塞，直到 GPU 完成上一个操作才会启动这个操作。

一般情况下，我们在核函数的执行后紧接的操作就是将核函数运算的结果通过 cudaMemcpy 传回主机，由于 cudaMemcpy 会阻塞主机并且是空流中在核函数后的下一个操作，因此主机实际上不得不等到核函数执行完才执行下一步操作。

但由于核函数的启动是异步的，我们不难想到这样的加速思路：在核函数启动以后和得到核函数计算的结果之前，我们可以让主机完成另一部分计算任务，最后，核函数计算的时间如果刚好能够覆盖主机计算的时间，主机完成的计算部分就完全是我们的收益，该思路如下图：

在核函数启动以后和得到核函数计算的结果之前，可以让主机完成另一部分计算任务

尽管一些情况下，和 GPU 相比，CPU 的计算速度并不算快，但 CPU 较大的内存使得其适合计算一些需要大量内存的任务，主机和核函数的并行事实上也是一些顶会工作的思路。

(8.2) 多个核函数的并行

考虑 8.1 的规则 3，同一个 CUDA 流中的两个核函数必然是串行的，如果我们希望在同一时刻运行两个核函数，就需要将它们放到不同的 CUDA 流中运行，我们可以在尖括号 <<< >>> 内设置第四个参数来指定一个核函数运行在哪个流上：

kernel_function<<<grid_size, block_size, shared_size, stream_id>>>(parameters);

注意，由于 <<< >>> 中的参数都是根据位置确定的，因此，指定 stream_id 的核函数必须指定 shared_size，如果不需要使用共享内存，也必须在第三个参数的位置放一个 0。

利用 CUDA 流来并发多个核函数是有好处的，它可以提升 GPU 硬件的利用率，减少同一时刻限制的 SM，在整体上提升系统性能。

(8.3) 流水线并行：数据传输和核函数的并行

上文提到到，基于 cudaMemcpy 函数的数据传输会阻塞主机代码，这样，在数据传输还没有完成时，主机也无法发出调用下一个 CUDA 操作的指令。即任何其它 CUDA 操作都无法在调用 cudaMemcpy 函数后与之并发。

考虑下面的应用场景：我们可以将 Host2Device Memcpy，Call Kernel Function，和 Device2Host Memcpy 三个操作组合成流水线并行的工作流（如下图示意），每个时刻设备最多可以同时执行三个 CUDA 操作，提升了 GPU 硬件的并行度和利用率。

关于流水线并行带来性能提升，在很多教程中有过讨论，本处就不赘述了。相关信息可以参考附录。

CUDA 操作流水线并行

我们一般的做法是，将每条完整的流水线放到一个独立的 CUDA 流中执行，为完成这样的功能，我们需要主机在进行数据传输时不陷入阻塞，而是能够立刻抽身出来，完成同时刻的其它 CUDA 操作的调度。

比如在上图中红色虚线框内，同一时刻需要调度两个数据传输操作和一个核函数调用，显然，不能有任何一个数据传输操作使主机陷入阻塞。

因此我们需要不产生阻塞的数据传输函数，cudaMemcpyAsync 就能实现这样的功能：

cudaError_t cudaMemcpyAsync(
    void *dst,
    void *src,
    size_t num_bytes,
    enum cudaMemcpyKind kind,
    cudaStream_t stream
);

它的用法和 cudaMemcpy 几乎一致，但多了 stream 参数用于指定这个 CUDA 操作将在哪个流上运行。特别需要注意的是，在使用 cudaMemcpyAsync 函数时，必须使用不可分页主机内存，不可分页主机内存可以使用 cudaMallocHost 分配，不可分页主机内存必须要使用相应的 cudaFreeHost 释放：

cudaError_t cudaMallocHost(void **ptr, size_t size);
cudaError_t cudaFreeHost(void* ptr);

不熟悉操作系统知识的读者只需要记住这一规则即可，不可分页主机内存的概念可以参见附录。

cudaMemcpyAsync 函数是非阻塞的，这意味着我们可以在同一时刻，在不同的 CUDA 流中进行不同的数据传输任务，满足了我们实现流水线并行的需求。

(9) CUDA 错误检测

CUDA 已经了一套便于开发者进行错误检测的机制，对于错误检测，我们可以分以下两种情形处理：

返回 cudaError_t 的 CUDA API 函数

无返回值的 CUDA 核函数

(9.0) 返回 cudaError_t 的 CUDA API 函数的错误检测

对于返回 cudaError_t 的 CUDA API 函数的错误检测，我们可以采用以下宏函数实现：

#define CHECK(call)                     
do                                  \
{                                   \
    const cudaError_t error_code = call;                \
    if (error_code != cudaSuccess)                  \
    {                                   \
    printf("CUDA ERROR:\n");                    \
        printf("    FILE:   %s\n", __FILE__);           \
        printf("    LINE:   %d\n", __LINE__);           \
        printf("    ERROR CODE: %d\n", error_code);         \
        printf("    ERROR TEXT: %s\n", cudaGetErrorString(error_code)); \
    exit(1);                            \
    }                                   \
} while(0)

只需要在调用 CUDA API 函数套上该宏函数即可，比如：

CHECK(cudaMalloc(parameters...));

(9.1) 无返回值的 CUDA 核函数的错误检测

由于 CUDA 核函数不返回任何内容，因此不能使用上述方法进行错误检测，但是我们仍然可以在调用核函数后立即调用 cudaGetLastError 和 cudaDeviceSynchronize 两个函数并使用如上方法解析它们的 cudaError_t 返回值，来判断刚刚运行的核函数发生了什么错误：

kernel_function(parameters);
CHECK(cudaGetLastError());
CHECK(cudaDeviceSynchronize());

(10) CUDA 标准库

CUDA 提供了一系列的标准库，用于更加安全和高效地实现各种功能：

此处并不再对这些库的用法一一做介绍，因为这逾越了基础教程的功能，笔者也建议读者不必要在初学时过多地记忆这些库的用法，在实践和与 ChatGPT 的交流中积累技巧即可。

(11) 附录

Chat-0：关于 CUDA 虚拟架构和真实架构计算能力

关于 CUDA 虚拟架构和真实架构计算能力

Chat-1：关于线程束表决函数和线程束洗牌函数的同步功能

关于线程束表决函数和线程束洗牌函数的同步功能

Chat-2：关于流水线并行

关于流水线并行

Chat-3：关于不可分页内存

关于不可分页内存