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

Question

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

score 37 · Accepted Answer

必要的说明包含在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 无关。您可能想要验证并发内核的要求，和/或研究concurrentKernels 示例应用程序。
这里的大部分信息都是从这里完成的测试/工作中回收的，尽管这里使用单独的应用程序的演示与那里提供的 MPI 案例不同。

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

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

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

1 回答 1

Related

Reference