本人雖然研二開始接觸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的三種模式：

如上圖：

預設計算模式：多個程序在啟動時，驅動可以開啟多個上下文物件（context）分別繫結,比如一個程序繫結一個上下文物件，那麼這時候就涉及到單卡多程序內部是如何排程的，
獨佔程序計算模式：即驅動只開啟一個上下文物件，但是通常cpu測 程序之間是完全資源隔離的，那麼所謂開啟一個上下文，也估計只能對應一個程序（MPS除外，MPS就是工作在此模式下）；
禁止計算模式：即在裝置上不建立上下文（不明白這個模式的使用場景）；

那麼如上面介紹的三種模式，最常接觸的就是預設模式，這時候不論是使用者開啟一個tensorflow-gpu程式，還是看nvidia-smi顯示比較空閒去開啟多個gpu的程式都會有個疑問：

1:我開多程序能更好的利用單卡麼？
2:以及為什麼nvidia又有個東西叫MPS？

3 上下文切換的時間粒度

接著看3.2.5.2章節下面截圖的最後一句

官方也說了來自不同的上下文的kernel是不能同時執行的。那麼針對這個問題就有疑問了，是整塊卡不能同時執行，還是針對一個SM不能同時執行，還是針對SM中一個core不能同時執行。因為有傳統作業系統知識的同學就知道了，cpu支援多程序（多工）是通過時間片輪詢的方式去搶佔正在執行的任務的。但是cpu一個core我們理解就是一個單元啊，預設不可拆了啊，但是cuda可不是啊，一塊卡內部一堆SM，然後每個SM內部一堆core，我們編寫cuda程式碼時候是可以在一個執行緒裡面操作的，然後外部寫個<<<grid，block>>>去申請資源的，那假設我寫2個程序，內部分別只申請不到50%的資源，那到底一塊卡能不能同一時刻同時執行2個程序呢？帶著這個問題又找到了另一個地方

我們看到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發現一個有趣的現象


就是都是顯示利用率為100%，至於為什麼mem那一列為0.一切都在stackoverflow上找到了原因
nvidia-smi-volatile-gpu-utilization-explanation

即因為nvidia-smi是通過取樣，然後所謂gpu的使用率是從時間線角度，當前程式SM是否在使用，那mem為什麼為0，估計是獲取資料太小，取樣時刻沒法監測到。但是這裡的確說明一個問題，看nvidia-smi來衡量你的GPU當前是否繁忙，SM是否全都用上了，內部全域性視訊記憶體和L2到L1以及暫存器的IO傳輸使用率啥的一概不準。而再加上是基於整卡進行上下文切換，那更隱藏了很多資源使用率的有效資訊。