nvidia[单卡内部的调度原理]

本人虽然研二开始接触cuda,但是终究未从事cuda开发,故而皆为零零散散的知识,虽然看了好几本cuda编程的书籍以及官网的文档(肯定没看全啊,我也不是专门从事cuda开发),市面上几乎都是关于如何cuda编程的书籍,而这些书籍中也看过不少《CUDA C编程权威指南》,《CUDA专家手册》,《CUDA并行程序设计 GPU编程指南》,《高性能CUDA应用设计与开发 方法与最佳实践 》等等,以及官网《CUDA_C_Programming_Guide》此类文档,还有论文《GPU Scheduling on the NVIDIA TX2: Hidden Details Revealed》,但是都没有完全的系统的去介绍底层的调度原理(当然本博文也没法做到完全系统的挖掘)。


1 引言

先介绍几个概念:
上下文(context):gpu也学着cpu的设计模式,创建了所谓上下文的概念,在cpu中:

CPU寄存器,是CPU内置的容量小、但速度极快的内存。程序计数器,则是用来存储CPU正在执行的指令的位置,或者即将执行的下一条指令的位置。他们都是CPU在运行任何任务前,必须依赖的环境,因此也被叫做CPU上下文。

那么(猜测,待验证)gpu的上下文也差不多就是内置的寄存器状态,L1缓存,以及指令计数器啥的。
进程:这里指host侧的进程
线程:这里指device侧的线程
任务:这里指linux系统下的线程

2 nvidia的gpu的三种模式

首先《CUDA_C_Programming_Guide》的3.5章节,介绍了gpu的三种模式:
nvidia[单卡内部的调度原理]
如上图:

默认计算模式:多个进程在启动时,驱动可以开启多个上下文对象(context)分别绑定,比如一个进程绑定一个上下文对象,那么这时候就涉及到单卡多进程内部是如何调度的,
独占进程计算模式:即驱动只开启一个上下文对象,但是通常cpu测 进程之间是完全资源隔离的,那么所谓开启一个上下文,也估计只能对应一个进程(MPS除外,MPS就是工作在此模式下);
禁止计算模式:即在设备上不创建上下文(不明白这个模式的使用场景);

那么如上面介绍的三种模式,最常接触的就是默认模式,这时候不论是用户开启一个tensorflow-gpu程序,还是看nvidia-smi显示比较空闲去开启多个gpu的程序都会有个疑问:

1:我开多进程能更好的利用单卡么?
2:以及为什么nvidia又有个东西叫MPS?

3 上下文切换的时间粒度

接着看3.2.5.2章节下面截图的最后一句
nvidia[单卡内部的调度原理]
官方也说了来自不同的上下文的kernel是不能同时执行的。那么针对这个问题就有疑问了,是整块卡不能同时执行,还是针对一个SM不能同时执行,还是针对SM中一个core不能同时执行。因为有传统操作系统知识的同学就知道了,cpu支持多进程(多任务)是通过时间片轮询的方式去抢占正在运行的任务的。但是cpu一个core我们理解就是一个单元啊,默认不可拆了啊,但是cuda可不是啊,一块卡内部一堆SM,然后每个SM内部一堆core,我们编写cuda代码时候是可以在一个线程里面操作的,然后外部写个<<<grid,block>>>去申请资源的,那假设我写2个进程,内部分别只申请不到50%的资源,那到底一块卡能不能同一时刻同时执行2个进程呢?带着这个问题又找到了另一个地方
nvidia[单卡内部的调度原理]
我们看到3.6章节紧邻的最上部分,说在之前开普勒和麦克斯韦等架构上,抢占是线程块级别的,在后续帕斯卡等架构上,是能指令级别的。开始想那不就是指令级别可以互相抢占么,可是转念一想,这说的是时间上的抢占粒度,和空间上是卡级别?还是SM级别?还是core级别(当然这个粒度是不可能的,毕竟那么多书籍文档都说明是按照warp去调度的,最小粒度也就是一半warp)?没关系啊?

4 上下文切换的空间粒度

这里随便写个代码:

// nvcc test.cu -std=c++11 -o test
#include<iostream>
#include<chrono>
#include<stdio.h>

using namespace std;
using namespace chrono;

