CPU+GPU异构计算

GPU计算并不是指单独的GPU计算，而是指CPU+GPU的异构计算。一块单独的GPU是无法独立的完成所有计算任务的，它必须在CPU的调度下才能完成特定的任务。CPU更适合进行逻辑复杂低并行的程序，GPU更适合逻辑简单高并行的任务。这主要是由于两种处理器的硬件特性不同。

• 左图：CPU的结构，逻辑控制单元占据了大部分的空间，逻辑计算部分相对较小，片上还包括缓存部分，DRAM 是内存，一般不在片上，CPU通过总线访问内存。
• 右图：GPU的结构，绿色代表了逻辑计算单元，占据了GPU的大部分空间，而控制单元和缓存相对较小。这一组逻辑控制单元、缓存、计算单元，相当于一个完整的计算核心，称为SM。一个GPU包含多个SM，SM越多，GPU的计算能力越强。
• CPU和GPU通过PCIe总线链接，用于传输指令和数据。
• CPU线程是重量级实体，操作系统交替执行线程，线程上下文切换花销很大
• GPU线程是轻量级的，GPU应用一般包含成千上万的线程，多数在排队状态，线程之间切换基本没有开销。

CUDA：一种异构计算平台

CUDA平台不是单单指软件或者硬件，而是建立在Nvidia GPU上的一整套平台，并扩展出多语言支持

CUDA 提供了两层 API给程序员使用，即 CUDA 驱动（driver）API 和 CUDA 运行时（runtime）API。其中，CUDA 驱动 API 是更加底层的 API，它为程序员提供了更为灵活的编程接口；CUDA 运行时 API 是在 CUDA 驱动 API 的基础上构建的一个更为高级的 API，更容易使用。这两种 API 在性能上几乎没有差别。从程序的可读性来看，使用 CUDA 运行时 API 是更好的选择。这两种API是互斥的，也就是你只能用一个，两者之间的函数不可以混合调用，只能用其中的一个库。

一个CUDA应用通常可以分解为两部分，

• CPU 主机端代码
• GPU 设备端代码

CUDA nvcc编译器会自动分离你代码里面的不同部分，如图中主机代码用C写成，使用本地的C语言编译器编译，设备端代码，也就是核函数，用CUDA C编写，通过nvcc编译，链接阶段，在内核程序调用或者GPU设备操作时，添加运行时库。

nvcc 是从LLVM开源编译系统为基础开发的。

编程模型

“Hello World!”

/*
*hello_world.cu
*/
#include<stdio.h>
__global__ void hello_world(void)
{
  printf("GPU: Hello world!\n");
}
int main(int argc,char **argv)
{
  printf("CPU: Hello world!\n");
  hello_world<<<1,10>>>();
  cudaDeviceReset();//if no this line ,it can not output hello world from gpu
  return 0;
}

简单介绍其中几个关键字

__global__

CUDA 中的核函数与 C++ 中的函数是类似的，但一个显著的差别是：它必须被限定词

（qualifier）_global_ 修饰。其中 global 前后是双下划线。另外，核函数的返回类型必

须是空类型，即 void。

hello_world<<<1,10>>>();

这句话C语言中没有’<<<>>>’是对设备进行配置的参数，也是CUDA扩展出来的部分。

cudaDeviceReset();

这句话如果没有，则不能正常的运行，因为这句话包含了隐式同步，GPU和CPU执行程序是异步的，核函数调用后成立刻会到主机线程继续，而不管GPU端核函数是否执行完毕，所以上面的程序就是GPU刚开始执行，CPU已经退出程序了，所以我们要等GPU执行完了，再退出主机线程。

一般CUDA程序分成下面这些步骤：

1. 分配GPU内存
2. 拷贝内存到设备
3. 调用CUDA内核函数来执行计算
4. 把计算完成数据拷贝回主机端
5. 内存销毁

CUDA中的线程组织

