0 软件抽象和硬件结构对应关系的例子

把GPU跟一个学校对应起来，学校里有教学楼、操场、食堂，还有老师和学生们；很快有领导（CPU）来检查卫生（需要执行的任务Host程序），因此这个学校的学生们要完成打扫除的工作（Device程序）。

软件抽象资源包括Thread、Warp、Block和Grid

硬件资源包括SP和SM

0.1 软件抽象

Grid对应的是年级

是抽象的划分组织方式

根据年级划分任务，Grid可以分为多个不同的班级

Block对应的是班级

是抽象的划分组织方式

每个班级有若干的同学（线程），可能一个两个不同的年级会出现在同一层楼（SM），或者一层楼只有一个班级，或者没有班级，但是每一层楼的班级最大数量是固定的

Warp对应的是兴趣小组

每个小组有32个学生；（同一时间他们一定是一个班级下的小组）

并且数量固定，即使凑不满这么多学生需要加进来不干活的学生，凑够一个小组

只要求他们有着一样的兴趣爱好（能执行相同的任务）

Thread对应的是学生

一个Thread对应一个SP

每个学生都有个课桌 ，放自己的物品，不能让别人用，表示每个Thread在软件上都有自己的空间（寄存器等）

0.2 硬件资源

SM对应的是教学楼的一个楼层

是实际存在的资源

一个楼层上可以有多个班级，年级和楼层并没有确定的对应关系，一个楼层中可以有很多来自不同的年级的Block

SM中的SP会被分成兴趣小组，承接不同的任务

SP对应的是学生

一个SP对应一个Thread

是实际存在的资源

每个学生都有个课桌 ，放自己的物品，不能让别人用，表示每个SP在硬件上都有自己的空间（local memory + registers)；

在楼层中，有公共的空间（走廊、厕所等），这一层楼的所有同学都可以停留，表示一个SM中有shared memory，这个SM上的Block都可以访问；（shared memory是不是所有的block都可以访问）

学校里的公共区域，比如操场、食堂等，所有同学都可以去运动、吃饭，表示GPU中有一些公共的存储空间供所有的Grid访问。

0.3 执行任务

虽然GPU是并行运行，但也并不是我们理想中所有的Thread一起工作，在打扫卫生时，并不是所有学生一起干活，学生经过老师（这里我们理解为Wrap Scheduler）安排后，分为一组一组的小组，每一个小组都只会做一件一样的事情，如果有人先做完了或者不需要做，那么他也会在旁边等他的组员，处于等待状态idle。

1 GPU不同存储的辨析

1.1 总述

这一点跟CPU比较像，就是存储空间越大，访问速度越慢。

GPU越靠近SM的内存就越快。内存的访问速度从快到慢依次为：

Registers->Caches->Shared Memory->Gloabl Memory（Local Memory）。

1.1.1 GPU存储与CPU存储

CPU的典型存储结构如下：

一般来说，CPU和内存之间的带宽只有数十GB/s。比如对于Intel Xeon E5-2699 v3，内存带宽达到68GB/s（(2133 * 64 / 8)*4 MB/s）:

GPU的存储结构一般如下:

GPU的高速缓存较小，上图的Memory实际上是指GPU卡内部的显存。但是与显存之间的带宽可以达到数百GB/s，比如P40的显存带宽为346GB/s，远远大于CPU的内存带宽，但是，相对于GPU的计算能力，显存仍然是瓶颈所在。

1.1.3 CPU与GPU交互

在现代的异构计算系统中，GPU是以PCIe卡作为CPU的外部设备存在，两者之间通过PCIe总线通信:

 ----------           ------------
|___DRAM___|         |___GDRAM____|
      |                    |
 ----------           ------------
|   CPU    |         |    GPU     |
|__________|         |____________|
      |                    |
  ---------            --------
 |___IO____|---PCIe---|___IO___|

对于PCIe Gen3 x1理论带宽约为1000MB/s，所以对于Gen3 x32的最大带宽为32GB/s，而受限于本身的实现机制，有效带宽往往只有理论值的2/3还低。所以，CPU与GPU之间的通信开销是比较大的。

1.2 Registers

• 寄存器是访问速度最快的空间。
• 当我们在核函数中不加修饰的声明一个变量，那该变量就是寄存器变量，如果在核函数中定义了常数长度的数组，那也会被分配到Registers中；寄存器变量是每个线程私有的，当这个线程的核函数执行完成后，寄存器变量也就不能访问了。
• 寄存器是比较稀缺的资源，空间很小，Fermi架构中每个线程最多63个寄存器，Kepler架构每个线程最多255个寄存器；一个线程中如果使用了比较少的寄存器，那么SM中就会有更多的线程块，GPU并行计算速度也就越快。
• 如果一个线程中变量太多，超出了Registers的空间，这时寄存器就会发生溢出，就需要其他内存（Local Memory）来存储，当然程序的运行速度也会降低。
• 因此，在程序中，对于那种循环操作的变量，我们可以放到寄存器中；同时要尽量减少寄存器的使用数量，这样线程块的数量才能增多，整个程序的运行速度才能更快。