__global__ void kernel(int *a){

  printf("grid:%d,block:%d,thread:%d\n",gridDim.x,blockIdx.x,threadIdx.x);
  for(int i0=0;i0<1000;i0++)
  for(int i=0;i<200000;i++){
    for(int i=0;i<100;i++)
      //a[i]=a[i]%3;
      int b=i%3;
  }
}

int main(){
  auto st = high_resolution_clock::now();
  int *a;
  cudaMalloc((void**)&a,sizeof(int)*100000000);
  kernel<<<1,1>>>(a);
  cudaDeviceSynchronize();
  auto ed = high_resolution_clock::now();
  cout<<"take time: "<<duration_cast<milliseconds>(ed-st).count()<<" ms"<<endl;
  return 0;
}

所以想法就是,创建一个1块1线程的程序,用它去跑(比如耗时100ms),然后同时运行好几个(比如10个),如果是卡级别的,那么几乎是大于10*100ms的时间(上下文切换的开销),而如果是sm级别的,那总的时间差不多稍大于等于100ms。
所以分别执行如下2个shell命令:

执行10次 test这个程序
for i in `seq 1 10`; do echo $i; done|xargs -n1 ./test

grid:1,block:0,thread:0
take time: 132 ms
grid:1,block:0,thread:0
take time: 116 ms
grid:1,block:0,thread:0
take time: 115 ms
grid:1,block:0,thread:0
take time: 114 ms
grid:1,block:0,thread:0
take time: 118 ms
grid:1,block:0,thread:0
take time: 117 ms
grid:1,block:0,thread:0
take time: 115 ms
grid:1,block:0,thread:0
take time: 115 ms
grid:1,block:0,thread:0
take time: 115 ms
grid:1,block:0,thread:0
take time: 115 ms

同时开启10个进程去执行test
for i in `seq 1 10`; do echo $i; done|xargs -n1 -P10 ./test

grid:1,block:0,thread:0
take time: 782 ms
grid:1,block:0,thread:0
take time: 790 ms
grid:1,block:0,thread:0
take time: 796 ms
grid:1,block:0,thread:0
take time: 802 ms
grid:1,block:0,thread:0
take time: 801 ms
grid:1,block:0,thread:0
grid:1,block:0,thread:0
take time: 811 ms
take time: 812 ms
grid:1,block:0,thread:0
take time: 968 ms
grid:1,block:0,thread:0
take time: 983 ms
grid:1,block:0,thread:0
take time: 983 ms

可以看出,几乎是10倍的时间,那为什么不是完全的大于等于10*100ms,就是因为nvcc和内部gcc自带一堆优化(真实原理只是猜测),
即使for循环改成:

  for(int i1=0;i1<10000;i1++)
  for(int i0=0;i0<10000;i0++)
  for(int i=0;i<10000;i++){
    for(int i=0;i<100;i++)
      //a[i]=a[i]%3;
      int b=i%3;
  }

grid:1,block:0,thread:0
take time: 129 ms
也和没加一样,着实佩服。
从这里可以得出结论,所谓上下文切换,是基于整个卡而言的,即一块卡同一时刻只能运行一个上下文的指令。

5 nvidia-smi的gpu利用率解读

当然上述代码并未涉及到IO传输,一切都是在寄存器,core内部就执行完了。所以耗时很短,如果将上面注释去掉,让他有访问全局显存的操作,这时候可以通过nvidia-smi发现一个有趣的现象
nvidia[单卡内部的调度原理]
nvidia[单卡内部的调度原理]
就是都是显示利用率为100%,至于为什么mem那一列为0.一切都在*上找到了原因
nvidia-smi-volatile-gpu-utilization-explanation
nvidia[单卡内部的调度原理]
即因为nvidia-smi是通过采样,然后所谓gpu的使用率是从时间线角度,当前程序SM是否在使用,那mem为什么为0,估计是获取数据太小,采样时刻没法监测到。但是这里的确说明一个问题,看nvidia-smi来衡量你的GPU当前是否繁忙,SM是否全都用上了,内部全局显存和L2到L1以及寄存器的IO传输使用率啥的一概不准。而再加上是基于整卡进行上下文切换,那更隐藏了很多资源使用率的有效信息。

上一篇:requests模块中通用的请求方法,即requests.request


下一篇:Python+requests库发送接口入参为xml格式的接口请求【多测师】