本人雖然研二開始接觸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傳輸，一切都是在寄存器，SM內部就執行完了。所以耗時很短，從nvidia-smi可以看出，都不到100%，就一下子結束了，為了讓nvidia-smi抓取到，直接讓他運行100次。
如果將上面注釋去掉，讓他有訪問全局顯存的操作，這時候可以通過nvidia-smi發現一個有趣的現象
就是都是顯示利用率為100%（而且跑了好久好久都沒停，等了一分多鍾還沒結束，手動停止了）.一切都在stackoverflow上找到了原因
nvidia-smi-volatile-gpu-utilization-explanation
即因為nvidia-smi是通過采樣，然后所謂gpu的使用率是從時間線角度，當前程序SM是否在使用，那mem為什么為0？估計是獲取數據太小，采樣時刻沒法監測到。但是這里的確說明一個問題，看nvidia-smi來衡量你的GPU當前是否繁忙，SM是否全都用上了，內部全局顯存和L2到L1以及寄存器的IO傳輸使用率啥的一概不准。而再加上是基於整卡進行上下文切換，那更隱藏了很多資源使用率的有效信息。