1.认识CUDA编程
2006年，NVIDIA公司发布了CUDA(Compute Unified Device Architecture, 统一计算设备架构)，是建立在NVIDIA的CPUs上的一个通用并行计算平台和编程模型。 基于CUDA编程可以利用GPUs的并行计算引擎来更加高效地解决比较复杂的计算难题，广泛应用于深度学习领域，基于GPU的并行计算已经成为训练深度学习模型的标配。
GPU并不是一个独立运行的计算平台，而需要与CPU协同工作，可以看成是CPU的协处理器，因此当我们在说GPU并行计算时，其实是指的基于CPU+GPU的异构计算架构。在异构计算架构中，GPU与CPU通过PCIe总线连接在一起来协同工作，CPU所在的位置称为主机端(host)，而GPU所在的位置称为设备端(device)，如下图所示：
可以看到，GPU包含更多的运算核心，其特别适合数据并行的计算密集型任务，如大型矩阵运算，而CPU的运算核心较少，但是其可以实现复杂的逻辑运算，因此其适合控制密集型任务。另外，CPU的线程是重量级的，上下切换开销大，但是GPU由于存在很多核心，其线程是轻量级的。因此，基于CPU+GPU的异构计算平台可以优势互补，CPU负责处理逻辑复杂的串行程序，GPU重点处理数据密集型的并行计算程序，从而发挥最大功效。
CUDA是NVIDIA公司开发的GPU编程模型，它提供了GPU编程的简易接口，基于CUDA编程可以构建基于GPU计算的应用程序。CUDA提供了对其它编程语言的支持，如C/C++，Python，Fortran等语言。
2.CUDA编程基础
2.1 host和device
CUDA异构计算架构中，CPU和GPU是协同工作的。host和device是两个重要的概念，我们用host指代CPU及其内存，而用device指代GPU及其内存。host程序在CPU上运行，device程序在GPU上运行，host和device之间可以相互通讯(进行数据拷贝)。
CUDA编程中，通过3个函数类型限定词(__global__、__device__、__host__)来区分host函数和device函数。具体如下：
-
__global__：在device上执行，从host中调用(一些特定的GPU也可以从device上调用)，返回类型必须是void，不支持可变参数，不能成为类成员函数。注意：下文会提到CUDA中的核函数，它就是用__global__声明的，并且是异步的，host不会等待kernel执行完就执行下一步。 -
__device__：在device上执行，从device中调用。 -
__host__：在host上执行，从host中调用，一般省略不写(默认)。
典型的CUDA程序执行流程如下：
1.分配host内存，并进行数据初始化；
2.分配device内存，并从host将数据拷贝到device上；
3.调用CUDA的核函数在device上完成指定的运算；
4.将device上的运算结果拷贝到host上；
5.释放device和host上分配的内存。
2.2 kernel
在CUDA的执行流程中，最重要的一个过程是调用CUDA的核函数来执行并行计算。kernel是CUDA中一个十分重要的概念，kernel是在device上实现并行计算的函数，核函数用__global__符号声明，调用时用<<<grid, block>>>来指定kernel要执行的线程数量。在CUDA中，每个线程都要执行核函数，并且会为每个线程分配一个唯一的线程号thread ID，这个ID值可以通过核函数的内置变量threadIdx来获得。
让我们进一步理解kernel的线程层次结构。GPU上有很多并行化的轻量级线程，kernel在device上执行时，实际上是启动很多线程，一个kernel所启动的所有线程称为一个网格(grid)，同一个网格上的线程共享相同的全局内存空间，grid是线程的第一个层次结构，而网格又可以分为很多线程块(block)，一个线程块中包含多个线程，是线程的第二个层次结构。线程的两层层次结构，如下图所示(这是一个grid和block均为2-dim的线程结构示例)：
其中，grid和block都定义为dim3类型的变量，dim3是包含3个无符号整数 (x, y, z) 成员的结构体，在定义时，缺省值为1。grid和block可以灵活的定义为1-dim、2-dim以及3-dim结构，kernel在调用时通过执行配置<<<grid, block>>>来指定kernel所使用的线程数及结构。以上图为例，定义grid和block、核函数调用的代码如下：
dim3 grid(3, 2);
dim3 block(5, 3);
kernel_fun<<<grid, block>>>(params...);
对于每个线程，通过两个内置变量 (blockIdx, threadIdx) 来唯一标识，它们都是uint3类型。其中，blockIdx指明线程所在grid中的位置，而threadIdx指明线程所在block中的位置。以上图中的Thread(3, 1)为例，坐标满足：
blockIdx.x = 1
blockIdx.y = 1
threadIdx.x = 3
threadIdx.y = 1
值得注意的是，一个线程块上的线程是放在同一个SM(Streaming Multiprocessor, 流式多处理器)上的，但是单个SM的资源有限，这导致线程块中的线程数是有限制的，现代GPUs的线程块中的线程数上限为1024个。
通过内置变量gridDim、blockIdx、blockDim、threadIdx，可以计算每一个线程的唯一标识ID——threadId。其中，gridDim用于获取grid各个维度的大小，blockDim用于获取block各个维度的大小。对于任意一个线程，它的blockId、threadId的计算公式：
上面的计算公式，适用于任意维度(1维、2维、3维)，以grid为2维、block为2维为例，此时gridDim.z=1、blockDim.z=1，blockIdx.z=0、threadIdx.z=0，带入化简得：
kernel的这种线程层次结构，天然适合vector、matrix等运算，以2维grid+2维block核结构为例，实现一个二维矩阵(N * N)的加法运算，每个线程负责处理每个位置的两个元素相加，代码见下。
// kernel函数定义
__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的线程配置
// 在这里，总线程的数量与矩阵大小一致，为N * N
dim3 blockSize(16, 16); // block的大小
dim3 gridSize(N / blockSize.x, N / blockSize.y); // grid的大小
// kernel函数调用-->多线程并行执行矩阵加法操作
MatAdd<<<gridSize, blockSize>>>(A, B, C);
...
}
2.3 CUDA的内存模型(Memory Model)
CUDA的内存模型分为6类：
- 1.Global Memory，全局内存：速度普通，读写。
- 2.Local Memory，本地内存(其实是全局内存)：速度普通，读写。
- 3.Shared Memory，共享内存：速度快，读写。
- 4.Register，寄存器()：速度“最快”，读写。
- 5.Constant Memory，常量内存：速度快，只读。
- 6.Texture Memory，纹理内存：速度快，只读。