1.3 Local Memory

Local Memory也是每个线程私有的，但却是存储在于Global Memory中的。在核函数中符合存储在寄存器中但不能进入核函数分配的寄存器空间中的变量将被存储在Local Memory中，Local Memory中可能存放的变量有以下几种：

• 使用未知索引的本地数组
• 较大的本地数组或结构体
• 任何不满足核函数寄存器限定条件的变量

1.4 Shared Memory

每个SM中都有共享内存，使用__shared__关键字（CUDA关键字的下划线一般都是两个）定义，共享内存在核函数中声明，生命周期和线程块一致。

同样需要注意的是，SM中共享内存使用太多，会导致SM上活跃的线程数量减少，也会影响程序的运行效率。

数据的共享肯定会导致线程间的竞争，可以通过同步语句来避免内存竞争，同步语句为：

void __syncthreads();

当所有线程都执行到这一步时，才能继续向下执行；频繁调用__syncthreads()也会影响核函数的执行效率。

共享内存因为需要分配给不同的线程所以被分成了不同个Bank，一个Warp中有32个线程，在比较老的GPU中，16个Bank可以同时互相访问，即一条指令就可以让半个Warp同时访问16个Bank，这种并行访问的效率可以极大的提高GPU的效率。比较新的GPU中，一个Warp即32个SP可以同时访问32个Bank，效率又提升了一倍。

下面这个图中：

左边的图每个线程访问一个Bank，不存在内存冲突，通过一个指令即可完成访问所有的访问操作；
中间的图虽然看起来有些乱，但还是一个线程对应一个Bank，也不存在冲突，一个指令即可完成。
右边的图中，存在多个Thread访问一个Bank的情况，如果是读操作，那么GPU底层可以通过广播的方式将数据传给各个Thread，延迟不会很大，但如果是写操作，就必须要等上一个线程写完成后才能进行下一个线程的写操作，延时会比较大。

1.5 Constant Memroy

常量内存驻留在设备内存中，每个SM都有专用的常量内存空间，使用__constant__关键字来声明，可以用来声明一些滤波系数等常量。

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

1.6 Texture Memory

纹理内存的使用并不多，它是为了GPU的显示而设计的，这里不多讲了。纹理内存也是存在于片外。

1.7 Global Memory

全局内存，就是我们常说的显存，就是GDDR的空间，全局内存中的变量，只要不销毁，生命周期和应用程序是一样的。

在访问全局内存时，要求是对齐的，也就是一次要读取指定大小（32、64、128）整数倍字节的内存，数据对齐就意味着传输效率降低，比如我们想读33个字节，但实际操作中，需要读取64字节的空间。

对于Global和Constant，Host可以通过下面的函数访问：

cudaGetSymbolAddress()
cudaGetSymbolSize()
cudaMemcpyToSymbol()
cudaMemcpyFromSymbol()

1.8 GPU缓存

每个SM都有一个一级缓存，所有SM公用一个二级缓存，GPU读操作是可以使用缓存的，但写操作不能被缓存。

L1 Cache：Pascal架构上，L1 Cache和Texture已经合为一体（Unified L1/Texture Cache），作为一个连续缓存供给warp使用。

L2 Cache：用来做Global Memory的缓存，容量大，给整个GPU使用。

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

讲到缓存，我们就必须要提一点，CPU和GPU在缓存上的一个重要差别就是“缓存一致性”（cache coherency） 问题。缓存一致是指一个内存的写操作需要通知所有核的各个级别的缓存，因此，无论何时，所有处理器核看到的内存视图是完全一样的。随着处理器中核数量的增多，这个“通知”的开销迅速增大，使得“缓存一致性”成为限制一个处理器中核数不能太多的一个重要因素。“缓存一致”系统中最坏的情况是，一个内存写操作会强迫每个核的缓存都进行更新，进而每个核都要对相邻的内存单元进行写操作。

CPU遵循“缓存一致”原则，而GPU不是。在GPU中系统不会自动的更新其他核的缓存。所以GPU能扩展到一个芯片内具有大数量的核心。它需要由程序员写清楚每个处理器核输出的各自不同的目标区域。从程序的视角看，这支持一个核仅负责一个输出或者一个小的输出集。

1.9 总结如下：

2 内存与软件硬件的一一对应

2.1 Thread

• 每一个Thread都有自己的local memory和Registers

即每个同学都可以把自己的东西放到自己的课桌上，别的同学不可以使用；

Local Memory，它是每个线程专有的线程，但却是存在于Global Memory中的，结合我们在第0节例子中拿学校和学生举的例子，可以理解为：学生的课桌都放满了，只能在操场里给他再找个地方放东西，所以访问速度是很慢的，但是这部分还是属于他的local memory，别的线程应该是访问不了的。

2.2 Block

