市面上有很多GPU厂家，他们产品的软硬件架构各不相同，但是核心往往差不多，整明白了一个基本上就可以触类旁通了。针对当前gpu底层的一些架构以及硬件层一些调度策略的话估计大部分人就很难说的上熟悉了，这个不是大家的错，主要是因为Nv gpu的整个生态都是闭源的，所以大家了解起来就会有一些障碍。下面的行文将基于以下三个层面进行阐述：CUDA编程模型、GPU 底层硬件架构与硬件层的调度策略、CUDA调度框架。
通过行业CUDA标杆做基本的调度管理与分析，在实际的CUDA代码执行过程中需要CPU和GPU的协同工作，在CPU上运行的称为Host程序，在GPU上运行的称为Device程序。比方说对于一个CUDA程序的可以分为两个部分（两者拥有各自的存储器）。
CUDA 最基本的执行单位是线程（Thread），图中每条曲线可视为单个线程，大的网格（Grid）被切分成小的网格，其中包含了很多相同线程数量的块（Block），每个块中的线程独立执行，可以通过本地数据共享实现数据交换同步。因此对于 CUDA 来讲，就可以将问题划分为独立线程块，并行解决的子问题，子问题划分为可以由块内线程并行协作解决。
CUDA 引入主机端（host）和设备（device）概念，CUDA 程序中既包含主机（host）程序也包含设备（device）程序，host 和 device 之间可以进行通信，以此来实现数据拷贝，主机负责管理数据和控制程序流程，设备负责执行并行计算任务。在 CUDA 编程中，Kernel 是在 GPU 上并行执行的函数，开发人员编写 Kernel 来描述并行计算任务，然后在主机上调用 Kernel 来在 GPU 上执行计算。
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[i] = x[i] + y[i];
}
int main(void)
{
int N = 1<<25; // 30M elements
float *x = new float[N];
float *y = new float[N];
// initialize x and y arrays on the host
for (int i = 0; i < N; i++) {
x[i] = 1.0f;
y[i] = 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[i] = x[i] + y[i];
}
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[i] = 1.0f;
y[i] = 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 通信。
因此 CUDA 和英伟达硬件架构有以下对应关系，从软件侧看到的是线程的执行，对应于硬件上的 CUDA Core，每个线程对应于 CUDA Core，软件方面线程数量是超配的，硬件上 CUDA Core 是固定数量的。Block 线程块只在一个 SM 上通过 Warp 进行调度，一旦在 SM 上调用了 Block 线程块，就会一直保留到执行完 kernel，SM 可以同时保存多个 Block 线程块，多个 SM 组成的 TPC 和 GPC 硬件实现了 GPU 并行计算。
1.2 CUDA多维度编程
示例：一维数组的求和计算
代码中注释的一、二处究竟该怎么来写？
------------------------------------------------------------
线程参数设置 情况1：一维grid，一维block （线程分配）
grid(1,1,1): block数量=1*1*1
block(length,1,1): thread数量=length*1*1
总thread数量 = (1*1*1)*(length*1*1)
-------------------------------------------------------------------------------------------
线程参数设置 情况二2：一维grid，二维block （线程分配）
grid(1,1,1): block数量=1*1*1
block(8,2,1): thread数量=8*2*1
总thread数量 = 16
我们一定要有并行思想，这里有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分配）
---------------------------------------------------
线程参数设置情况4: block和thread都分配
------------------------------------------------------------------
线程参数设置 情况5：二维grid，二维thread
示例：倒推其线程参数设置
它的线程参数设置是怎样的？线程索引怎么计算？
参数设置为：
总Thread数量： 8*4*1*8*2*1 = 512
一维数组的线程索引计算方法：
二维数组的线程索引计算方法：
根据CUDA算力不同thread，block，gird在不同维度的大小是有限制的:
Cuda Wrap的限制：
1.3 Stream
中文翻译为"流"，它主要是通过提升kernel函数的并发性来提升整个计算的运行效率。下面我们来看一下在cuda编程模型当中具体是如何使用stream的。
cudaStream_t?stream[nStreams];??
for?(int?i?=?0;?i?<?nStreams;?i?++)??
{??
????checkCuda(cudaStreamCreate(&stream[i]));??
}??
for?(int?i?=?0;?i?<?nStreams;?i?++)??
{??
????checkCuda(cudaStreamDestroy(stream[i]));??
}
上面所展示的是stream的创建和销毁，接下来我们来看一下如何使用stream
for?(int?i?=?0;?i?<?nStreams;?i?++)???
{??
????int?offset?=?i?*?streamSize;??
????checkCuda(cudaMemcpyAsync(&d_a[offset],?&a[offset],?streamBytes,?cudaMemcpyHostToDevice,?stream[i]));??
????kernel_function<<<streamSize/blockSize,?blockSize,?0,?stream[i]>>>(d_a,?offset);??
????checkCuda(cudaMemcpyAsync(&a[offset],?&d_a[offset],?streamBytes,?cudaMemcpyDeviceToHost,?stream[i]));??
}
stream具体用法如上面sample所示，如果你不显示的申请stream的话系统也会有一个default的stream0。大家可以从下面的这张图比较直观地看到两者在执行效率上的区别：
图3 cuda stream 串行和并行执行
1.4Graph
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在不同维度的大小产生影响。
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 软件抽象和硬件结构的一一对应关系
硬件结构可以参考之前的一篇文章
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计算模式的对比
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 = 7682432=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)间接缓冲特定引擎的命令缓冲区。与直接向队列中写入命令不同，您可以将命令写入一块内存，然后将指向该内存的指针放入队列中。然后，硬件将跟随指针并执行内存中的命令，然后返回到环中的其余命令。
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的指针。
图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提交之后呢？这就涉及到如何将这些任务进行调度和执行了，下面我们先上一个整体调度架构图
图14 gpu scheduler
gpu的整个调度结构如图14所示，从左到右依次为Application scheduler、stream scheduler、thread block scheduler和warp scheduler。下面我们来一一对他们进行介绍。
3.2 不同层次调用
NVIDIA GPU 硬件结合 CUDA 编程模型，提供了许多不同的并发机制，以提高 GPU 的利用，用户可以根据自身需求选择不同的技术方案：
3.2.1 K8S
2、Time Slicing
英伟达的Time Slicing是一种基于时间片的GPU共享调度策略，这种策略能让多个任务在同一个GPU上进行，而不是每个任务都独占一个GPU。这种策略的核心原理就是将时间分割成一系列的小片段，然后将这些时间片轮流分配给不同的任务。
3.2.2User 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.3Application 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的使用率。
图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 warpscheduler
通常情况下一个warp包含了32个thread，warpscheduler的主要作用就是从wrap中获取准备好的待执行的instruction，并把这些instruction分配给sm上的DisaptchUnit。接着DispatchUnit会把这些指令发送到SM的SIMDcore上执行。
总结
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
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）
优势
- 共享控制逻辑可以有更多的空间面基去分配给计算单元
- 大量的并行操作，不需要进行复杂的控制编程
SIMD VS SIMT
- CPU中通过SIMD来处理矢量数据；纯粹使用SIMD不能并行的执行有条件跳转的函数，很显然条件跳转会根据输入数据不同在不同的线程中有不同表现。
- GPU则使用SIMT，无需开发者费力把数据凑成合适的矢量长度，并且SIMT允许每个线程有不同的分支，利用SIMT 才能做到不同分支的并行操作。