GPU 调度策略架构与CUDA运行机制（二）
市面上有很多GPU厂家，他们产品的软硬件架构各不类似，但是核心往往差不多，整明白了一个基本上就可以触类旁通了。针对当前gpu底层的一些架构以及硬件层一些调度策略的话估计大部门人就很难说的上熟悉了，这个不是大家的错，主要是由于Nv gpu的整个生态都是闭源的，以是大家了解起来就会有一些障碍。下面的行文将基于以下三个层面举行论述：CUDA编程模子、GPU 底层硬件架构与硬件层的调度策略、CUDA调度框架。通过行业CUDA标杆做基本的调度管理与分析，在实际的CUDA代码执行过程中需要CPU和GPU的协同工作，在CPU上运行的称为Host程序，在GPU上运行的称为Device程序。比方说对于一个CUDA程序的可以分为两个部门（两者拥有各自的存储器）。
CUDA 最基本的执行单元是线程（Thread），图中每条曲线可视为单个线程，大的网格（Grid）被切分成小的网格，其中包含了很多类似线程数目的块（Block），每个块中的线程独立执行，可以通过当地数据共享实现数据交换同步。因此对于 CUDA 来讲，就可以将题目划分为独立线程块，并行解决的子题目，子题目划分为可以由块内线程并行协作解决。
https://i-blog.csdnimg.cn/direct/4758c08b8d0c45b895d800fe80920f8b.png
CUDA 引入主机端（host）和设备（device）概念，CUDA 程序中既包含主机（host）程序也包含设备（device）程序，host 和 device 之间可以举行通讯，以此来实现数据拷贝，主机负责管理数据和控制程序流程，设备负责执行并行计算任务。在 CUDA 编程中，Kernel 是在 GPU 上并行执行的函数，开发人员编写 Kernel 来描述并行计算任务，然后在主机上调用 Kernel 来在 GPU 上执行计算。
https://i-blog.csdnimg.cn/direct/c2705dd28c114e54baeb6a8eb2efb02c.png
1. CUDA编程模子
1.1 CUDA编程样例对比
代码 cuda_host.cpp 是只利用 CPU 在 host 端实现两个矩阵的加法运算，其中在 CPU 上计算的 kernel 可看作是加法运算函数，代码中包含内存空间的分配和开释。
#include <iostream>
#include <math.h>
#include <sys/time.h>
// function to add the elements of two arrays
void add(int n, float *x, float *y)
{
for (int i = 0; i < n; i++)
y = x + y;
}
int main(void)
{
int N = 1<<25; // 30M elements
float *x = new float;
float *y = new float;
// initialize x and y arrays on the host
for (int i = 0; i < N; i++) {
x = 1.0f;
y = 2.0f;
}
struct timeval t1,t2;
double timeuse;
gettimeofday(&t1,NULL);
// Run kernel on 30M elements on the CPU
add(N, x, y);
// Free memory
delete [] x;
delete [] y;
return 0;
}
在 CUDA 程序架构中，host 代码部门在 CPU 上执行，是普通的 C 代码。当遇到数据并行处置惩罚的部门，CUDA 会将程序编译成 GPU 能执行的程序，并传送到 GPU，这个程序在 CUDA 里称做核(kernel)。device 代码部门在 GPU 上执行，此代码部门在 kernel 上编写(.cu 文件)。
kernel 用 __global__ 符号声明，在调用时需要用 <<<grid, block>>> 来指定 kernel 要执行及结构。代码 cuda_device.cu 是利用 CUDA 编程实现 GPU 计算，代码涉及到 host（CPU）和 device（GPU）相关计算，利用 __global__ 声明将 add 函数转变为 GPU 可执行的 kernel。
#include <iostream>
#include <math.h>
// Kernel function to add the elements of two arrays
// __global__ 变量声明符，作用是将 add 函数变成可以在 GPU 上运行的函数
// __global__ 函数被称为 kernel
__global__
void add(int n, float *x, float *y)
{
for (int i = 0; i < n; i++)
y = x + y;
}
int main(void)
{
int N = 1<<25;
float *x, *y;
// Allocate Unified Memory – accessible from CPU or GPU
// 内存分配，在 GPU 或者 CPU 上统一分配内存
cudaMallocManaged(&x, N*sizeof(float));
cudaMallocManaged(&y, N*sizeof(float));
// initialize x and y arrays on the host
for (int i = 0; i < N; i++) {
x = 1.0f;
y = 2.0f;
}
// Run kernel on 1M elements on the GPU
// execution configuration, 执行配置
add<<<1, 1>>>(N, x, y);
// Wait for GPU to finish before accessing on host
// CPU 需要等待 cuda 上的代码运行完毕，才能对数据进行读取
cudaDeviceSynchronize();
// Free memory
cudaFree(x);
cudaFree(y);
return 0;
}
因此 CUDA 编程流程总结为：
[*] 编写 Kernel 函数描述并行计算任务。
[*] 在主机上配置线程块和网格，将 Kernel 发送到 GPU 执行。
[*] 在主机上处置惩罚数据传输和效果处置惩罚，以及控制程序流程。
为了实现以上并行计算，对应于 GPU 硬件在举行实际计算过程时，CUDA 可以分为 Grid，Block 和 Thread 三个条理结构：
[*] 线程条理结构Ⅰ-Grid：kernel 在 device 上执行时，实际上是启动很多线程，一个 kernel 所启动的所有线程称为一个网格（grid），同一个网格上的线程共享类似的全局内存空间，grid 是线程结构的第一条理。
[*] 线程条理结构Ⅱ-Block：Grid 分为多个线程块（block），一个 block 里面包含很多线程，Block 之间并行执行，并且无法通讯，也没有执行顺序，每个 block 包含共享内存（shared memory），可以共享里面的 Thread。
[*] 线程条理结Ⅲ-Thread：CUDA 并行程序实际上会被多个 threads 执行，多个 threads 会被群组成一个线程 block，同一个 block 中 threads 可以同步，也可以通过 shared memory 通讯。
https://i-blog.csdnimg.cn/direct/0303622428ae4cbe89ac5383c34ac500.png因此 CUDA 和英伟达硬件架构有以下对应关系，从软件侧看到的是线程的执行，对应于硬件上的 CUDA Core，每个线程对应于 CUDA Core，软件方面线程数目是超配的，硬件上 CUDA Core 是固定命量的。Block 线程块只在一个 SM 上通过 Warp 举行调度，一旦在 SM 上调用了 Block 线程块，就会不停保存到执行完 kernel，SM 可以同时生存多个 Block 线程块，多个 SM 组成的 TPC 和 GPC 硬件实现了 GPU 并行计算。
https://i-blog.csdnimg.cn/direct/cf26e780709d41a892807a6e6b0a80b9.png
1.2 CUDA多维度编程
示例：一维数组的求和计算
https://i-blog.csdnimg.cn/blog_migrate/c4fadcfd089bfeb566d293bacb186f25.png
代码中解释的一、二处究竟该怎么来写？
------------------------------------------------------------
线程参数设置 环境1：一维grid，一维block （线程分配）
https://i-blog.csdnimg.cn/blog_migrate/491f6f5f9ccb57b752ca17047158c087.png
grid(1,1,1): block数目=1*1*1
block(length,1,1): thread数目=length*1*1
总thread数目 = (1*1*1)*(length*1*1)
https://i-blog.csdnimg.cn/blog_migrate/6d3e16eb831e4f2bfbe66f35c2609bc8.png
https://i-blog.csdnimg.cn/blog_migrate/04f4c70495fa4171dbfee0226c267170.png
-------------------------------------------------------------------------------------------
线程参数设置 环境二2：一维grid，二维block （线程分配）
https://i-blog.csdnimg.cn/blog_migrate/e11643fc5e2155554a2b46043fb2ca83.png
grid(1,1,1): block数目=1*1*1
block(8,2,1): thread数目=8*2*1
总thread数目 = 16
https://i-blog.csdnimg.cn/blog_migrate/de9b221918d132bbae05059d0e91ca21.png
https://i-blog.csdnimg.cn/blog_migrate/acff7e31650c92095b9e46c6448f0cfa.png
我们肯定要有并行头脑，这里有16个线程，kernel启动后，每个线程都有本身的索引号，好比某个线程位于grid中哪个维度的block（即blockIdx.x,blockIdx.y,blockIdx.z），又位于该block的哪个维度的线程（即threadIdx.x,threadIdx.y,threadIdx.z），利用这些线程索引号映射到对应的数组下标，我们要做的工作就是将保证这些下标不重复（如果重复的话，那就惨了），最初那种一维的计算方式就不行了。因此，通过利用threadIdx，blockDim来举行映射（偏移）。blockDim.x=8,blockDim.y=2
--------------------------------------------------------------------------------------
线程参数设置 环境3：一维grid，一维block （block分配）
https://i-blog.csdnimg.cn/blog_migrate/1a9ceb6caedbe274353bb14bd852b3a5.png
https://i-blog.csdnimg.cn/blog_migrate/16146cb2ffdb679b241fcdbb1d2700f3.png
https://i-blog.csdnimg.cn/blog_migrate/98979b66ed7e5a310fd88ca4425a9a13.png
---------------------------------------------------
线程参数设置环境4: block和thread都分配
https://i-blog.csdnimg.cn/blog_migrate/efd6a8a9b441af5759851093c6814dde.png
https://i-blog.csdnimg.cn/blog_migrate/2cda1af7fe3a3d7d77b1e21261b7569d.png
https://i-blog.csdnimg.cn/blog_migrate/4b45d971c7513f4a0f3b3be73ca6789f.png
------------------------------------------------------------------
线程参数设置 环境5：二维grid，二维thread
https://i-blog.csdnimg.cn/blog_migrate/ad5682fbc765f7cdbb38f34b8b4c6757.png
https://i-blog.csdnimg.cn/blog_migrate/8a27fcec97db4ea1bcc780877af55b37.png
https://i-blog.csdnimg.cn/blog_migrate/d19d7e8f300480947c70ac0011d3607e.png
示例：倒推其线程参数设置
https://i-blog.csdnimg.cn/blog_migrate/89036bdd468f80545ed295758e46f9e5.png
它的线程参数设置是怎样的？线程索引怎么计算？
参数设置为：
https://i-blog.csdnimg.cn/blog_migrate/5f1c6fafcd65e2e2e2a61ffe165e9115.png
总Thread数目： 8*4*1*8*2*1 = 512
一维数组的线程索引计算方法：
https://i-blog.csdnimg.cn/blog_migrate/97b2313216786d529e0d3bc4e8708b09.png
二维数组的线程索引计算方法：
https://i-blog.csdnimg.cn/blog_migrate/0d0fed32d504d45dd0809a076231c242.png
根据CUDA算力不同thread，block，gird在不同维度的巨细是有限制的:
https://i-blog.csdnimg.cn/blog_migrate/cfce397c9736f4d0f3b914a0aa60d178.png
Cuda Wrap的限制：
https://i-blog.csdnimg.cn/blog_migrate/028f0c707fb946423df54bbb4ad7d6eb.png
1.3 Stream
中文翻译为"流"，它主要是通过提拔kernel函数的并发性来提拔整个计算的运行效率。下面我们来看一下在cuda编程模子当中具体是如何利用stream的。
cudaStream_t stream;
for (int i = 0; i < nStreams; i ++)
{
checkCuda(cudaStreamCreate(&stream));
}
for (int i = 0; i < nStreams; i ++)
{
checkCuda(cudaStreamDestroy(stream));
}
上面所展示的是stream的创建和销毁，接下来我们来看一下如何利用stream
for (int i = 0; i < nStreams; i ++)
{
int offset = i * streamSize;
checkCuda(cudaMemcpyAsync(&d_a, &a, streamBytes, cudaMemcpyHostToDevice, stream));
kernel_function<<<streamSize/blockSize, blockSize, 0, stream>>>(d_a, offset);
checkCuda(cudaMemcpyAsync(&a, &d_a, streamBytes, cudaMemcpyDeviceToHost, stream));
}
stream具体用法如上面sample所示，如果你不显示的申请stream的话系统也会有一个default的stream0。大家可以从下面的这张图比力直观地看到两者在执行效率上的区别：
https://i-blog.csdnimg.cn/blog_migrate/ca3ac94ff05efa35d4a3f1c94f7d2f50.png
图3 cuda stream 串行和并行执行
1.4 Graph
2 GPU 底层硬件架构与硬件层的调度策略
2.1 GPU的软件抽象
软件资源的抽象即为GPU的线程模子，可以分为Grid、Block、Thread和Warp。
Grid、Block、Thread是一种软件构造结构，是线程构造的三个条理，并不是硬件的，因此理论上我们可以以任意的维度（一维、二维、三维）去排列Grid，Block，Thread；在硬件上就是一个个的SM或者SP，并没有维度这一说，只是软件上抽象成了具有维度的概念。
thread，block，gird在不同维度的巨细根据算力不同是有限制的：以是在不同CUDA版本或在编译时没有指定架构的环境下，大概CUDA版本也会对thread，block，grid在不同维度的巨细产生影响。
https://i-blog.csdnimg.cn/direct/1290946bce3e451ca731148e80335367.png
2.1.1 Grid（线程网格）
一个Kernel函数对应一个Grid。
一个Grid中会分成若干个Block。同一Grid下的不同Block大概会被分发到不同的SM上执行。
Grid跑在GPU上的时候，大概是独占一个GPU，也大概是多个kernel函数并发占用一个GPU（后面这种实现需要fermi及更新的GPU架构支持）。
2.1.2 Block
数个threads会被群组成一个block，同一个block中的threads可以同步，也可以通过shared memory通讯
2.1.3 Thread
一个CUDA的并行程序会被以很多个Thread来执行
每个Thread中的局域变量被映射到SM的寄存器上，而Thread的执行则由CUDA核心也就是SP来完成。
2.1.4 Warp
Warp是GPU执行程序时的调度单元，同一个Warp里的线程执行类似的指令，即SIMT。
一个SM的CUDA core会分成几个Warp（即CUDA core在SM中分组)，由Warp scheduler负责调度。尽管Warp中的线程从同一程序地址，但大概具有不同的活动，好比分支结构。由于GPU规定同一Warp中所有线程在同一周期执行类似的指令，Warp发散分支过多会导致有效分支减少性能下降。
一个SM同时并发的Warp是有限的，由于资源限制，SM要为每个线程块分配共享内存，也要为每个线程束中的线程分配独立的寄存器，以是SM的配置会影响其所支持的线程块和Warp并发数目。
一个Warp中的线程一定在同一个block中，如果block所含线程数目不是Warp巨细的整数倍，那么多出的那些thread地点的Warp中，会剩余一些inactive的thread，也就是说，即使凑不够Warp整数倍的thread，硬件也会为Warp凑足，只不外那些thread是inactive状态，需要留意的是，即使这部门thread是inactive的，也会消耗SM资源。由于warp的巨细一般为32，以是block所含的thread的巨细一般要设置为32的倍数。
例：如果一个块中有128个线程，那么线程0-31将在一个Warp中，32-63将在下一个Warp中
Warp非常紧张，缘故原由如下：
[*]Warp中的线程是被绑定在一起的。如果Warp中的一个线程沿着if-else块的if侧走，而其他线沿着else侧走，那么实际上所有32条线程都会沿着两侧走。在执行功能上是没有题目的，那些不应该被执行分支的线程会被禁用，因此始终得到精确的效果，但是如果双方都很长，那么性能损失就很紧张。
[*]Warp内的线程（实际上是半纠缠的（self-warp））一起从内存中获取数据，是一起访问共享内存中的同一段数据同一段的。也就是说如果可以确保Warp中的所有线程都从同一段内获取数据，就只需要实现一次内存转换。
[*]如果它们都从随机地址获取数据，那么就需要排队去实现32次内存转换。
2.2 软件抽象和硬件结构的逐一对应关系
硬件结构可以参考之前的一篇文章
https://i-blog.csdnimg.cn/direct/61fd2fa2b6c745cd8ecb7254b741c8c7.png
2.2.1 Block对应于SM
[*]SM上可以同时存在多个Block被执行，这些Block不肯定来自同一个kernel函数。
[*]SM设备有Device Limit，Warp和Block的数目不能超过对应的上限。
[*]除了受到设备界说的限制之外，还受到硬件资源的限制：
[*]SP的寄存器数目
[*]线程块消耗的共享内存量
每个线程会占用肯定命量的寄存器和Shared Memory，因此SM上同时存活的Block数目不应当超过这些硬件资源的限制。由于SM上可以同时有来自不同kernel的Block存在，因此有时候即便SM上剩余资源不足以再容纳一个kernel A的Block，但却仍大概容纳下一个kernel B的Block。
[*]一个线程块的thread只能在一个SM上调度
2.2.2 Block与Thread之间的接洽Warp 对应于 SM与SP之间的接洽
[*]软件抽象里，以为任务分配到Block之后，所有的线程是并行执行的，这只是个逻辑上无懈可击的抽象，究竟上我们不大概对一个任意巨细的Block都给出一个划一巨细的CUDA核心阵列去推动它的并行计算，来真正并行的执行它们。因而有了Warp这个概念。物理上，Block被划分成一块块的warp分别映射到CUDA核心阵列上执行，每一个warp就都可以理解为是一个线程的集装箱，为的是线程数目固定统一可以给他分配统一的硬件资源，每个集装箱只装一种货品，也就是下面同步执行的意思。
[*]现在，CUDA中的Warp都是从threadIdx = 0开始，以threadIdx连续的32个线程为一组划分得到，即便最后剩下的线程不足32个，也将其作为一个Warp。CUDA kernel的配置中，我们经常把Block的size设置为32的整数倍，正是为了让它能够精确划分为整数个Warp（更深刻的缘故原由和存储器访问性能有关，但这种环境下仍然和Warp的size脱不了干系）。
[*]Warp是SM调度和执行的基础概念。Block被划分成32个线程组成的Warp。这样，大量的Warp生存在SM上，等待被调度到CUDA核心阵列去执行。
[*]Warp中的活动线程由Warp Scheduler驱动。每一块SM中有单独的一个或者多个Warp Scheduler（举例：GM204中32个CUDA核心共享一个Warp Scheduler），以及多个CUDA核心。
[*]当一个Warp执行中出现等待（存储器读写延迟等）后，Warp Scheduler就迅速切换到下一个可执行的Warp，对其发送指令直到这个Warp又一次出现等待，周而复始。这就是常说“用多线程掩盖延迟”。SM会从驻留在SM上的所有Warp中举行指令调度。（这里的驻留表示已经可以被执行的Warp，会从这里挑选，这时候挑选出来的Warp能来自于驻留在SM上的任何线程块）。
[*]通常一个SM中的SP会分成几个Warp（也就是SP在SM中是举行分组的，物理上举行的分组）。
[*]同步执行：Warp中的32个SP是一起工作的，执行类似的指令，如果没有这么多thread需要工作，那么这个Warp中的一些SP是不工作的，处于闲置状态。
2.2.3 Thread对应于SP
[*]Thread在SP也就是CUDA Cores上执行
[*]Thread会被分配Register/Local Memory，数据存在这里
[*]SM上的CUDA核心是有限的，它们代表了能够在物理上真正并行的线程数（也就是优化到最佳环境下所能最大到达同一时刻在运行的并行数目）
[*]每一个线程都有本身的寄存器内存和local memory，一个warp中的线程是同时执行的，也就是当举行并行计算时，线程数尽量为32的倍数，如果线程数不上32的倍数的话；如果是1，则warp会生成一个掩码，当一个指令控制器对一个warp单元的线程发送指令时，32个线程中只有一个线程在真正执行，其他31个 历程会进入静默状态。
2.3 软件抽象和硬件结构对应关系的例子
把GPU跟一个学校对应起来，学校里有讲授楼、操场、食堂，另有老师和门生们；很快有向导（CPU）来查抄卫生（需要执行的任务Host程序），因此这个学校的门生们要完成打扫除的工作（Device程序）。
[*]软件抽象资源包罗Thread、Warp、Block和Grid
[*]硬件资源包罗SP和SM
2.3.1 软件抽象
Grid对应的是年级
是抽象的划分构造方式
根据年级划分任务，Grid可以分为多个不同的班级
Block对应的是班级
是抽象的划分构造方式
每个班级有若干的同砚（线程），大概一个两个不同的年级会出如今同一层楼（SM），或者一层楼只有一个班级，或者没有班级，但是每一层楼的班级最大数目是固定的
Warp对应的是兴趣小组
每个小组有32个门生；（同一时间他们肯定是一个班级下的小组）
并且数目固定，即使凑不满这么多门生需要加进来不干活的门生，凑够一个小组
只要求他们有着一样的兴趣爱好（能执行类似的任务）
Thread对应的是门生
一个Thread对应一个SP
每个门生都有个课桌 ，放本身的物品，不能让别人用，表示每个Thread在软件上都有本身的空间（寄存器等）
2.3.2 硬件资源
SM对应的是讲授楼的一个楼层
是实际存在的资源
一个楼层上可以有多个班级，年级和楼层并没有确定的对应关系，一个楼层中可以有很多来自不同的年级的Block
SM中的SP会被分成兴趣小组，承接不同的任务
SP对应的是门生
一个SP对应一个Thread
是实际存在的资源
每个门生都有个课桌 ，放本身的物品，不能让别人用，表示每个SP在硬件上都有本身的空间（local memory + registers)；
在楼层中，有公共的空间（走廊、茅厕等），这一层楼的所有同砚都可以停顿，表示一个SM中有shared memory，这个SM上的Block都可以访问；（shared memory是不是所有的block都可以访问）
学校里的公共地区，好比操场、食堂等，所有同砚都可以去活动、用饭，表示GPU中有一些公共的存储空间供所有的Grid访问。
2.3.3 执行任务
虽然GPU是并行运行，但也并不是我们理想中所有的Thread一起工作，在打扫卫生时，并不是所有门生一起干活，门生经过老师（这里我们理解为Wrap Scheduler）安排后，分为一组一组的小组，每一个小组都只会做一件一样的事情，如果有人先做完了或者不需要做，那么他也会在旁边等他的组员，处于等待状态idle。
4 用多线程掩盖延迟
Global Memory访存延迟可以到达数百个时钟周期，即便是最快的Shared Memory和寄存器在有写后读依赖时也需要数十个时钟周期。这好像和CUDA强大的处置惩罚能力完全相悖。
为什么GPU具有这么高的计算能力？如果连寄存器都这么慢，怎么会有高性能呢？岂非这不会成为最大的瓶颈吗？
由于这个高延迟的开销被掩盖了，掩盖在大量线程之下。更清楚的说，控制单元（Warp Scheduler）在多组线程之间快速切换，当一组线程Warp（一个线程组，在CUDA里叫做Warp）由于访存或其他缘故原由出现等待时，就将其挂起，转而执行另一组线程，GPU的硬件体系允许同时有大量线程存活于GPU的SM（流多处置惩罚器）之中，这种快速切换保证资源的最大利用率——控制单元始终有指令可以发放，执行单元始终有任务可以执行，仍然可以保持最高的指令吞吐，每个单元基本都能保持充实的忙碌。
这就是GPU硬件设计中非常有特色的基本头脑：用多线程掩盖延迟。这一设计区别于CPU的特点是，大量高延迟寄存器取代了少量低延迟寄存器，寄存器的数目保证了可以有大量线程同时存活，且可以在各组线程间快速切换。尽管每个线程是慢的，但庞大的线程数成就了GPU的数据吞吐能力。
下面图片可以分析：GPU用多个Warp掩盖延迟 / 与CPU计算模式的对比
https://i-blog.csdnimg.cn/direct/d145843be18148cab220118afd42aed7.png
GPU由于多个Warp可以快速切换来掩盖延迟，而CPU用快速的寄存器来减小延迟。两者的紧张区别是寄存器数目，CPU的寄存器快但少，因此Context Switch代价高；GPU寄存器多而慢，但寄存器数目保证了线程Context Switch非常快。同时也是由于GPU对高延迟的容忍度比力高，他只追求在长时间内比力稳定的较大吞吐量，而不在意相应时间。
4.1 多少线程才能够掩盖掉常见的延迟呢？
对于GPU，最常见的延迟大概要数寄存器写后读依赖，即一个局域变量被赋值后接着不久又被读取，这时候会产生大约24个时钟周期的延迟。为了掩盖掉这个延迟，我们需要至少24个Warp轮流执行，一个Warp遇到延迟后的空闲时间里执行其余23个Warp，从而保持硬件的忙碌。在Compute Capability 2.0，SM中有32个CUDA核心，平均每周期发射一条指令的环境下，我们需要24 ∗ 32 = 768 24*32 = 76824∗32=768个线程来掩盖延迟。
保持硬件忙碌，用CUDA的术语来说，就是保持充实的Occupancy，这是CUDA程序优化的一个紧张指标。
5 关于现代GPU如此举行软件抽象和硬件设计的一些思考
整个设计逻辑关系我觉得可以归结为如下的环境
[*]目标是实现任务
[*]发现任务具有如下的特性：允许肯定的延迟；需要大吞吐量；有大量同样的操作或者计算
[*]以是设计了现有的硬件体系架构，软件抽象模子
那么为什么这样的计算或者说任务可以被如上所说的硬件软件更好的完成呢？
实在是由于我们是在已知任务特性的环境下（我们实际利用中所需要完成的任务大概率属于这些，或者说这些任务在CPU上比力容易有掣肘），才把结构设计成这样的。
[*]第一方面：
[*]现实世界中应用在大规模数据上的计算，通常都涵盖在这一计算模式之中，因而思量更复杂的模式本质上是不必要的。
好比计算大气的流动，每一点的风速仅仅取决于该点邻域上的密度和压强分布；
好比计算图像的卷积，每一个输出像素都仅是对应源点邻域和一个卷积核的内积。
[*]从这些例子中我们可以看到，除了各个数据单元上举行的计算是一样的，计算中数据之间的相互影响也具有某种“局域性”，一个数据单元上的计算最多需要它某个邻域上的数据。这一点意味着线程之间是弱耦合的，邻近线程之间会有一些共享数据（或者是计算效果），远距离的线程间则独立无关。
这个性子反映在CUDA里，就是Block划分的两重天地：Block内部具有Shared Memory，线程间可以共享数据、通讯和同步，Block外部则完全独立，Block间没有通讯机制，相互执行顺序不影响计算效果。这一划分使得我们既可以利用线程间通讯做一些复杂的应用和算法加快，又可以在Block的粒度上自由调度计算任务，在不同计算能力的硬件平台上自适应的调整任务安排。
[*]第二方面：
多个线程同步执行划一的运算，使得我们可以用单路指令流对多个执行单元举行控制，大幅度减少了控制器的个数和系统的复杂度
[*]第三方面：
把留意力放在“险些划一”这里。最简单的并行计算方案是多路数据上同时举行完全划一的计算，即SIMD（单指令多数据流）。这种方案是非常受限的。究竟上我们可以看出，“完全划一”是不必要的。只要这些计算在大多数时候完全划一，就可以对它们做类似于SIMD的加快，不同点是在计算分叉时候，各个线程不划一的特别环境下，只需要分支内并行，分支间串行执行即可，究竟这些只是很少出现的环境。 这样，把“完全划一”这个限制轻微放松，就可以得到更广阔的应用范围和不输于SIMD的计算性能，即SIMT（单指令流多线程）的一个紧张环节，这是GPU强大处置惩罚能力的缘故原由。
3. CUDA调度框架
随着科研和商业领域对于高性能计算需求的日益增长，GPU作为一种提供了大量并行处置惩罚能力的硬件设备，得到了广泛应用。然而，GPU设备通常价格昂贵，且大概并非全时段都在举行高负载的运算，因此如何进步GPU的利用率，最大限度减小浪费，对于开发者们而言是一大寻衅。
别的在很多场景下，一个用户的应用大概并不需要占用整个GPU，或者同一时段有多个用户或任务需要利用GPU资源。如果每个用户或任务独占一个GPU，大概会导致资源浪费和效率低下。相反，如果可以让多个任务共享同一个GPU，则可以大大进步GPU的利用效率。
由此可见，GPU共享调度的目标主要有以下几点：
[*]进步资源利用率：让多个任务或用户可以共享同一GPU；
[*]降低成本：进步GPU利用率以降低单元任务的计算成本；
[*]进步性能：通过合理的调度策略，减少任务之间的辩论，进步团体运行性能；
[*]提供公平性：在多用户或多任务的环境下，保证每个用户或任务都能公平地获取到GPU资源；
[*]保证任务的隔离性：虽然多个任务共享一个GPU可以进步资源利用率，但也需要保证任务之间的隔离性，防止一个任务影响到其他任务的运行。
共享调度技术主要包含共享和隔离两种技术
1、共享
要在k8s集群中实现GPU共享调度，即多个Pod共享利用同一张显卡，需要集群拥有细粒度分配GPU资源的机制，将整卡的资源拆分成多份，并分配给Pod。要做到这一点，一般是通过扩展资源的方式将GPU注册到节点信息中，调度器根据这些扩展资源信息分配资源，到达共享调度的目的。
2、隔离
现在GPU隔离主要分为三种：
[*]①显存隔离：指将 GPU 的显存资源举行隔离，按部署服务的配置文件中所声明的资源界说分配给对应服务，每个服务所分配的显存资源之间互不影响。
[*]②算力隔离：指将 GPU 的计算能力举行隔离，按比例分配给共享 GPU 的任务上。
[*]③故障隔离：fatal exception发生时会影响其他应用。
3.1 Indirect buffer
IB (Indirect Buffer)间接缓冲特定引擎的命令缓冲区。与直接向队列中写入命令不同，您可以将命令写入一块内存，然后将指向该内存的指针放入队列中。然后，硬件将跟随指针并执行内存中的命令，然后返回到环中的其余命令。
https://i-blog.csdnimg.cn/direct/52a56c55bb0c488aaf883ec61a6f63b9.pnghttps://i-blog.csdnimg.cn/direct/c9ff254f9dab4fb783ca21987692de38.png
GPU Resource Management：GPU channel是GPU与CPU之间的桥接接口，通过CPU向GPU发送GPU指令的唯一通道，GPU channel包含了两类用于存储GPU指令的buffer：
GPU command buffer (也称之为FIFO push buffer)
Ring buffer (也称之为indirect buffer)，从上图中看出，这个buffer是环形结构的，即其容量是固定的，这也是为什么叫Ring buffer的缘故原由吧
当GPU指令被写入到GPU command buffer时，系统还会向Ring buffer中写入与此指令所对应的packet，packet包含了此指令在GPU command buffer中的偏移位置与长度数据。
在执行指令的时候，GPU不是直接从GPU command buffer中读取数据，而是先经过Ring buffer读取出当前待处置惩罚指令的相关信息，再据此读取GPU command（这也是为什么Ring buffer被称之为indirect buffer的缘故原由）。
3.1 基本概念
再聊调度之前，我们照旧先来重点介绍几个相关的概念：channel、tsg、runlist、pbdma。
[*] channel
这是nv driver层的才有的概念，每一个gpu应用程序会创建一个或者多个channel。而channel也是gpu硬件(在gpu context 层面来说)操作的最小单元。
[*] tsg
全称为timeslice group，通常环境下一个tsg含有一个或者多个channel，这些channel 共享这个tsg的timeslice。
[*] runlist
多个tsg或者channel的集合，gpu硬件就是从runlist上选取channel来举行任务执行。
[*] pbdma
全称为pushbuffer dma。push buffer可以简单的理解为一段主机内存，这段内存主要有cpu写然后gpu来读。gpu通过从pushbuffer 里面拿到的数据生成相应的command(也叫methods) 和data(address) 。而上面讲到的channel里面包含有指向pushbuffer的指针。
https://i-blog.csdnimg.cn/blog_migrate/2ed1a860b048f1c59580feecd9b209ca.png
图13
结合图13再给大家理一下上面几个概念之前的一些关联。起首，runlist里面的每个entry就是一个channel，每个channel里面有Inst Blk Ptr 也即instance块指针，这些指针分别指向生存gpu上下文的内存和push buffer也即上图当中的PB seg。
接着我们先来简单的描述一下gpu应用是如何通过channel来提交任务的，具体流程如下：
Submitting new work to a channel involves the following steps:
1. Write methods to a pushbuffer segment
2. Construct a new GP entry pointing to that pushbuffer segment
3. Update GP_PUT in USERD( User-Driver Accessible RAM) to indicate the
new GP entry is ready
4. Request the doorbell handle from RM, given the channel ID
5. Write the channel's handle to the NOTIFY_CHANNEL_PENDING register
相信大家结合上面的一些讲述应该比力容易看懂上面的提交流程这里就不再赘述了，接下来我们回到调度正题上来。上面说到了应用提交work的相关流程，那这个work提交之后呢？这就涉及到如何将这些任务举行调度和执行了，下面我们先上一个团体调度架构图
https://i-blog.csdnimg.cn/blog_migrate/38424cca82de69154a230244de292143.png
图14 gpu scheduler
gpu的整个调度结构如图14所示，从左到右依次为Application scheduler、stream scheduler、thread block scheduler和warp scheduler。下面我们来逐一对他们举行介绍。
3.2 不同条理调用
NVIDIA GPU 硬件结合 CUDA 编程模子，提供了很多不同的并发机制，以进步 GPU 的利用，用户可以根据自身需求选择不同的技术方案：
https://img-blog.csdnimg.cn/img_convert/397992ee1133192d5d9eefb7ef2763ff.webp?x-oss-process=image/format,png
3.2.1 K8S
2、Time Slicing
英伟达的Time Slicing是一种基于时间片的GPU共享调度策略，这种策略能让多个任务在同一个GPU上举行，而不是每个任务都独占一个GPU。这种策略的核心原理就是将时间分割成一系列的小片断，然后将这些时间片轮流分配给不同的任务。
3.2.2 User scheduler
3、多实例GPU（ MIG ）
迄今为止讨论的机制要么依赖于利用 CUDA 编程模子API（如 CUDA 流）对应用程序的更改，要么依赖于CUDA系统软件（如时间切片或 MPS ）。
利用MIG，基于 NVIDIA 安培体系结构的 GPU ，例如 NVIDIA A100 ，可以为 CUDA 应用程序安全划分多达七个独立的 GPU 实例，为多个应用程序提供专用的GPU资源。这些包罗流式多处置惩罚器（SMs）和GPU引擎，如复制引擎或解码器，为不同的客户端如历程、容器或虚拟机（ VM ）等提供界说的QoS和故障隔离。
当对GPU举行分区时，可以在单个MIG实例中利用之前的CUDA流、CUDA MPS和时间切片机制。
4、vGPU
NVIDIA vGPU 使具有完全输入输出内存管理单元（ IOMMU ）掩护的虚拟机能够同时直接访问单个物理 GPU 。除了安全性之外， NVIDIA vGPU 还存在其他优势，如通过及时虚拟机迁徙举行虚拟机管理，能够运行混合的 VDI 和计算工作负载，以及与很多行业虚拟机监控程序的集成。值得留意的是，利用vGPU需要license，购买license的费用需要思量在技术选型里面。
利用基于PCIE的
3.2.3 Application scheduler
MPS —— 它通过将多个历程的 CUDA Context，归并到一个 CUDA Context 中，省去了 Context Switch 的开销，也在 Context 内部实现了算力隔离。如前所述，MPS 的致命缺陷，是把很多历程的 CUDA Context 归并成一个，从而导致了额外的故障传播。以是尽管它的算力隔离效果极好，但恒久以来工业界利用不多，多租户场景尤其如此。
通常环境下两个不同的gpu应用是不能同时占用gpu的计算单元的，他们只能通过期分复用的方法来利用gpu。具体来讲就是gpu按照FIFO的策略依次从runlist上拿取channel举行执行，每一个channel只能运行肯定的时间，等时间片用完之后就会举行切换来运行其他的channel。但是这种时分复用的调度算法有一个缺陷就是如果App每次提交的任务都比力小就无法占满gpu SM从而导致了gpu 团体利用率比力低。为了解决这个题目，nvidia 又提出了一别的一种调试算法叫Multi-Process Service，我们也叫空分。在MPS的场景下它允许两个不同的应用能够在同一时刻去占用不同的gpu sm，从而来进步gpu的利用率。
https://i-blog.csdnimg.cn/blog_migrate/6ddc20b14e03a0ee9353f27613c42c6b.png
图15 MPS
3.2.4 stream scheduler
当gpu从runlist里面取出channel之后会生成相应的command和数据，而每个stream里面包含了一系列的commands。由于不同的应用的stream是可以设置不同的优先级的，以是stream scheduler主要负责不同应用的stream的调度和抢占。
3.2.5 Thread Block scheduler
它主要负责将thread block assign给gpu的sm，完成thread block跟gpu sm之间的逐一映射。通常能不能将一个 kernel的thread block assign给某个sm主要看SM上的计算能力。举个例子，如果说一个sm支持 2048 threads和32 blocks，那么如果某个kernel有64个threads和64个blocks则scheduler也只能选这个kernel一半的blocks去运行。
3.2.6 warp scheduler
通常环境下一个warp包含了32个thread，warp scheduler的主要作用就是从wrap中获取准备好的待执行的instruction，并把这些instruction分配给sm上的Disaptch Unit。接着Dispatch Unit会把这些指令发送到SM的SIMD core上执行。
总结
gpu上的其他细节另有很多，笔者这篇文章就看成抛砖引玉了。如果大家想更加深入的研究的话可以去看看nv的一些open gpu doc，别的就是官方放出来的一些开源代码。这些都是非常紧张的研究材料，仔细研读之后应该会有一些启发
参考文献
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环球同步评测 - 超能网
————————————————
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
原文链接：https://blog.csdn.net/qq_41554005/article/details/119760698
深入GPU原理：线程和缓存关系【AI芯片】GPU原理01_哔哩哔哩_bilibiliwww.bilibili.com/video/BV1bm4y1m7Ki/?spm_id_from=333.999.0.0&vd_source=bf331b9ca4fb3b040bf9d1e87899c075编辑
AISystem/02Hardware/03GPUBase at main · chenzomi12/AISystem (github.com)github.com/chenzomi12/AISystem/tree/main/02Hardware/03GPUBase
特别分析：本文是对开源项目AISystem 的内容贡献
@ZOMI酱
1.4 Single Instruction Multiple Threads（SIMT）
[*]GPU中的SIMT体系结构相对于CPU的SIMD（单指令多数据，Single Instruction Multiple Data）。中文翻译：单指令多线程。SIMT对于可编程性的利益使得NVIDIA的GPU架构师为这种架构命名，而不是将其描述为 SIMD 。
[*]为了有效地管理和执行多个单线程，流多处置惩罚器（SM）采用了SIMT架构。此架构在第一个unified computing GPU中由NVIDIA公司生产的GPU引入。
[*]GPU利用SIMT执行 32 个并行线程的 Warp ，实现单指令、多线程，这使得每个线程能够访问本身的寄存器，从不同的地址加载和存储，并遵循不同的控制流路径。CUDA编译器和GPU一起工作，以确保Warp的线程组尽大概频繁地被分配到SM中，一起执行类似的指令序列，从而最大限度地进步性能。
[*]每个线程可以包含控制流指令（控制流指令为标量指令）
[*]同组Warp中的这些线程可以执行不同的控制流路径
[*]当一个Warp中的线程分支到不同的执行路径时，产生分支发散（Branch divergence）
https://i-blog.csdnimg.cn/direct/129e5d2204d2494b803b485c00732c5b.png
https://i-blog.csdnimg.cn/direct/9f38ab3ee97f436da665e94b0e4a6707.png
优势
[*]共享控制逻辑可以有更多的空间面基去分配给计算单元
[*]大量的并行操作，不需要举行复杂的控制编程
SIMD VS SIMT
[*]CPU中通过SIMD来处置惩罚矢量数据；纯粹利用SIMD不能并行的执行有条件跳转的函数，很显然条件跳转会根据输入数据不同在不同的线程中有不同表现。
[*]GPU则利用SIMT，无需开发者费力把数据凑成合适的矢量长度，并且SIMT允许每个线程有不同的分支，利用SIMT 才能做到不同分支的并行操作。
免责声明：如果侵犯了您的权益，请联系站长，我们会及时删除侵权内容，谢谢合作！更多信息从访问主页：qidao123.com:ToB企服之家，中国第一个企服评测及商务社交产业平台。
页:
[1]