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

我可以在带有MPS的NVIDIA Kepler GPU上同时运行非MPI CUDA应用程序吗？我想这样做是因为我的应用程序无法充分利用GPU,所以我希望它们能够共同运行.有没有代码示例呢？

必要的说明包含在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 
#include 

#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 
#include 
#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.对于在(例如Kepler)上测试的GPU,上述测试结果仍然是正确的,但是当在Pascal或更新的GPU上运行上述测试用例时,将在非MPS情况下观察到不同的结果.调度程序在最新的MPS文档中被描述为"时间分片"调度程序,并且似乎正在发生的是,不是等待来自一个进程的内核完成,调度程序可以根据一些未发布的规则选择预先 - 运行内核,以便它可以从另一个进程切换到另一个内核.这仍然不意味着来自单独进程的内核在CUDA文档中对该单词的传统用法"并发"运行,但上面的代码被时间分片调度程序(在Pascal和更新)上"欺骗",因为它取决于使用SM时钟设置内核持续时间.时间片调度器的组合加上SM时钟的这种使用使得该测试用例看起来"同时"运行.但是,如MPS doc中所述,当A和B源自非MPS情况下的单独进程时,来自内核A的代码不在与来自内核B的代码相同的时钟周期中执行.

一种替代的方法来证明这一点使用上述的一般的方法可能是使用由一个数字环路,而不是通过读取SM时钟设定,如所述的内核持续时间设置内核持续时间在这里.在这种情况下必须小心,以避免编译器使循环"优化".