计算机架构分类

Flynn于1972年提出了计算平台的Flynn分类法，主要根据指令流和数据流来分类，共分为四种类型的计算平台，如下图所示：

单指令流单数据流机器（SISD）

SISD机器是一种传统的串行计算机，它的硬件不支持任何形式的并行计算，所有的指令都是串行执行。并且在某个时钟周期内，CPU只能处理一个数据流。因此这种机器被称作单指令流单数据流机器。早期的计算机都是SISD机器，如冯诺.依曼架构，如IBM PC机，早期的巨型机和许多8位的家用机等。

单指令流多数据流机器（SIMD）

SIMD是采用一个指令流处理多个数据流。这类机器在数字信号处理、图像处理、以及多媒体信息处理等领域非常有效。GPU就是这样的架构。

多指令流单数据流机器（MISD）

MISD是采用多个指令流来处理单个数据流。由于实际情况中，采用多指令流处理多数据流才是更有效的方法，因此MISD只是作为理论模型出现，没有投入到实际应用之中。

多指令流多数据流机器（MIMD）

MIMD机器可以同时执行多个指令流，这些指令流分别对不同数据流进行操作。最新的多核计算平台就属于MIMD的范畴，例如Intel和AMD的双核处理器等都属于MIMD。

CPU架构和GPU架构

通用的cpu是SISD架构的，而GPU架构是SIMD的。

比如，有以下指令：

对于没有SIMD的处理单元，需要4条指令将4个float数值相加，汇编伪代码如下：

但有了SIMD技术，只需一条指令即可处理完：

SIMT（Single Instruction Multiple Threads，单指令多线程）是SIMD的升级版，可对GPU中单个SM中的多个Core同时处理同一指令，并且每个Core存取的数据可以是不同的。

上述指令会被同时送入在单个SM中被编组的所有Core中，同时执行运算，但a、b 、c的值可以不一样：

GPU架构分析

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

上图包括关键组件：

CUDA核心

共享内存/一级缓存

寄存器文件

加载/存储单元

特殊功能单元

线程束调度器

Fermi 架构

Fermi架构是第一个完整的GPU架构，所以了解这个架构是非常有必要的。

Fermi架构逻辑图如上，具体数据如下:

512个加速核心，CUDA核

每个CUDA核心都有一个全流水线的整数算数逻辑单元ALU，和一个浮点数运算单元FPU

CUDA核被组织到16个SM上

6个384-bits的GDDR5 的内存接口

支持6G的全局机栽内存

GigaThread疫情，分配线程块到SM线程束调度器上

768KB的二级缓存，被所有SM共享

而SM则包括下面这些资源：

执行单元（CUDA核）

调度线程束的调度器和调度单元

共享内存，寄存器文件和一级缓存

每个多处理器SM有16个加载/存储单元所以每个时钟周期内有16个线程（半个线程束）计算源地址和目的地址

特殊功能单元SFU执行固有指令，如正弦，余弦，平方根和插值，SFU在每个时钟周期内的每个线程上执行一个固有指令。

每个SM有两个线程束调度器，和两个指令调度单元，当一个线程块被指定给一个SM时，线程块内的所有线程被分成线程束，两个线程束选择其中两个线程束，在用指令调度器存储两个线程束要执行的指令

GPU编程模型

异构架构

GPU作为外设一般插在主板的PCIe卡口上，运行程序的时候，CPU像是一个控制者，指挥两台Titan完成工作后进行汇总，和下一步工作安排，所以CPU我们可以把它看做一个指挥者，主机端，host，而完成大量计算的GPU是我们的计算设备，device。

上面这张图能大致反应CPU和GPU的架构不同。

左图：一个四核CPU一般有四个ALU，ALU是完成逻辑计算的核心，也是我们平时说四核八核的核，控制单元，缓存也在片上，DRAM是内存，一般不在片上，CPU通过总线访问内存。

右图：GPU，绿色小方块是ALU，我们注意红色框内的部分SM，这一组ALU公用一个Control单元和Cache，这个部分相当于一个完整的多核CPU，但是不同的是ALU多了，control部分变小，可见计算能力提升了，控制能力减弱了，所以对于控制（逻辑）复杂的程序，一个GPU的SM是没办法和CPU比较的，但是对了逻辑简单，数据量大的任务，GPU更高效，一个GPU有好多个SM，而且越来越多。

