MPS 多进程服务（Multi-Process Scheduling）是 CUDA 应用程序编程接口（API）的替代二进制兼容实现。从 Kepler 的 GP10 架构开始，NVIDIA 就引入了 MPS（基于软件的多进程服务），这种技术在当时实际上是称为 HyperQ ，允许多个 stream 或者 CPU 的进程同时向 GPU 发射 Kernel 函数，结合为一个单一应用程序的上下文在 GPU 上运行，从而实现更好的 GPU 利用率。在单个进程的任务处理，对 GPU 利用率不高的情况下是非常有用的。实际上，在 Pascal 架构出现之后的 MPS 可以认为是 HyperQ 的一种实现方式。

MPS 作用

多容器共享 GPU 会引发资源竞争和浪费的问题，不仅 GPU 利用率不高，也无法保证 QoS。当使用 NVIDIA MPS 时，MPS Sever 会通过一个 CUDA Context 管理 GPU 硬件资源，多个 MPS Clients 会将它们的任务通过 MPS Server 传入 GPU，从而越过了硬件时间分片调度的限制，使得它们的 CUDA kernels 实现真正意义上的并行。特别地，Volta MPS 可以兼容 docker 容器，并且支持执行资源配置（即每个 Client context 只能获取一定比例的 threads），提高了多容器共享 GPU 的 QoS。

  • 增加 GPU 的利用率
  • 减少多个 CUDA 进程在 GPU 上的上下文空间。该空间主要是用于存储和调度资源；
  • 减少 GPU 的上下文的切换。

MPS 实例

假设在 CPU 端有 A、B、C 三个进程，每个进程都要发射 CUDA Kernel 的任务到 GPU 上去，并且假设它们每一个独立的任务对 GPU 利用率都不高。

在不使用 MPS 服务的情况下，A、B、C 三个进程实际上也可以同时把 CUDA 任务发射到 GPU 上去，但是默认采用时间片轮转调度的方式。首先第一个时间片，A 任务被执行，接着第二个时间片，执行 B 任务，第三个时间片， C 任务将被执行。时间片是依次进行轮转调度的，分别执行 A、B、C 中的任务。

可以直观地看到， 在 GPU 中，每一个时刻只有一个任务在执行。这种情况下，CPU 中的 process（进程）发射的 CUDA 任务对 GPU 的利用率很低。

MPS 与 Hyper-Q 区别

MPS 与 Hyper-Q 的区别是：来自单个进程的 CUDA 任务与来自多个进程的 CUDA 任务之间存在差异。 Hyper-Q 消除了来自单个进程的请求并发性的一些人为障碍。但是由于与 Hyper-Q 无关的 CUDA 行为，来自多个进程的请求仍然序列化。 MPS 作为一个“漏斗”，用于从多个进程/级别收集 CUDA 任务，并将它们发布到 GPU，就好像它们来自单个进程一样，以便 Hyper-Q 可以生效。

Pascal 架构和 Volta 架构

我们常用的 GTX 1080 ti 和 GP100 是 Pascal 架构的，NVIDIA 基于深度学习的任务特点推出了 Volta 架构的 Tesla V100 数据中心级显卡。我们简单低看一下两种架构的 GPU 区别。

上图是 Pascal SM 架构图，可以看到一个 GP100 SM 分成两个处理块，每块有 32768 个 32 位寄存器 + 32 个单精度 CUDA 核心 + 16 个双精度 CUDA 核心 + 8 个特殊功能单元（SFU） + 8 个存取单元 + 一个指令缓冲区 + 一个 warp 调度器 + 两个分发单元。LD/ST 加载存储单元，SFU 为特殊功能单元，用来执行超越函数指令，如正弦函数。

