市面上有很多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
#include
#include
// 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__ 符号声明，在调用时需要用 <<
#include
#include
// 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