CPU和GPU之间通过PCIe总线连接，用于传递指令和数据。

一个异构应用包含两种以上架构，所以代码也包括不止一部分：

主机代码

设备代码

主机代码在主机端运行，被编译成主机架构的机器码，设备端的在设备上执行，被编译成设备架构的机器码，所以主机端的机器码和设备端的机器码是隔离的，自己执行自己的，没办法交换执行。

主机端代码主要是控制设备，完成数据传输等控制类工作，设备端主要的任务就是计算。

GPU基本概念

硬件层面：

SP(Streaming Processor):流处理器， 是GPU最基本的处理单元，在fermi架构开始被叫做CUDA core。

SM(Streaming MultiProcessor): GPU中最小的执行单元。一个SM由多个CUDA core组成。

软件层面：

Kernel：在GPU上执行的函数。

thread：一个CUDA的并行程序会被以许多个threads来执行。

Block：数个threads会被群组成一个block，同一个block中的threads可以同步，也可以通过shared memory通信。

Grid：Kernel 所对应的线程集合，由多个Block组成。

Warp：是最基本的执行单元，一个warp包含32个并行thread，这些thread以不同数据资源执行相同的指令。

GPU编程模型

在CUDA代码中，用__global__修饰的是核函数(Kernel)，即在GPU设备中执行的代码。

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

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

同步

共享内存

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

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

依靠下面两个内置结构体确定线程标号：

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

threadIdx（线程在线程块内的位置索引）

网格和块的维度一般是二维和三维的，也就是说一个网格通常被分成二维的块，而每个块常被分成三维的线程。

CUDA编程模型从逻辑角度和硬件角度的对应关系：

一个线程在一个CUDA Core执行，一个线程块被分配到一个SM上面执行，一个Grid在GPU设备执行。

线程束

上述我们知道了，CUDA代码在执行的时候会被划分成很多个块，然后不同的块可以被调度到不同的SM上面执行。而在GPU物理单元的调度最小单元是线程束。线程束是SM中基本的执行单元，线程块会被分配到某个SM上，并划分成多个线程束，每个线程束一般是32个线程，同一个线程束中的线程执行相同的指令。

即线程块中的线程会以32个线程划分成不同的线程束，然后每一个SM上会有线程束调度器，将SM上的线程束调度到执行单元上执行。

SM上线程束的调度：

线程束和线程块的区别

线程块Block

线程块是一个逻辑产物，写程序时是用二维三维的方式进行，所以线程划分成多维。

在块中，每个线程有唯一编号(一维二维三维都有)

在网格中，每个线程块也有唯一的编号(一维二维三维都有)

线程束Warp

线程束是SM中的基本执行单元，是一个物理定义的产物，每一个线程束按照SIMT执行。

当线程块分配到一个SM上以后，以32为单位划分为多个线程束，向上取整。

线程束是通过线性地址来划分线程块的。

GPU内存模型

典型的内存层次结构图如下所示，从上到下分别是寄存器、缓存、主存和磁盘，从上往下访问速度是越来越慢但存储空间是越来越大。

CUDA内存模型相对于CPU来说那是相当丰富了，GPU上的内存设备有：

寄存器

共享内存

本地内存

常量内存

纹理内存

全局内存

上述各种都有自己的作用域，生命周期和缓存行为。CUDA中每个线程都有自己的私有的本地内存；线程块有自己的共享内存，对线程块内所有线程可见；所有线程都能访问读取常量内存和纹理内存，但是不能写，因为他们是只读的；全局内存，常量内存和纹理内存空间有不同的用途。对于一个应用来说，全局内存，常量内存和纹理内存有相同的生命周期。

寄存器

寄存器无论是在CPU还是在GPU都是速度最快的内存空间，但是和CPU不同的是GPU的寄存器储量要多一些，而且当我们在核函数内不加修饰的声明一个变量，此变量就存储在寄存器中，但是CPU运行的程序有些不同，只有当前在计算的变量存储在寄存器中，其余在主存中，使用时传输至寄存器。在核函数中定义的有常数长度的数组也是在寄存器中分配地址的。

