如何使用 Nvidia 多进程服务 (MPS) 运行多个非 MPI CUDA 应用程序？

Posted 2023-04-15

【中文标题】如何使用 Nvidia 多进程服务 (MPS) 运行多个非 MPI CUDA 应用程序？

【英文标题】：How do I use Nvidia Multi-process Service (MPS) to run multiple non-MPI CUDA applications?

【发布时间】：2016-04-15 01:35:28

【问题描述】：

我能否在具有 MPS 的 NVIDIA Kepler GPU 上同时运行非 MPI CUDA 应用程序？我想这样做是因为我的应用程序无法充分利用 GPU，所以我希望它们共同运行。有没有代码示例可以做到这一点？

【参考方案1】：

documentation 中包含 MPS 服务的必要说明。您会注意到，这些指令并不真正依赖或调用 MPI，因此它们确实没有任何 MPI 特定的内容。

这是一个演练/示例。

阅读上述链接文档的第 2.3 节，了解各种要求和限制。我建议为此使用 CUDA 7、7.5 或更高版本。与以前版本的 CUDA MPS 存在一些配置差异，我不会在这里介绍。此外，我将演示仅使用单个服务器/单个 GPU。我用于测试的机器是使用 K40c (cc3.5/Kepler) GPU 和 CUDA 7.0 的 CentOS 6.2 节点。节点中还有其他 GPU。在我的例子中，CUDA 枚举顺序将我的 K40c 放置在设备 0 上，但 nvidia-smi 枚举顺序恰好将它作为 id 2 放置在顺序中。在具有多个 GPU 的系统中，所有这些细节都很重要，会影响下面给出的脚本。

我将创建几个帮助 bash 脚本以及一个测试应用程序。对于测试应用程序，我们想要一些具有明显可以与来自应用程序的其他实例的内核同时运行的内核的东西，并且我们还想要一些能够在这些内核（来自单独的应用程序/进程）时变得明显的东西是否同时运行。为了满足演示目的的这些需求，让我们有一个应用程序，它的内核只在单个 SM 上的单个线程中运行，并且在退出并打印之前简单地等待一段时间（我们将使用约 5 秒）信息。这是一个执行此操作的测试应用程序：

$ 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
$

我们将使用 bash 脚本来启动 MPS 服务器：

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

还有一个 bash 脚本“同时”启动我们的测试应用的 2 个副本：

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

我们也可以有一个 bash 脚本来关闭服务器，尽管本演练不需要它：

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

现在，当我们使用上面的 mps_run 脚本启动我们的测试应用程序，但没有实际启用 MPS 服务器时，我们会得到预期的行为，即应用程序的一个实例需要大约 5 秒，而另一个实例需要大约两倍（约 10 秒），因为它不会与来自另一个进程的应用程序同时运行，它会在另一个应用程序/内核运行时等待 5 秒，然后花 5 秒运行自己的内核，总共约 10 秒：

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

另一方面，如果我们先启动 MPS 服务器，然后重复测试：

$ su
Password:
# ./start_as_root.bash
Set compute mode to EXCLUSIVE_PROCESS for GPU 0000:82:00.0.
All done.
# exit
exit
$ ./mps_run
kernel duration: 6.167079s
kernel duration: 6.263062s
$

我们看到两个应用程序需要相同的时间来运行，因为内核同时运行，这是由于 MPS。

欢迎您进行您认为合适的实验。如果此序列似乎对您正常工作，但运行您自己的应用程序似乎没有给出预期的结果，一个可能的原因可能是您的应用程序/内核无法与应用程序/内核的其他实例同时运行，因为与内核的构造无关，与 MPS 无关。您可能想验证the requirements for concurrent kernels，和/或研究concurrentKernels sample app。

这里的大部分信息都是从 here 完成的测试/工作中回收的，尽管这里使用单独的应用程序的演示与那里的 MPI 案例不同。

更新：当从多个进程运行内核时，非 MPS 情况下的调度程序行为似乎随着 Pascal 和更新的 GPU 而改变。上述测试结果对于测试的 GPU（例如 Kepler）仍然是正确的，但是在 Pascal 或更新的 GPU 上运行上述测试用例时，在非 MPS 情况下会观察到不同的结果。调度程序在the latest MPS doc 中被描述为“时间片”调度程序，似乎正在发生的事情是，调度程序不是等待来自一个进程的内核完成，而是根据一些未发布的规则，选择预先清空正在运行的内核，以便它可以从另一个进程切换到另一个内核。这仍然不意味着来自不同进程的内核在 CUDA 文档中该词的传统用法中“同时”运行，但上述代码被时间片调度程序（在 Pascal 和更新版本上）“欺骗”，因为它取决于关于使用 SM 时钟设置内核持续时间。时间分片调度程序加上 SM 时钟的这种用法的组合使这个测试用例看起来“并发”运行。但是，如 MPS 文档中所述，当 A 和 B 在非 MPS 情况下源自不同的进程时，来自内核 A 的代码与来自内核 B 的代码在相同的时钟周期内执行。

使用上述通用方法演示这一点的另一种方法可能是使用由多个循环设置的内核持续时间，而不是通过读取 SM 时钟设置的内核持续时间，如 here 所述。在这种情况下必须小心，以避免编译器“优化”循环。

【讨论】：

是否可以通过CUDA从头实现MPS？

我不这么认为。

@RobertCrovella 你能分享你对这个问题的看法吗？为什么你认为这是不可能的？似乎 MPS 使用了一些未记录的驱动程序 API 或类似的东西？

MPS 接受来自不同进程的工作（例如 CUDA 内核启动），并在设备上运行它们，就好像它们来自单个进程一样。就好像它们在单个上下文中运行一样。我不知道如何使用我熟悉的当前公开的 API 来做到这一点。除此之外，我没有任何进一步的想法，并且不太可能回答有关该主题的进一步问题。