前言
《并行编程》系列是学习《Intro to Parallel Programming》过程中所做的笔记记录以及个人一些所思所想。
GPU 与 CPU
衡量一个高性能处理器的时候，采用两个指标。
- 执行时间（Latency）：执行一项任务所花时间，采用时间单位。
- 吞吐量（Throughput）：单位时间完成的任务量。
而非常遗憾的是，这两项指标并不总是一致的，它们通常是矛盾的。比如说：
A地和B地相距4500 KM，从A到B可以有两种选择。一种方法是开跑车，车上乘坐2个人，以200 KM/H的速度开到B地；另一种方法是乘坐客车，车上乘坐着40个人，以50 KM/H的速度开到B地。
| 方案 | Latency(hour) | Throughput(people/hour) |
|---|---|---|
| 开跑车 | 4500 / 200 = 22.5 | 2 / 22.5 = 0.0889 |
| 客车 | 4500 / 50 = 90 | 40 / 90 = 0.444 |
虽然这个例子并不是很合理，但是它展示了 Latency 和 Throughput 的计算方式。
传统的 CPU 设计就是尝试去最优化执行时间，使其在每一项任务上的处理时间都能够达到最优。而 GPU 的设计与 CPU 不同，它的目标最大化吞吐量。因为在计算机图形学中，我们更加关心每秒能处理的像素数量，而不是每个像素需要花多少时间处理，甚至只要每秒能处理的像素数量只要能增加，即便单个像素处理的时间需要增加两倍也是可以接受的。
GPU 设计原则
- GPU 有许多简单的计算单元，它们组合在一块可以执行大量的计算，GPU 通常会牺牲更多的控制力来换取更强大的计算能力。
- GPU 采用显式并行编程模型(课程一大重点)，即编程的时候就是按多处理器的思路进行，而不是假设只有单个处理器，然后将程序交给编译器映射到多个处理上。
- GPU 设计的优化目标是 Throughput 而不是 Latency ，它可以接受单个任务执行时间延长，只要每秒能处理的任务总数能增加，也因此 GPU 适用于以 Throughput 为最重要衡量指标的应用程序中。
CUDA 编程模型
异构型计算机拥有两种不同的处理器，它们是 CPU 和 GPU。如果只是简单地写一个 C 程序，那么它只使用到了 CPU，而如果想要使用 GPU 就要借助 CUDA。CUDA 编程模型允许我们通过一个程序同时对两个处理器进行编程，另外虽然 CUDA 支持多门编程语言，但是本课程中主要使用 C 语言。
CUDA 中普通 C 语言部分的程序会运行在 CPU （也称为"HOST"）中，而另外一部分将在 GPU （相对于"HOST"被称为"DEVICE"）中运行。然后 CUDA 编译器会将 CPU 部分的代码和 GPU 部分的代码分开编译，为每个处理器生成各自的编译结果。
CUDA 将 GPU 当做 CPU 的协处理器（co-processor）来对待，并且假设 HOST 和 DEVICE 各自拥有独立的内存用于存储数据，GPU 通常采用高性能的内存块来作为内存。当谈到 GPU 和 CPU 的关系时，CPU 则处于主导地位。CPU 负责运行主程序，并向 GPU 发送控操作指令。
操作内容包含有：
- 将数据从 CPU 内存中移动到 GPU 内存中。
- 将数据从 GPU 内存中移动到 CPU 内存中。
- 向 GPU 中的内存申请空间。
- 调用 GPU 中的程序，以并行的方式进行运算，这些程序也称为内核，所以 HOST 能够启动 DEIVCE 中的内核。
操作1和2涉及的命令是 cudaMemcpy ，操作3涉及的指令是 cudaMalloc。
CUDA 程序流程
一个典型的 CUDA 程序流程是：
- CPU 为 GPU 申请存储内存空间（cudaMalloc）。
- CPU 将输入数据复制到 GPU 内存中（cudaMemcpy）。
- CPU 启动 GPU 内核处理数据（Kernel launch）。
- CPU 将结果从 GPU 内存中复制回来（cudaMemcpy）。
容易发现，步骤2与4属于数据传输的过程。在程序中我们通常都希望能尽量减少数据传输所消耗的时间，而使更多时间花在计算上。所以对于 I/O 密集型的程序，便不适用于 CUDA 或者 GPU 编程。事实上，成功的 GPU 程序在计算时间与传输通信时间的比率上通常具有较高的值。
GPU 的优点
GPU 擅长处理以下两个事项：
- 高效地启动大量的线程
- 并行地运行大量的线程
举个例子，比如说要对一个长64的数组进行求平方运算。
CPU 的做法
首先是只运行于 CPU 中的做法。
程序中对数组进行遍历，然后依次对每一个元素都执行相同的乘积操作。这些操作是在一个线程中串行执行的，所以该线程将会执行循环64次。
注：此处的线程指的是执行完整代码的一条独立路径。
GPU 的做法
理论知识
之前介绍过，CUDA 的代码需要分成两部分，一部分运行于 CPU，一部分运行于 GPU。GPU 部分所要实现的逻辑很简单，这里是使得输出等于输入的平方，但是这部分并没有说明并行运算的程度（或者是线程数量）。事实上，指明并行运算程度的任务将交给 CPU 进行。所以 CPU 需要为 GPU 分配内存空间，再将数据复制到 GPU 内存中，然后再启动 GPU 计算平方数的内核（此处声明了64个线程）。
同时创建64个线程用于执行平方运算的好处是，每个线程都有一个唯一的线程索引，所有就可以将数组的第 n 个元素分配给第 n 个线程进行处理。
代码实践
定义 GPU 内核代码。
__global__ void square(float *d_out, float *d_in){
// 获取线程索引，将线程索引也作为数组的元素索引
int idx = threadIdx.x;
float f = d_in[idx];
d_out[idx] = f * f;
}
然后通过内核启动语句配置并启动内核。
...
const int ARRAY_SIZE = 64;
...
square<<<1, ARRAY_SIZE>>>(d_out, d_in);
...
所以，这里启动了一个含有64个线程的块，每个线程各自负责计算数组中的一个元素。
配置启动内核的参数
启动内核的语句形式如：
kernel_function<<<blocks_number, thread_per_block>>>(d_out, d_in)
kernel_function 是自定义的内核函数名称。<<<....>>> 则是 CUDA 定义的特殊符号，其中接受两个参数，分别用于启动的块的数量以及每个块的线程数。
例如，SQUARE<<<1, 64>>>(d_out, d_in) 语句启动了一个内核，并指定了内核具有一个块，每个块存在64个线程。
如果需要使用到更多的计算资源，那么便可以通过 <<<...>>> 的参数进行配置。其中关于内核的配置具有两个特点需要知道：
- 一个内核可以同时运行多个块。
- 每个块可以运行多个线程，不过线程具有上限，通常而言在较旧的 GPU 中这个上限为 512，在较新的 GPU 中上限是 1024。
所以当我们想要启动128个线程计算128个数的平方时，代码可以改为 SQUARE<<<1, 128>>>(...) 。那么如果想要想要启动1280个线程呢？这时便有多种策略。比如 SQUARE<<<5, 256>>> 或者 SQUARE<<<10, 128>>> ，但是需要注意不能写成 SQUARE<<<1, 1280>>> ，因为这样超出了最大线程限制。
但是当前的块和线程都是一维的，而如果我们需要处理二维或者是三维结构的数据，这样显然就不方便了，所以 CUDA 也支持二维和三维的块与线程布局方式。
借助 dim3(x, y, z) 函数可以创建指定维度的布局方式。默认情况下，每个维度的值都为1，所以 dim3(w, 1, 1) == dim3(w) == w 。
多维内核
启动内核的最一般形式是：
kernel<<<dim3(bx, by, dz), dim3(tx, ty, tz), shmem>>>(...)
dim3(bx, by, bz) 指定了块的维度。dim3(tx, ty, tz) 指定了每个块的线程维度。shmem 这个参数不太常用，默认为 0，它以字节为单位指定了每个线程块分配的共享内存量，关于该参数的具体使用方法，后续会介绍到。
上图展示了一个二维的布局，其中 thread 是最基本的单位，若干个 thread 组成了一个 block，而若干个 block 组合成一个 grid。CUDA 还提供了多个属性用于实现获取线程索引、块索引和块大小等。
-
threadIdx：获取块内的线程索引，最多具有三个维度x, y, z。 -
blockDim：获取块大小，最多具有三个维度x, y, z。 -
blockIdx：获取块索引，与线程索引一样，最多具有三个维度。 -
gridDim：获取 grid 大小，也是最多具有三个维度x, y, z。
注意事项：这边有一个在编程处理图像的时候经常容易犯错的坑。在处理图像数据的时候，通常会将
grid of blocks定义为与图像大小一致，然后每个block中的thread数定义为1。 这样对于一个分辨率为n*m的图像，则使用了n*m个块，每个块中只有一个线程负责处理当前像素值。但是需要注意的是，在定义gridSize的时候需要定义成dim3(m, n)。但是在编程时出于习惯，我们很可能会写成dim3(n, m)，然而这样是错误的 。这是因为dim3()接受的三个参数依次对应了x, y, z的取值范围，所以在使用图像的宽高参数指定dim3参数的时候应该是先制定宽度再指定高度。否则处理之后的图像很可能出现有一半是黑色的情况。
