Nvidia 并行计算架构 CUDA 分析(三)——CUDA架构下运行程序的基础

1.程序的分割

在 CUDA 的架构下，一个程序分为两个部份：Host 端和 Device 端。Host 端是指在 CPU 上执行的部份，而 Device 端则是在GPU上执行的部份。Device 端的程序又称为 "kernel"。通常 Host 端程序会将数据准备好后，复制到GPU的内存中，再由GPU执行 Device 端程序，完成后再由 Host 端程序将结果从GPU的内存中取回。CPU 存取 GPU 内存时只能通过 PCIe 接口，速度有限。

2.CUDA 处理流程

通过前面对程序的分割的描述，CUDA处理程序流程总体来说包含以下步骤：

从系统内存中复制数据到GPU内存；

CPU指令驱动GPU运行；

GPU 的每个CUDA核心并行处理；

GPU 将CUDA处理的最终结果返回到系统的内存。

具体一点，CUDA程序执行的基本流程是： 

分配内存空间和显存空间

初始化内存空间

将要计算的数据从Host内存上复制到GPU内存上

执行kernel计算

将计算后GPU内存上的数据复制到Host内存上

处理复制到Host内存上的数据

3.程序执行的各种组件

在 CUDA 架构下，GPU芯片执行时的最小单位是thread。数个 thread 可以组成一个线程块（block）。一个 block 中的 thread 能存取同一块共享的内存，可以快速进行同步和通信操作。

每一个 block 所能包含的 thread 数目是有限的。执行相同程序的 block，可以组成grid。不同 block 中的 thread 无法存取同一个共享的内存，因此无法直接通信或进行同步。不同的 grid可以执行不同的程序（即 kernel）。 Grid、block 和 thread 的关系，如下图所示：

4.CUDA程序的执行

GPU包含N个流式多处理器（MP，同一个MP中的多个线程块可以并行执行），每个MP具有M个标量处理器（SP），每个MP处理批量的线程块（每个线程仅能被一个MP处理）。每个线程块划分成SIMD的线程组，称为wrap。调度器在各个wrap之间切换，wrap中的包含了连续的，递增的线程ID，当前一个wrap包含了32个线程。

5.Stream创建

CUDA的一个Stream表示一个按特定顺序执行的GPU操作序列。诸如kernel启动、内存拷贝、事件启动和停止等操作可以排序放置到一个Stream中。按指定顺序加入到Stream中的操作也是其执行的顺序。

首先，选择的设备必须支持称为device overlap的能力。支持该特性的GPU具有CUDA kernel操作的同时执行设备和Host内存之间拷贝操作的能力。

如果设备支持overlapping，那么Stream的创建应使用命令cudaStreamCreate()。然后，分配Stream的内存，并以随机整数填充。函数调用cudaMemcpyAsync()发出请求执行一个到由参数stream指定的Stream的内存拷贝。当调用返回后，并不能保证该拷贝操作在同一个Stream中的下一个操作执行前被执行。使用cudaMemcpyAsync()需要使用cudaHostAlloc()。当然，kernel的启动也使用了参数stream。当（上述）循环终止，可能仍然会有一点工作队列等候GPU来完成。需要与Host进行同步，以保证任务的完成。在同步操作完成之后，定时器可以停止，内存可以被清除掉。在应用程序退出之前，Stream需要销毁。

6.Kernel的启动参数

任何对_global_函数的调用必须指定执行配置参数。执行配置参数定义了网格的维数以及在设备上执行该函数的线程块以及与其相关的Stream。当使用runtime API时，执行配置参数是通过在函数名和圆括号括起的参数列表之间插入<<>>形式的表达式的方式实现的。

Dg的类型是dim3，指定了网格的维数和大小，如Dg.x * Dg.y的值等于启动的线程块个数，Dg.z必须等于1。

Db的类型是dim3，指定了每个线程块的维数和大小，如Db.x * Db.y* Db.z等于每个线程块中的线程数。

Ns的类型是size_t，指定了除静态分配给这次函数调用的内存外，在共享内存中动态分配给每个线程块的字节数（Bytes）；该动态分配的内存供任意以外部数组（external array）的形式声明的变量使用；Ns是一个可选参数，默认值是0。

S的类型是cudaStream_t，指定了相关的Stream，S是一个可选参数，默认值是0。

7.计算网格

计算网格是由线程块组成的网格。每个线程都执行该kernel，应用程序指定了网格和线程块的维数，网格的布局可以是1维、2维或3维的。每个线程块有一个唯一的线程块ID，线程块中的每个线程具有唯一的线程ID。同一个线程块中的线程可以协同访问共享内存，实现线程之间的通信和同步。每个线程块最多可以包含的线程的个数为512个，线程块中的线程以32个线程为一组的Wrap的方式进行分时调度（具体内容可以参考《GPU调研文档终稿》中线程调度部分的内容），每个线程在数据的不同部分并行地执行相同的操作。

8.CUDA程序调度

线程的 CUDA 层次结构映射到 GPU 上处理器的层次结构；一个 GPU 执行一个或多个 Grid； Grid调度线程块在多处理器上执行。流式多处理器（ Fermi 上的 SM / Kepler 上的 SMX）执行一个或多个线程块；SMX 中的 CUDA Core和其他执行单元执行线程指令。每个多处理器一批接一批地处理块批次。一个块仅在一个多处理器内处理。在一个批次内并被一个多处理器处理的线程块称为活动（Active）线程块。SMX 以 32 个线程为一组的形式执行，这 32 个线程叫做 Warp，并以SIMD方式由多处理器执行。并且同一个warp内是严格串行的，因此在warp内是无须同步的。活动的wrap（活动线程块中的wrap）是分时调度的：线程调度器定期从一个wrap切换到另一个wrap，以便最大化地利用多处理器的计算资源。Wrap的发射（issue）顺序是不确定，但是可以实现同步。线程块网格中的线程块的发射也是不确定的，且块之间没有同步机制，在网格执行期间，同一网格的两个不同线程块中的线程无法通过全局内存互相通信。

在执行时，GPU 的任务分配单元（global block scheduler）将网格分配到GPU 芯片上。启动CUDA 内核时，需要将网格信息从CPU 传输到GPU。任务分配单元根据这些信息将块分配到SM 上。任务分配单元使用的是轮询策略：轮询查看SM 是否还有足够的资源来执行新的块，如果有则给SM 分配一个新的块，如果没有则查看下一个SM。决定能否分配的因素有：每个块使用的共享内存数量，每个块使用的寄存器数量，以及其它的一些限制条件。任务分配单元在SM 的任务分配中保持平衡，但是程序员可以通过更改块内线程数，每个线程使用的寄存器数和共享存储器数来隐式的控制，从而保证SM 之间的任务均衡。