在 CUDA 编程模型中，线程是进行计算或内存操作的最低级别的抽象。一个 GPU 往往有几千个计算核心，而总的线程数必须至少等于计算核心数时才有可能充分利用 GPU 中的全部计算资源。实际上，总的线程数大于计算核心数时才能更充分地利用 GPU 中的计算资源，因为这会让计算和内存访问之间及不同的计算之间合理地重叠，从而减小计算核心空闲的时间。

一个核函数只能有一个grid，一个grid可以有很多个块，每个块可以有很多的线程（目前最大为1024个线程），这种分层的组织结构使得我们的并行过程更加自如灵活：

一个线程块block中的线程可以完成下述协作：

• 同步
• 共享内存

不同块内线程不能相互影响！他们是物理隔离的！

接下来就是给每个线程一个编号了，我们知道每个线程都执行同样的一段串行代码，那么怎么让这段相同的代码对应不同的数据呢？首先第一步就是让这些线程彼此区分开，才能对应到相应从线程，使得这些线程也能区分自己的数据。如果线程本身没有任何标记，那么没办法确认其行为。
依靠下面两个内置结构体确定线程标号：

• blockIdx（线程块在线程网格内的位置索引）
• threadIdx（线程在线程块内的位置索引）

注意这里的Idx是index的缩写（我之前一直以为是identity x的缩写），这两个内置结构体基于 uint3 定义，包含三个无符号整数的结构，通过三个字段来指定：

• blockIdx.x
• blockIdx.y
• blockIdx.z
• threadIdx.x
• threadIdx.y
• threadIdx.z

上面这两个是坐标，当然我们要有同样对应的两个结构体来保存其范围，也就是blockIdx中三个字段的范围threadIdx中三个字段的范围：

• blockDim
• gridDim

他们是dim3类型(基于uint3定义的数据结构)的变量，也包含三个字段x,y,z.

• blockDim.x
• blockDim.y
• blockDim.z

网格和块的维度可以是一维、二维、三维的。CUDA 中对能够定义的网格大小和线程块大小做了限制。对任何从开普勒到图灵架构的 GPU 来说，网格大小在 x、y 和 z 这 3 个方向的最大允许值分别为 2 31−1、65535 和 65535；线程块大小在 x、y 和 z 这 3 个方向的最大允许值分别为 1024、1024 和 64。另外还要求线程块总的大小，即 blockDim.x、blockDim.y 和 blockDim.z 的乘积不能大于 1024。也就是说，不管如何定义，一个线程块最多只能有 1024 个线程。

核函数

核函数就是在CUDA模型上诸多线程中运行的那段串行代码，这段代码在设备上运行，用NVCC编译，产生的机器码是GPU的机器码，所以我们写CUDA程序就是写核函数，第一步我们要确保核函数能正确的运行产生正确的结果，第二优化CUDA程序的部分，无论是优化算法，还是调整内存结构，线程结构都是要调整核函数内的代码，来完成这些优化的。

kernel_name<<<grid,block>>>(argument list);

这个三个尖括号’<<>>’内是对设备代码执行的线程结构的配置（或者简称为对内核进行配置），也就是我们上一篇中提到的线程结构中的网格，块。 通过指定grid和block的维度，我们可以配置：

• 内核中线程的数目
• 内核中使用的线程布局

我们可以使用dim3类型的grid维度和block维度配置内核，也可以使用int类型的变量，或者常量直接初始化：

kernel_name<<<4,8>>>(argument list);

上面这条指令的线程布局是：

执行模型

下图从逻辑角度和硬件角度描述了CUDA编程模型对应的组件。

GPU结构

GPU架构是围绕一个流式多处理器（SM）的扩展阵列搭建的。通过复制这种结构来实现GPU的硬件并行。

上图包括关键组件：

• CUDA核心
• 共享内存/一级缓存
• 寄存器文件
• 加载/存储单元
• 特殊功能单元
• 线程束调度器

GPU中每个SM都能支持数百个线程并发执行，每个GPU通常有多个SM，当一个核函数的网格被启动的时候，多个block会被同时分配给可用的SM上执行。

注意: 当一个blcok被分配给一个SM后，他就只能在这个SM上执行了，不可能重新分配到其他SM上了，多个线程块可以被分配到同一个SM上。