• 每一个Block有自己的shared memory，构成Block的所有Thread都可以访问。可以被线程中所有的线程共享，其生命周期与线程块一致

即每个班所在的教室里的走道、讲台等，是这个班里同学们的公共区域，别的班级的同学不能进入；

2.3 Grid

• Grid之间会有Global memory和Cache

所有的Grid都可以访问，即学校里的操场、餐厅等，是全校同学的公共区域，所有年级的同学都可以共享。

所有的thread（包括不同block的thread）都共享一份 global memory、constant memory、和texture memory。所有的线程都可以访问全局内存（Global Memory）

2.4 Warp

每一个时钟周期内，Warp（一个block里面一起运行的thread，其中各个线程对应的数据资源不同，因为指令相同但是数据不同）现在规定的thread数量是32个。一个block中最多含有16个warp。所以一个block中最多含有512个线程。

2.5 其他存储

• 只读内存块：常量内存（Constant Memory）和纹理内存（Texture Memory）
• 每个SM有自己的L1 cache，SM通过L2 cache连接到Global Memory

3 SM中的存储实现细节

GPU基本处理单元是流多处理器，有关处理单元介绍，之前的文章中有叙述。这篇主要讲存储结构

3.1 寄存器

• 流处理器先与寄存器交换数据，寄存器负责存储指令，以及指令计算所需的变量。寄存器速度最高，与核心同频。可以无延迟的让Core执行指令运算。

3.2 L1、L2和SMEM

• L1和SMEM（Shared Memory）其实算作一个东西，或者准确来讲，在一同一块芯片区域当中。只是根据程序需要，决定哪一部分划分成L1缓存，哪一部分划分成Shared Memory，这个操作是动态的。
• L1 cache负责缓存内存地址，而SMEM负责缓存计算所需的变量（包括顶点数据，纹理数据，以及计算引入的中间变量等）。
• L1/SMEM缓存频率非常高，延迟非常低，而且访问这个东西所需要的指令周期是可预测的。这个可预测十分重要，这种情况下算法可以优化其计算和访问SMEM里面数据的策略，比如，遇到数据访问操作，可以切换执行其他计算指令，待SMEM数据被取回到寄存器后，就切回刚刚执行数据访问操作的位置，继续执行后续的指令。
• L1/SMEM每个SM独享一份，不与其他SM共享。如果有共享需求，则这时候应该L2出场。L2相比L1，其频率没那么高，相应的访问其中的数据需要付出一定的延迟代价。当然了，GPU执行当中为了最大化执行效率，都有动态缓存执行状态和切换执行其他指令的策略。
• L2缓存保存了显存的一部分拷贝。在必要情况下（GPU没有独立显存，或者系统也在内存当中划分了一部分区域提供给显存使用），它也可以保存一部分内存的拷贝。所谓拷贝，意味着这是一份连续的显存/内存拷贝（所以指望tex2DLod这样的操作可以减少显存读写以提升性能的可以省省了）。
• 对于CPU而言，多线程访问L2可能是不安全的。所以CPU一般是提供了额外的指令集去确保L2读-改-写操作是安全的。但是现在某些GPU（比如NVidia Fermi系列）就可以保证这套流程是安全的。GPU相比CPU还有一个好处，它往往对一套数据（比如数组，纹理等）访问是并行的，所以在GPU侧进行这个访问操作就可以安全很多。

为了最大化执行效率动态切换执行指令策略。

3.3 显存

我们经常考虑优化一些东西，包括OverDraw，贴图纹理，遮挡剔除，批次合并，不少都是为了照顾显存带宽的。显存虽然大，但是它频率不高，访问还有延时，而且这个延时往往是三位数甚至四位数的GPU核心指令周期。

通常交给GPU计算东西的时候，都是要把数据从内存拷贝到显存，GPU计算完成之后，从显存拷贝回来。当然如果这些数据用于显示，直接划分一块显存区域给帧缓存，然后让显示器读取这块显存即可，这时候不需要往回拷贝。

LAST 参考文献

GPU 初理解 - 简书

GPU架构之处理模块 - 知乎

GPU中的基本概念 - 云+社区 - 腾讯云

CUDA, 软件抽象的幻影背后 之二 | 奇点视觉

CUDA, 软件抽象的幻影背后 | 奇点视觉

GPU编程1–GPU中的基本概念 - 知乎

(3条消息) gpu的单位表示_GPU中的基本概念_weixin_39717121的博客-CSDN博客

CUDA的thread,block,grid和warp - 知乎

GPU编程3–GPU内存深入了解 - 知乎

GPU架构之Hierarchy Memory多级存储 - 知乎

cuda编程(一)：GPU概念与架构 - 知乎

GPU计算 – GPU体系结构及CUDA编程模型

Nvidia GPU架构 - Cuda Core，SM，SP等等傻傻分不清？_咚咚锵的博客-CSDN博客_cuda sm

Fermi威力完美呈现，GeForce GTX 580全球同步评测 - 超能网