上图是 Volta SM 架构图，每个 SM 有 64 个 FP32 核、64 个 INT32 核、32 个 FP64 核与 8 个全新的 Tensor Core。新的 Tensor Core 是 Volta GV100 最重要的特征，有助于提高训练神经网络所需的性能。Tesla V100 的 Tensor Core 能够为训练、推理应用的提供 120 Tensor TFLOPS。相比于在 P100 FP 32 上，在 Tesla V100 上进行深度学习训练有 12 倍的峰值 TFLOPS 提升。而在深度学习推理能力上，相比于 P100 FP16 运算，有了 6 倍的提升。与前一代 Pascal GP100 GPU 类似，GV100 GPU 由多个图形处理集群（Graphics Processing Cluster，GPC）、纹理处理集群（Texture Processing Cluster，TPC）、流式多处理器（Streaming Multiprocessor，SM）以及内存控制器组成。一个完整的 GV100 GPU 由 6 个 GPC、84 个 Volta SM、42 个 TPC（每个 TPC 包含了 2 个 SM）和 8 个 512 位的内存控制器（共 4096 位）。

不同架构上的 MPS 实现

上图是基于 Pascal 架构的 MPS 服务对任务的处理情况。可以看到 A、B、C 三个进程分别提交各自的任务到 MPS 的服务端，并在服务端整合为一个统一的上下文，并将三个任务同时发射到 GPU 中执行，这就有效地提升了 GPU 的利用率。在 Pascal 架构下，MPS 是最多可以支持 16 个进程或者说 16 个用户同时提交任务。

上图是 Volta 架构 MPS 的执行情况，Volta 架构对 MPS 的实现做了改进，主要是基于硬件加速的方式来实现。此时不同的进程是可以直接穿过 MPS 服务器，提交任务到 GPU 的硬件，并且每个进程客户端有隔离的地址空间，这样可以进一步减少 Launch（发射进程）时带来的延迟，也可以通过限制执行资源配置来提升服务质量。这里所说提升服务质量是指怎么样平衡多个 process（进程）发射任务对计算和存储资源的占用情况。比如我们现在可以去设定每一个 process 上下文，最多可以使用多少个资源。Volta 下的 MPS 服务最多可以允许同时 48 个 Client（客户端）。

MPS 基准测试

官方给出的一个 benchmark（基准）测试。我们知道，对于单个任务占用 GPU 资源比较少的情形，MPS 服务是非常有用的，比如在深度学习中做 Inference（推理）应用。相比 Training（训练），Inference 对于计算和存储资源的要求比较小，这个时候会出现我们之前看到的情况，单一的 Kernel 任务是没法有效利用 GPU 的。从上面的 benchmark 可以看到图中最左侧灰色的柱状图，在不使用 MPS 的情况下，Inference 的吞吐性能很小；而中间绿色的柱状图，使用 MPS 允许多个 Client 同时发射计算任务到 GPU，此时 GPU 吞吐性能直接提升了七倍；最后一个柱状图表示，如果我们使用 MPS，并结合 Batching 操作，吞吐性能还能继续再提升 60%左右。由此可见，对于像 Deep Learning 的 Inference 这样的应用，MPS 技术是可以有效地帮助我们优化 GPU 利用率以及程序的吞吐性能。

MPS 的使用

MPS 组成

MPS 主要包括控制守护进程（MPS Control Daemo）、客户端运行时（Client Runtime）和服务进程（Server Process）。CUDA contexts 可以通过 MPS Server 提交作业，这样就越过了硬件的时间切片调度的限制，实现多个进程在同一个 GPU 上并行执行。如果没有 MPS，多个进程只是在同一个 GPU 上并发执行，即每个进程在一个时间分片里独享 GPU。

默认情况下，GPU 是没有开启 MPS 的，每个 CUDA 程序会创建自己的 CUDA Context 来管理 GPU 资源，并以时间分片的方式共享 GPU。开启 MPS 后，在需要的时候，MPS control daemon 会启动一个 MPS Server，监听任务请求。

MPS 执行过程

当 CUDA 首次在程序中初始化时，CUDA 驱动程序将尝试连接到 MPS 控制守护进程。 如果连接尝试失败，程序继续运行，正常情况下没有 MPS。 但是，如果连接尝试成功，则 MPS 控制守护程序将继续执行，以确保在与客户端连接之前启动的 MPS 服务器与连接客户端的用户标识相同。 MPS 客户端然后继续连接到服务器。 MPS 客户端，MPS 控制守护程序和 MPS 服务器之间的所有通信都使用命名 pipe 道完成。默认情况下，命名 pipe 道被放置在/tmp/nvidia-mps/下

开启与关闭 MPS

启动 mps-control

