在使用CUDA Runtime API進行編程時，會在第一次調用CUDA runtime函數時自動（隱式的）創建CUDA Context（Primary Context）。CUDA Context中包含了Process中所有的狀態，例如顯存的映射/分配，Kernel函數定義，以及其他相關狀態信息。每個Process都有自己的CUDA Context並對應1個GPU設備，並且其不能在多個Process之間共享。
對於1個Process內使用multi-stream或者multi-thread方式進行多Kernel並發的情況，因為它們在同一CUDA context內，Hyper-Q特性允許同一時刻有多個kernel同時使用GPU資源且並發地執行。當單個Kernel佔用的GPU資源較少時，這種方式能夠有效的提升GPU的利用率，並且在時間維度上縮短整個程序執行時間（相比於多個Kernel串行執行），提升任務吞吐能力。
對於多個Process進行多Kernel並發的情況，由於它們都有各自的CUDA Context，雖然Hyper-Q特性也支持將多個Context同時放在GPU上執行，但GPU並不能在同一時間內同時執行多個context內的Kernel，它們需要在GPU上進行時間片輪轉調度，同時要做Context switch。程序總的執行時間，與多個kernel串行執行的時間總和相等或更長。
MPS是Multi-Process Service的縮寫。它是CUDA API中的一項特殊功能，支持在具有Hyper-Q功能的GPU（Kepler/Maxwell/Pascal/Volta…）上實現多進程CUDA程序（特別是基於MPI的多任務程序）的並發執行。
也就是說，對於上述multi-process無法真正實現多個kernel並行執行的情況，開啟MPS服務能夠幫助我們解決該問題。MPS能夠將多個Process的CUDA context整合成一個，也就是說不需要再進行GPU上的時間片輪轉以及context switch，這與使用單個Process內multi-stream方式做kernel並發的效率是一樣的。而且，不能忽視的是，由於使用了多進程，CPU端代碼的執行效率也提高了。這對於很多線上服務來說非常簡單高效。
MPS服務開啟的步驟如下：
1) 只可見需要運行MPS服務的GPU，例如0號設備：export CUDA_VISIBLE_DEVICES=0
2) 設置GPU為exclusive process計算模式（需要sudo權限）：sudo nvidia-smi -i 0 -c EXCLUSIVE_PROCESS
註：此步驟並非必須的。不開啟exclusive模式也可以使用MPS。開啟該功能的目的只是確保1塊GPU上只有1個process context。
3) 啟動MPS服務：nvidia-cuda-mps-control -d
4) 之後，在另一個窗口中（client端），使用mpirun執行多進程。
5) 查看mps deamon及服務：ps -ef | grep mps
如果執行過process，則server會被啟動，並且一直留在系統中：
1) 停止MPS服務：echo quit | nvidia-cuda-mps-control
2) 查看deamon及server是否已關閉：ps -ef | grep mps
3) 將compute mode恢復為默認：sudo nvidia-smi -i 0 -c 0
OS：Ubuntu 16.04
CUDA：9.1.85
GPU：TITAN V *
*在不同架構不同型號GPU以下測試取得的結果略有不同。
__global__ void testMaxFlopsKernel(float * pData, int nRepeats, float v1, float v2)
{
int tid = blockIdx.x* blockDim.x+ threadIdx.x;
float s = pData[tid], s2 = 10.0f - s, s3 = 9.0f - s, s4 = 9.0f - s2;
for(int i = 0; i < nRepeats; i++)
{
s=v1-s*v2;
s2=v1-s1*v2;
s3=v1-s2*v2;
s4=v1-s3*v2;
}
pData[tid] = ((s+s2)+(s3+s4));
}
int nRepeats = 1000000000; // repeat computing number per thread
dim3 blockSize(1, 1, 1); // block size
dim3 gridSize(1,1,1); // grid size
testMaxFlopsKernel<<<gridSize, blockSize>>>(fpData_d, nRepeats, 10.0f, 9.9899f);
為了簡單易理解，我們只設置kernel發射1個thread，設置的「nRepeats」較大的目的是讓kernel函數執行時間儘可能長，這樣利用觀察多個process的kernel之間進行overlap。
2個Kernel的執行時間分別為6854ms，6838ms，幾乎是單個Kernel執行時長的2倍。正如前文所述，因為kernel在GPU上是被時間片輪轉調度執行的，相當於2個kernel是在交替串行執行。4個Kernel執行時間分別為13952，13947ms，13934ms，13940ms。結論同上。
2個Kernel執行時間分別為3467ms，3467ms，與單個Kernel執行的時間基本相同。多個Kernel之間真正實現了並行執行。
4個Kernel的執行時間分別為3505ms，3492ms，3492ms，3492ms。結論同上。
分析工具：nvprof + nvvp
1) 啟動nvprof：nvprof --profile-all-processes -o output.%p
打開另一個ssh窗口，執行mpirun啟動多進程執行：回到nvprof窗口，可以看到已經生成了2個process的output文件，此時按「ctrl+c」可以退出nvprof（不退出亦可）：
2) 打開nvvp，在file菜單選擇import，然後選擇nvprof；
Next後，選擇multi process：
Next，Browse選擇output文件：
點擊Finish，即可得到timeline結果。具體參見下面的具體分析結果。
結論：可以看到4個Process發射的4個Kernel之間可以concurrent執行，但並不是真正意義上的overlap。實際上，每個Process發射的Kernel在自己的執行周期內都佔用了整個GPU（雖然只有1個thread），而多個Kernel之間採用時間片輪轉調度的方式在GPU上執行，所以每個Kernel的執行都被「拉長」了，而多個Kernel並發後總的執行時間約等於每個Kernel串行執行時間的總和。
結論：可以看到4個Process發射的4個Kernel之間可以concurrent執行，並且是overlap的。每個Kernel的執行時間基本等於單個Process發射該kernel的執行時間。而多個Kernel並發後（前提是每個Lernel佔用的GPU計算和存儲資源都不大）總的執行時間約等於單個Kernel執行時間，真正實現了並發執行，充分利用GPU資源的同時，減少了多進程總的執行時間，提升了吞吐性能。