本节首先将介绍图像处理器的发展历史，然后介绍使用CUDA语言进行GPU编程，最后简单介绍GPU具体的硬件实现。
1 Graphics 101 & GPU history
查了一下101是美国一个俚语，意思大概是基础，因此标题的意思是"图形基础以及GPU历史"。
简单回顾一下之前讲到的GPU知识。
GPU是一个多核的芯片，在每个核上有多个ALUs用于执行SIMD，一个核上可以同时执行多个CUDA线程。
我们之前介绍过CPU使用超线程技术在单核上实现线程并行，GPU的单核多线程并行技术有很大的区别，之后会讲到。
同一时代时，GPU的内存比CPU的相对小，cache也少，但带宽快很多(具体原因可看本文最后扩展部分)。
最早期的GPU适用于渲染图形，如3D建模，动画、游戏等。
从以前的课我们知道，从2000年开始CPU已经陷入了瓶颈，因此人们更加热衷于各种专用芯片，如GPU。
早期的GPU编程只用于图形方面的东西，后来，人们意识到很多问题都可以转化为基于图像操作的一系列operation的组合。
如2002年的GPGPU（"general purpose" computation on GPUs）可以做一些科学相关的计算。
2004年斯坦福的研究生创建了一门流编程语言(Brook)，它将GPU硬件抽象为数据并行处理器。Brook编译器将流程序转化为OpenGL命令。 OpenGL是图形标准库。该语言也成为了CUDA语言的基础。
下图展示了GPU架构的演讲历史：
2 CUDA programming abstractions
我们本节介绍CUDA编程语言，下一节介绍具体的硬件实现。
在2007年，CUDA编程语言还没有出现之前，人们必须通过一种固定的pipeline来调用GPU的硬件，非常不方便。
2007年，NVIDIA Tesla架构上引入了CUDA编程语言，人们可以更加方便地使用GPU硬件，具体步骤如下：
- 应用在GPU中申请内存，并把数据从CPU拷贝到GPU内存
- 应用提供GPU一个核函数，该核函数是CPU想让GPU执行的程序。
- 应用告诉GPU通过SPMD的方式运行这个核函数
- CPU说go，然后GPU开始执行
OS不参与GPU上程序的调度，因为OS太慢了，都是GPU自己靠硬件调度的。
在后面的课程中，我们将GPU称作device，将CPU称作host
CUDA编程语言类似于C语言，运行在GPU上并使用了GPU的硬件接口。
CUDA很贴近现代GPU的架构特点，其抽象与GPU硬件比较接近。
OpenCL则类似于CUDA的一个开源版本，但它可以运行在除了NVIDIA的GPU上。本课介绍的CUDA特性，OpenCL也都支持。
建议在学习该课的过程中可以思考以下几个问题，在本文的最后还会再一次思考这些问题。
- Is CUDA a data-parallel programming model?
- Is CUDA an example of the shared address space model?
- Or the message passing model?
- Can you draw analogies to ISPC instances and tasks? What about pthreads?
我们将频繁提到CUDA线程(CUDA Thread)这一个术语，CUDA线程在抽象上与我们接触到的pthread相似，但在硬件实现上有大大地不同。
CUDA thread采用的NVIDIA提出的SIMT(single instruction multi thread)架构，SIMT和SIMD的区别可以看本文扩展部分。
接下来将开始介绍CUDA编程。
CUDA程序是一组具有层次结构的CUDA线程，层次从高到低分别是线程网格（grid）、线程块（block）、线程（thread）。
Grid是一个由block组成的3维矩阵，因此block在grid中的ID是3元组(bi,bj,bk).
block是一个由thread组成的3维矩阵，因此thread中block中的ID是3元组(ti,tj,tk).
在本文中，为了简便，bk和tk都为1，即将Grid和block化简为二维矩阵，如下图右边的代码所示.
如上图右侧所示，CUDA程序分为Host code(上面的代码)和cuda device code(下面的代码)
其中host code指在CPU上执行的C或C++代码，host code中会设置grid，block的3D维度参数, 并调用kernel.
cuda device code则是一些使用__global__修饰的核函数，他们会运行在GPU上。
cuda程序的线程数量是程序中显示指定的，上图中的72=12(Nx) x 6(Ny)，就是总的线程数量
cuda程序中，分为CPU内存和GPU内存，CPU和GPU只能操作他自己的内存，因此需要使用内存拷贝函数在CPU和GPU内存之间进行数据拷贝，如下面代码
int bytes = sizeof(float) * N//需要拷贝的字节数
float* deviceA;//指向gpu内存的指针
cudaMalloc(&deviceA, bytes);//在gpu内存中申请空间，并把地址给deviceA
cudaMemcpy(deviceA, A, bytes, cudaMemcpyHostToDevice);//将A指向的CPU内存的数据拷贝到deviceA所指的GPU内存中
需要注意的是，CPU代码中使用deviceA[i]是不合法的，因为cpu不能直接访问gpu的数据。
GPU内存分为三种，分别是device global memory, per-block shared memory 和per-thread private memory(如果将线程自己的寄存器算上，那就是4种了).
具体区分从名字就可以看出来，第一种是当前device中所有CUDA线程都可以访问的，第二种是只有同一个block的CUDA线程才能访问，最后一种只有当前线程才能访问。从访问速度上来说，第一种最慢，最后一种最快。因此，CUDA编程的时候就需要注意数据的本地性以减少内存延迟，如下面两种1D卷积的写法，其中需要卷积的数据长度为1024 x 1024:
该程序使用了100w个CUDA线程和8K个thread block。
核函数输入的参数input和ouput都指向了global device memory.
如果想声明一个per-block shared memory的变量，则需要使用
__shared__修饰该变量，即图第二个写法的第5行否则，声明的变量处于per-thread private memory。
我们可以看到方法1红色框内对global device memory的读取次数为
128 * 3，而方法二红色框中对global memory的读取次数为130，对per-block内存的读取次数为128 * 3。（但我没找到GPU三种memory的延迟数据，有知道的盆友可以留言区评论一下。）
CUDA中的同步分为3类：
-
__syncthreads- barrier：用于等待该block的所有线程达到这个位置
- 原子操作
- 如
flaot atomicAdd(float *addr, float amount) - 显然，原子操作只用于global内存和 block shared内存
- 如
- host/device 同步
- 隐式的barrier，在kernel函数的返回位置。
总结一下CUDA的抽象
- 层次结构的线程( grid，block，thread)
- 不同的地址空间(global , shared , private memory, 对应per program, per block, pre thread)
- 使用内置的内存拷贝原语在host和device之间拷贝数据
- barrier同步原语(
__syncthreads)- 用于block中不同线程执行进度的同步。
- 用于shared和global变量的原子操作
- private变量只在自己的线程用，不需要使用原子操作。
3 CUDA implementation on modern GPUS
接下来介绍CUDA在GPU上具体的硬件实现细节。
首先简单说一下CUDA编译后的二进制文件汇总有哪些信息：
- 程序的text段（包含指令）
- 然后就是一些程序的所需资源的信息，用于GPU调度器调度block和thread时参考，如下面这些信息
- 128 threads per block
- B bytes of local data per thread
- 130 floats(520 bytes) of shared space per thread block
一个GPU中有多个core。
第二节中计算1D卷积的程序启动了8K个thread block(也叫work)，由GPU的Thread block scheduler（GPU的硬件）根据block所需到资源信息，使用一个动态调度策略，将一个个thread block放置到不同的core。 block之间的执行顺序是随机的。
CUDA的实现中，有个术语叫做warp，它是一组包含32个线程的东西。这32个CUDA线程共享一个指令流(后面会再次解释这句话的)。warp是gpu硬件调度和执行的基本单位。
本文以GTX980为例介绍GPU的架构实现。(一些参数可以在这里找到ASUS GTX 980 Specs | TechPowerUp GPU Database)
一个GTX 980 包含16个SMM( Streaming Multiprocessor, 流多处理器，也称作core. 在GTX 980的Maxwell架构中SM称为SMM，在kepler架构中SM称为SMX), 一个GTX980 包含一个2MB L2 cache 和一个4GB global memory. 如下图：
一个SMM中包含2个48KB的L1 cache， 一个96KB的shared memory, 以及一个可以存储64个warp执行上下文的存储空间(256KB)，
课中的图片和SMM原始架构有点不同，下面再放一张官方的SMM的架构图，对照着看。
可以看到SMM中存在4个小的模块，每个模块包含1个 warp scheduler。Warp scheduler的作用就是管理一系列的warp，在那些满足条件的warp中选中一个来发射指令。因此每个时钟周期，一个SMM可以从64个warp中选4个同时运行。
每个模块中， 还有2个dispatch uint，每一个dispatch unit可以分派1条指令，因此每个模块可以从当前warp中同时执行2条指令。这里需要注意的是，当wrap进行访存等长延迟的操作时，warp scheduler会调度一个wrap替换这个模块的执行的wrap，从而隐藏延迟
每个模块中还有32个cuda Core, 这里的core指的是课中的ALU，有些教程里面也称为SP（streaming processor）
两个模块共享一个L1 cache，4个模块共享一个shared memory；
以上参考CUDA微架构与指令集（4）-指令发射与warp调度 - 知乎
总结一下GTX 980的属性，
1.1Ghz clock
16个SMM
16 x 128 = 2048 SIMD mul-add ALUs
计算能力= 1.1Ghz x 2 x 2048 = 4.6TFLOPS,这里的乘以2是因为GPU的ALU时钟频率是GPU其他部分频率的2倍，具体可以见5.3节。
最多16 x 64 = 1024个warp可以交叉执行，即1024 x 32 = 32768个线程
最多16 x 4 = 64个warp可以并行执行,即 64 x 32 = 2048个线程
一个block的线程必须在同一个SMM中执行，这点很好理解，block有per-block shared 内存，该内存在硬件上指SMM内存的shared memory，如果一个block可发散到多个SMM，岂不是要经常搬运各个SMM中搬shared memory？那也太低效了。
最后，将一个在device上运行kernel的例子，该kernel中1000个block，每个block有128个线程，520字节shared memory
该device有2个core，每个core最多交叉运行12个warp，有1.5KB shared memory。
- step1：block0 to core0
- step2 : block1 to core 1
- step3: block3 to core0
- step4: block4 to core 1
- step5: core0 , core1都没shared memory了，等等
- step6: core1中的block1结束计算， block5 to core 1
4 summary
warp 是CUDA的实现细节，指一个block中的一组32个的线程，可以使用32-wide的SIMD执行。这一组32个线程，共享一个指令流，因此他们在遇到分支时会影响程序性能（分支部分可以看本文最后的扩展）。
在GPU的一个core中，保存着多个warp的上下文，每个clock可以选择一个warp进行执行。
写CUDA代码需要知道GPU具体的core数目和每个core支持的block数目，这样才能设计出充分发挥GPU性能的程序，使得线程 block数刚好可以fill gpu。
CUDA中软硬件的对应关系如下图（图来自 cuda 线程调度_weiwei0319的博客-CSDN博客_cuda 调度)
我们可以看到，CUDA的抽象和GPU的硬件是互相对应的，也就是说CUDA中的线程要分为grid，block和thread这三个层级的原因是为了和硬件相对应，这也呼应了本文前部分说的CUDA是一种贴近硬件的语言。
最后，我们思考一下本节开头提出的几个问题，由于时间关系视频中没有回答这些问题，这里我只谈一下我的看法。
- Is CUDA a data-parallel programming model?
- 先回顾一下什么是data parallel 编程模型
- data parallel的基本结构是程序对每个元素执行相同的函数，典型的例子就是SPMD编程
- 因此我认为CUDA是属于data parallel，他会对每个元素执行相同的kernel函数，是符合data parallel的性质的。
- 先回顾一下什么是data parallel 编程模型
- Is CUDA an example of the shared address space model?
- 共享地址空间的概念是指不同线程可以通过共享地址空间进行通信，感觉也是符合的。比如CUDA的gobal memory以及各种同步原语，都基于一个共享地址空间。
- Or the message passing model?
- 不是消息传递。消息传递模型的特点是所有数据需要通过显式地send和recv才能通信。CUDA并没有这种机制。
- Can you draw analogies to ISPC instances and tasks? What about pthreads?
- 从cuda的抽象讲，cuda thread是类似于ISPC中的instances的概念的，
- cuda thread和pthreads有相同，也有不同。
- 相同点在于，他们都是一种可以并行执行的实例，并且通过共享地址空间通信。
- 不同点在于，cuda thread是data parallel模型，每个thread执行的是相同的代码，而pthreads的运行则更加灵活，可以是不同的代码。并且cuda thread是以warp为单位执行和调度的，而pthreads则以单个pthread为单位进行执行和调度。
5 扩展
5.1 显卡中的三种频率
显卡中有三种频率：Core/Shader/Memory Clock
分别是指显卡核心，显卡流处理器(本课程中指ALUs)，显存
核心时钟运行一些处理器层面的功能，如指令decoder
Shader 时钟则是运行独立的处理单元。
在AMD 的 Shader = Core
在NV 的 Shader = Core × 2
因此，本文第三节最后计算FLOPS时需要乘以2。
参考：
显卡的Core/Shader/Memory Clock对使用、计算、温度有怎样的影响 - 计算机技术讨论区 - 中国分布式计算论坛 - Powered by Discuz!
What's the difference between shader clock and core clock? - Graphics Cards - Linus Tech Tips
5.2 CPU和GPU的设计思路是怎么影响芯片设计的？
处理器性能中有2个指标是经常要考虑的：延迟和吞吐量。
所谓延迟，是指从发出指令到最终返回结果中间经历的时间间隔。而所谓吞吐量，就是单位之间内处理的指令的条数。
CPU 在设计时的导向就是减少指令的时延，而GPU在设计时的导向是增加吞吐量。
首先谈一下以减少指令时延为导向的CPU。为了减少指令的时延，CPU的架构有如下几个特点：
- 强劲的算术逻辑单元：实现整形浮点型复杂运算的低延时。
- 多级且大缓存：减少内存延时，加快数据访问速度。
- 复杂的控制器：分支预测，预测执行，数据推行计，乱序执行等技术
GPU架构的特点：
- 高效能的算术逻辑单元：大量长延时但高度流水线化的算术单元，实现高吞吐。
- 小缓存：提高存储的吞吐量，但与CPU相比不需要过分降低内存时延，所以缓存比较小。
- 简单控制器：没有分支预测，没有乱序执行等技术。
总结一下，CPU是典型的以空间换时间的思路，将芯片上的大量空间都给了缓存和复杂的控制器以减少指令执行的延迟。而GPU则是以空间换吞吐量，将芯片上的大量空间都给了逻辑运算单元以实现高吞吐量的并行计算。
单条指令，CPU的执行速度可以是GPU的10倍以上.
单位时间内，GPU执行指令数量可以是CPU的10倍以上。
当然，也因为设计导向的区别，同一时代的GPU的内存带宽会比CPU的内存带宽大好几倍。(因为CPU关注于内存延时，吞吐量相比GPU不大，因此没必要使用高内存带宽)
参考：
GPU底层技术、全球市场格局分析（中）_人工智能_Finovy Cloud_InfoQ写作社区
CUDA架构与应用杂谈 - 知乎
5.3 SIMD和SMIT
SIMT是NVIDIA提出的GPU新概念。二者都通过将同样的指令广播给多个执行单元来实现并行。一个主要的不同就是，SIMD的操作对象是data，而SIMT操作对象是cuda thread，因此SMIT中的thread：
- 每个thread拥有自己的寄存器
- 每个thread拥有自己的private内存
- 每个thread拥有自己的instruction address counter(2017Volta架构之前是一个warp中的线程共享一个PC)
- 每个thread拥有自己的堆栈(2017Volta架构之前是一个warp中的线程共享一个堆栈)
- 每个thread拥有自己的执行路径(2017Volta架构之前好像不可以，因为是共享PC的)
下面是Volta以及之后的架构，warp中的每个线程都有独立的PC和stack, 线程现在可以在 sub-warp 粒度上发散和重新收敛。
大部分thread只是逻辑上并行，并不是所有的thread可以在物理上同时执行。这就导致，同一个block，甚至同一个wrap中的线程可能会有不同步调，有时需要使用__syncthreads同步block中的多个线程。
参考：
CUDA 编程手册系列第四章: 硬件的实现 - NVIDIA 技术博客
新一代 Volta 架构解析 - 知乎
CUDA ---- GPU架构（Fermi、Kepler）_ooMelloo的博客-CSDN博客
5.4 分支条件
下面只讨论Volta之前的架构，它们都是只用
GPU以wrap为单位执行线程，但是仔细想想，每个线程真的一定执行相同的指令吗？程序可不是简简单单一条一条往下走，还有分支和循环等指令会改变程序流向，由于每个线程输入的数据不同，很有可能会进入不同的分支。
(该图来源：GPU并行计算的简单介绍_分支)
上图展示了一个warp里不同的线程随着时间进入不同的分支的情况
Bit Mask可以指定哪些线程干活，哪些空转。
因为SIMT始终是同一条指令的，从寄存器角度看就是PC指针始终是一样的。(volta架构之后就不是这样了，可看5.3节)
在这种情况下， warp 执行时需要通过多条分支路径。warp 按顺序通过这些路径，因此也会增加执行时间。
如上图中cycle 2中，只有执行分支1的线程在运行，其他线程暂停。在cycle 3中，只有执行分支2的的线程运行，其他线程暂停。因此，在分支部分运行的总时间是多个分支耗时的总和。
需要注意的是，分支之后的共同代码还是一个warp中的线程一起执行的。除了分支等不得不分开执行的情况，其他情况中一个warp中的线程一起执行。
由于同一个warp是同一个指令流的，这32个线程的进度都一样。但同一个block的不同warp会在不同的执行进度，因此需要barrier来同步。
当warp中出现发散行为时，warp会串行执行每个分支路径，并禁用其他非活动时的路径中的线程。这会造成性能显著下降。
5.5 其他一些问题
一个grid中最多有多个block？
一个block中最多有多少个thread？
下图是一些相关的限制：
CUDA的thread,block,grid和warp - 知乎
但一个SMM里面64warp，一个block不是最多应该有64 x 32 =2048 个线程吗，为什么才最大是1024？