1
2
3
4

export CUDA_VISIBLE_DEVICES=0  
export CUDA_MPS_PIPE_DIRECTORY=/tmp/nvidia-mps  
export CUDA_MPS_LOG_DIRECTORY=/tmp/nvidia-log   
nvidia-cuda-mps-control -d

关闭 mps-control

1

echo quit | nvidia-cuda-mps-control

注意以上两个命令需要管理员权限。

Volta MPS 资源配置

执行资源配置的方法如下：

1
2
3

nvidia-cuda-mps-control

set_default_active_thread_percentage 10

命令为每个 MPS Client 限制 10%的 threads。不是为每个 Client 预留专用资源，而是限制它们可以最多使用多少 threads。默认情况下，每个 Client 可以获取所有 threads（即 100%）。

公平性

关闭 MPS，多任务通过时间分片的调度方式共享 GPU；开启 MPS，多任务共享 Server 的 CUDA Context。无论哪种情况，在所有任务所占显存总容量不超出 GPU 容量时，每个任务都能公平地获得 GPU 的 threads。

MPS 程序示例

示例来自stackoverflow

编写主程序

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56

$ cat t1034.cu  
#include <stdio.h>  
#include <stdlib.h>  
  
#define MAX_DELAY 30  
  
#define cudaCheckErrors(msg) \  
  do { \  
    cudaError_t __err = cudaGetLastError(); \  
    if (__err != cudaSuccess) { \  
        fprintf(stderr, "Fatal error: %s (%s at %s:%d)\n", \  
            msg, cudaGetErrorString(__err), \  
            __FILE__, __LINE__); \  
        fprintf(stderr, "*** FAILED - ABORTING\n"); \  
        exit(1); \  
    } \  
  } while (0)  
  
  
#include <time.h>  
#include <sys/time.h>  
#define USECPSEC 1000000ULL  
  
unsigned long long dtime_usec(unsigned long long start){  
  
  timeval tv;  
  gettimeofday(&tv, 0);  
  return ((tv.tv_sec*USECPSEC)+tv.tv_usec)-start;  
}  
  
#define APPRX_CLKS_PER_SEC 1000000000ULL  
__global__ void delay_kernel(unsigned seconds){  
  
  unsigned long long dt = clock64();  
  while (clock64() < (dt + (seconds*APPRX_CLKS_PER_SEC)));  
}  
  
int main(int argc, char *argv[]){  
  
  unsigned delay_t = 5; // seconds, approximately  
  unsigned delay_t_r;  
  if (argc > 1) delay_t_r = atoi(argv[1]);  
  if ((delay_t_r > 0) && (delay_t_r < MAX_DELAY)) delay_t = delay_t_r;  
  unsigned long long difft = dtime_usec(0);  
  delay_kernel<<<1,1>>>(delay_t);  
  cudaDeviceSynchronize();  
  cudaCheckErrors("kernel fail");  
  difft = dtime_usec(difft);  
  printf("kernel duration: %fs\n", difft/(float)USECPSEC);  
  return 0;  
}  
  
$ nvcc -arch=sm_35 -o t1034 t1034.cu  
$ ./t1034  
kernel duration: 6.528574s  
$

编写开启 MPS 脚本

1
2
3
4
5
6

$ cat start_as_root.bash  
#!/bin/bash  
# the following must be performed with root privilege  
export CUDA_VISIBLE_DEVICES="0"  
nvidia-smi -i 0 -c EXCLUSIVE_PROCESS  
nvidia-cuda-mps-control -d

编写执行程序

1
2
3
4
5

$ cat mps_run  
#!/bin/bash  
./t1034 &  
./t1034  
$

编写关闭 MPS 脚本

1
2
3
4
5

$ cat stop_as_root.bash  
#!/bin/bash  
echo quit | nvidia-cuda-mps-control  
nvidia-smi -i 2 -c DEFAULT  
$

运行 MPS 脚本

1
2
3
4

$ ./mps_run  
kernel duration: 6.409399s  
kernel duration: 12.078304s  
$

运行主程序

1
2
3
4

$ ./start_as_root.bash  
$ ./mps_run  
kernel duration: 6.167079s  
kernel duration: 6.263062s

参考资料