在SM上同一个块内的多个线程进行线程级别并行，而同一线程内，指令利用指令级并行将单个线程处理成流水线。

 

线程束

线程束是最小的执行单位，通常由32个线程组成，是从机器的角度，在某时刻T，SM上只执行一个线程束，也就是32个线程在同时同步执行，线程束中的每个线程执行同一条指令

SIMD vs SIMT

单指令多数据的执行属于向量机，比如我们有四个数字要加上四个数字，那么我们可以用这种单指令多数据的指令来一次完成本来要做四次的运算。这种机制的问题就是过于死板，不允许每个分支有不同的操作，所有分支必须同时执行相同的指令，必须执行没有例外。
相比之下单指令多线程SIMT就更加灵活了，虽然两者都是将相同指令广播给多个执行单元，但是SIMT的某些线程可以选择不执行，也就是说同一时刻所有线程被分配给相同的指令，SIMD规定所有人必须执行，而SIMT则规定有些人可以根据需要不执行，这样SIMT就保证了线程级别的并行，而SIMD更像是指令级别的并行。
SIMT包括以下SIMD不具有的关键特性：

1. 每个线程都有自己的指令地址计数器
2. 每个线程都有自己的寄存器状态
3. 每个线程可以有一个独立的执行路径

而上面这三个特性在编程模型可用的方式就是给每个线程一个唯一的标号（blckIdx,threadIdx），并且这三个特性保证了各线程之间的独立

同步

cuda中的同步主要指的是线程块内的线程之间的同步。块级别的就是同一个块内的线程会同时停止在某个设定的位置，用

__syncthread();

可扩展性

对于不同的GPU硬件，差异在于SM的数量，我们将kernel划分为很多block，每个block划分的前提是每个块的执行顺序不影响最终的结果。

内存组织

GPU中有多种类型的内存，每种内存有不同的容量和延迟。

各类内存的分类和特征

内存组织示意图

全局内存

全局内存是核函数中所有的线程都可以访问的内存结构，该内存的特点就是，容量大，延迟高，对全部线程可见。全局内存的主要角色是为核函数提供数据，并在主机与设备及设备与设备之间传递数据。为了获得高性能的计算程序，尽量减少主机和设备之间的数据传输。

使用时，首先在主机端为全局内存变量分配设备内存。

cudaMalloc()

可以用 cudaMemcpy 函数将主机的数据复制到全局内存，或者反过来。M为需传输的字节数。

//主机到设备
cudaMemcpy(d_x, h_x, M, cudaMemcpyHostToDevice);
//设备到主机
cudaMemcpy(h_z, d_z, M, cudaMemcpyDeviceToHost);
//设备到设备
cudaMemcpy(d_x, d_y, M, cudaMemcpyDeviceToDevice);

常量内存

常量内存（constant memory）是有常量缓存的全局内存，数量有限，一共仅有 64 KB。它的可见范围和生命周期与全局内存一样。不同的是，常量内存仅可读、不可写。由于有缓存，常量内存的访问速度比全局内存高，但得到高访问速度的前提是一个线程束中的线程（一个线程块中相邻的 32 个线程）要读取相同的常量内存数据。

纹理内存和表面内存

纹理内存（texture memory）和表面内存（surface memory）类似于常量内存，也是一

种具有缓存的全局内存，有相同的可见范围和生命周期，而且一般仅可读（表面内存也可

写）。不同的是，纹理内存和表面内存容量更大，而且使用方式和常量内存也不一样。

寄存器

在核函数中定义的不加任何限定符的变量一般来说就存放于寄存器（register）中。核函数中定义的不加任何限定符的数组有可能存放于寄存器中，但也有可能存放于局部内存中。另外，以前提到过的各种内建变量，如 gridDim、blockDim、blockIdx、threadIdx 及 warpSize 都保存在特殊的寄存器中。在核函数中访问这些内建变量是很高效的。