寄存器对于每个线程是私有的，寄存器通常保存被频繁使用的私有变量，注意这里的变量也一定不能使共有的，不然的话彼此之间不可见，就会导致大家同时改变一个变量而互相不知道，寄存器变量的声明周期和核函数一致，从开始运行到运行结束，执行完毕后，寄存器就不能访问了。

寄存器是SM中的稀缺资源，Fermi架构中每个线程最多63个寄存器。Kepler结构扩展到255个寄存器，一个线程如果使用更少的寄存器，那么就会有更多的常驻线程块，SM上并发的线程块越多，效率越高，性能和使用率也就越高。

那么问题就来了，如果一个线程里面的变量太多，以至于寄存器完全不够呢？这时候寄存器发生溢出，本地内存就会过来帮忙存储多出来的变量，这种情况会对效率产生非常负面的影响，所以，不到万不得已，一定要避免此种情况发生。

本地内存

核函数中符合存储在寄存器中但不能进入被核函数分配的寄存器空间中的变量将存储在本地内存中，编译器可能存放在本地内存中的变量有以下几种：

使用未知索引引用的本地数组

可能会占用大量寄存器空间的较大本地数组或者结构体

任何不满足核函数寄存器限定条件的变量

本地内存实质上是和全局内存一样在同一块存储区域当中的，其访问特点——高延迟，低带宽。对于2.0以上的设备，本地内存存储在每个SM的一级缓存，或者设备的二级缓存上。

共享内存

在核函数中使用__share__修饰符的内存，称为共享内存。每个SM都有一定数量的由线程块分配的共享内存，共享内存是片上内存，跟主存相比，速度要快很多，也即是延迟低，带宽高。其类似于一级缓存，但是可以被编程。

使用共享内存的时候一定要注意，不要因为过度使用共享内存，而导致SM上活跃的线程束减少，也就是说，一个线程块使用的共享内存过多，导致更过的线程块没办法被SM启动，这样影响活跃的线程束数量。

共享内存在核函数内声明，生命周期和线程块一致，线程块运行开始，此块的共享内存被分配，当此块结束，则共享内存被释放。

共享内存是线程块内线程可见的，可以通过**__syncthreads()**同步语句来避免内存竞争问题

共享内存和SM中的一级缓存都使用64KB的片上内存，并且可以静态划分彼此的容量。

常量内存

常量内存驻留在设备内存中，每个SM都有专用的常量内存缓存，常量内存使用**constant**来修饰。

常量内存存在于核函数之外，在kernel函数外声明，即常量内存存在于内存中，并不在片上，常量内容的访问速度也是很快的，这是因为每个SM都有专用的常量内存缓存，会把片外的常量读取到缓存中；对所有的核函数都可见，在Host端进行初始化后，核函数不能再修改。

常量内存并不是在片上的，而是在DRAM上，而其有在片上对应的缓存，其片上缓存就和一级缓存和共享内存一样， 有较低的延迟，但是容量比较小。

纹理内存

纹理内存驻留在设备内存中，在每个SM的只读缓存中缓存，纹理内存是通过指定的缓存访问的全局内存。

从数学的角度，上图中的4个地址并非连续的，在一般的CPU缓存中，这些地址将不会缓存。但由于GPU纹理缓存是专门为了加速这种访问模式而设计的，因此如果在这种情况中使用纹理内存而不是全局内存，那么将会获得性能的提升。

全局内存

位于片外存储体中。容量大、访问延迟高、传输速度较慢。所有数据都必须先传入到这里，或者从这里传出。

一般在主机端代码里定义，也可以在设备端定义，不过需要加修饰符，只要不销毁，是和应用程序同生命周期的。

在CUDA中用修饰符 device 声明变量：

缓存

与CPU缓存类似，GPU缓存不可编程，其行为出厂是时已经设定好了。GPU上有4种缓存：

一级缓存

二级缓存

只读常量缓存

只读纹理缓存

每个SM都有一个一级缓存，所有SM公用一个二级缓存。一级二级缓存的作用都是被用来存储本地内存和全局内存中的数据，也包括寄存器溢出的部分。

每个SM有一个只读常量缓存，只读纹理缓存，它们用于设备内存中提高来自于各自内存空间内的读取性能。

总结