寄存器变量仅仅被一个线程可见。也就是说，每一个线程都有一个变量 n 的副本。虽然在核函数的代码中用了这同一个变量名，但是不同的线程中该寄存器变量的值是可以不同的。每个线程都只能对它的副本进行读写。寄存器的生命周期也与所属线程的生命周期一致，从定义它开始，到线程消失时结束。寄存器内存在芯片上（on-chip），是所有内存中访问速度最高的，但是其数量也很有限。

局部内存

我们还没有用过局部内存（local memory），但从用法上看，局部内存和寄存器几乎一样。核函数中定义的不加任何限定符的变量有可能在寄存器中，也有可能在局部内存中。寄存器中放不下的变量，以及索引值不能在编译时就确定的数组，都有可能放在局部内存中。这种判断是由编译器自动做的。

虽然局部内存在用法上类似于寄存器，但从硬件来看，局部内存只是全局内存的一部分。所以，局部内存的延迟也很高。每个线程最多能使用高达 512 KB 的局部内存，但使用过多会降低程序的性能。

共享内存

共享内存和寄存器类似，存在于芯片上，具有仅次于寄存器的读写速度，数量也有限。不同于寄存器的是，共享内存对整个线程块可见，其生命周期也与整个线程块一致。也就是说，每个线程块拥有一个共享内存变量的副本。共享内存变量的值在不同的线程块中可以不同。一个线程块中的所有线程都可以访问该线程块的共享内存变量副本，但是不能访问其他线程块的共享内存变量副本。共享内存的主要作用是减少对全局内存的访问，或者改善对全局内存的访问模式。

L1 和 L2 缓存

从费米架构开始，有了 SM 层次的 L1 缓存（一级缓存）和设备（一个设备有多个 SM）层次的 L2 缓存（二级缓存）。它们主要用来缓存全局内存和局部内存的访问，减少延迟。从编程的角度来看，共享内存是可编程的缓存（共享内存的使用完全由用户操控），而 L1 和 L2 缓存是不可编程的缓存（用户最多能引导编译器做一些选择）。

全局内存和共享内存的合理使用

全局内存的合理使用

在启用了 L1 缓存的情况下，对全局内存的读取将首先尝试经过 L1 缓存；如果未中，则接着尝试经过 L2 缓存；如果再次未中，则直接从 DRAM 读取。一次数据传输处理的数据量在默认情况下是 32 字节。

关于全局内存的访问模式，有合并（coalesced）与非合并（uncoalesced）之分。合并访问指的是一个线程束对全局内存的一次访问请求（读或者写）导致最少数量的数据传输，否则称访问是非合并的。

为简单起见，我们从全局内存拷贝时忽略L1和L2缓存，一个线程束请求32个单精度浮点数，每个浮点数4字节，共128字节。如果满足合并访问条件，128/32=4次传输就可完成访问，传输的数据均为线程束所需要的。下面是常见的数据访问模式举例：

• 合并访问
• 顺序的合并访问
• 乱序的合并访问
• 非合并访问
• 不对齐的非合并访问（地址32位对齐）
• 跨越式非合并访问
• 广播式的非合并访问

共享内存的合理使用

共享内存是一种可被程序员直接操控的缓存，主要作用有两个：一个是减少核函数中对全局内存的访问次数，实现高效的线程块内部的通信，另一个是提高全局内存访问的合并度。

关于共享内存，有一个内存 bank 的概念值得注意。为了获得高的内存带宽，共享内存在物理上被分为 32 个（刚好等于一个线程束中的线程数目，即内建变量 warpSize 的值）同样宽度的、能被同时访问的内存 bank。我们可以将 32 个 bank 从 0 到 31 编号。在每一个 bank 中，又可以对其中的内存地址从 0 开始编号。为方便起见，我们将所有 bank 中编号为 0 的内存称为第一层内存；将所有 bank 中编号为 1 的内存称为第二层内存。每个 bank 的宽度为 4 字节。当同一线程束内的多个线程试图访问同一个 bank 中不同层的数据时，就会发生 bank 冲突。

通常可以用改变共享内存数组大小的方式来消除或减轻共享内存的 bank 冲突。例如，将上述核函数中的共享内存定义修改为如下：

__shared__ real S[TILE_DIM][TILE_DIM + 1];