如何使用 Nvidia 多进程服务 (MPS) 来 运行 多个非 MPI CUDA 应用程序？
How do I use Nvidia Multi-process Service (MPS) to run multiple non-MPI CUDA applications?
我可以 运行 非 MPI CUDA 应用程序同时运行在带有 MPS 的 NVIDIA Kepler GPU 上吗？我想这样做是因为我的应用程序无法充分利用 GPU，所以我希望它们共同 运行 在一起。是否有任何代码示例可以执行此操作？
MPS 服务的 documentation 中包含必要的说明。您会注意到，这些指令并不真正依赖或调出 MPI，因此实际上没有任何关于它们的东西 MPI-specific。
这是一个walkthrough/example。
阅读 above-linked 文档的第 2.3 节了解各种要求和限制。为此，我建议使用 CUDA 7、7.5 或更高版本。与先前版本的 CUDA MPS 存在一些配置差异，我不会在此处介绍。此外，我将仅使用单个 server/single GPU 进行演示。我用于测试的机器是使用 K40c (cc3.5/Kepler) GPU 和 CUDA 7.0 的 CentOS 6.2 节点。节点中还有其他 GPU。在我的例子中，CUDA 枚举顺序将我的 K40c 放在设备 0 上，但 nvidia-smi 枚举顺序恰好将它作为 id 2 放在顺序中。所有这些细节在具有多个 GPU 的系统中都很重要，影响下面给出的脚本。
我将创建几个助手 bash 脚本和一个测试应用程序。对于测试应用程序，我们想要一些内核显然可以 运行 与来自应用程序其他实例的内核并发的东西，我们还想要一些东西，当这些内核（来自单独的apps/processes) 是否同时 运行ning。为了满足演示目的的这些需求，让我们有一个应用程序，其内核仅在单个 SM 上的单个线程中 运行s，并且只需等待一段时间（我们将使用 ~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 个副本 "simultaneously":
$ 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 秒，而另一个 app/kernel 运行ning，然后花费 5 秒 运行ning 自己的内核，总共约 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，内核 运行 同时 运行ning。
欢迎您根据自己的喜好进行试验。如果此序列对您来说似乎工作正常，但 运行 您自己的应用程序似乎没有给出预期的结果，一个可能的原因可能是您的 app/kernel 无法 运行 由于内核的构造而与 app/kernel 的其他实例同时出现，与 MPS 无关。您可能想要验证 the requirements for concurrent kernels, and/or study the concurrentKernels sample app.
这里的大部分信息都是从 test/work 完成的 here 中回收的，尽管这里使用单独应用程序的介绍与那里介绍的 MPI 案例不同。
UPDATE：当 运行 来自多个进程的内核似乎已随着 Pascal 和更新的 GPU 发生变化时，non-MPS 情况下的调度程序行为。以上测试结果对于测试的 GPU（例如开普勒）仍然是正确的，但是当 运行 在 Pascal 或更新的 GPU 上运行上述测试用例时，在 non-MPS 情况下将观察到不同的结果。调度程序被描述为 the latest MPS doc 中的 "time-sliced" 调度程序，似乎正在发生的事情是，调度程序可以根据一些未发布的规则选择等待一个进程的内核完成，而不是pre-empt 一个 运行ning 内核，以便它可以从另一个进程切换到另一个内核。这仍然并不意味着来自不同进程的内核在 CUDA 文档中该词的传统用法是 运行ning "concurrently"，但是上面的代码是 "tricked" time-sliced 调度程序（在 Pascal 和更新版本上），因为它依赖于使用 SM 时钟来设置内核持续时间。 time-sliced 调度程序加上这种 SM 时钟的使用使这个测试用例看起来 运行 "concurrently"。但是，如 MPS 文档中所述，当 A 和 B 源自 non-MPS 情况下的不同进程时，内核 A 的代码与内核 B 的代码不在相同的时钟周期内执行。
使用上述一般方法来证明这一点的另一种方法可能是使用由多个循环设置的内核持续时间，而不是通过读取 SM 时钟设置的内核持续时间，如所述here。在这种情况下必须小心，以避免编译器出现循环 "optimized out"。
我可以 运行 非 MPI CUDA 应用程序同时运行在带有 MPS 的 NVIDIA Kepler GPU 上吗？我想这样做是因为我的应用程序无法充分利用 GPU，所以我希望它们共同 运行 在一起。是否有任何代码示例可以执行此操作？
MPS 服务的 documentation 中包含必要的说明。您会注意到，这些指令并不真正依赖或调出 MPI，因此实际上没有任何关于它们的东西 MPI-specific。
这是一个walkthrough/example。
阅读 above-linked 文档的第 2.3 节了解各种要求和限制。为此，我建议使用 CUDA 7、7.5 或更高版本。与先前版本的 CUDA MPS 存在一些配置差异，我不会在此处介绍。此外，我将仅使用单个 server/single GPU 进行演示。我用于测试的机器是使用 K40c (cc3.5/Kepler) GPU 和 CUDA 7.0 的 CentOS 6.2 节点。节点中还有其他 GPU。在我的例子中，CUDA 枚举顺序将我的 K40c 放在设备 0 上，但 nvidia-smi 枚举顺序恰好将它作为 id 2 放在顺序中。所有这些细节在具有多个 GPU 的系统中都很重要，影响下面给出的脚本。
我将创建几个助手 bash 脚本和一个测试应用程序。对于测试应用程序，我们想要一些内核显然可以 运行 与来自应用程序其他实例的内核并发的东西，我们还想要一些东西，当这些内核（来自单独的apps/processes) 是否同时 运行ning。为了满足演示目的的这些需求，让我们有一个应用程序，其内核仅在单个 SM 上的单个线程中 运行s，并且只需等待一段时间（我们将使用 ~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 个副本 "simultaneously":
$ 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 秒，而另一个 app/kernel 运行ning，然后花费 5 秒 运行ning 自己的内核，总共约 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，内核 运行 同时 运行ning。
欢迎您根据自己的喜好进行试验。如果此序列对您来说似乎工作正常，但 运行 您自己的应用程序似乎没有给出预期的结果，一个可能的原因可能是您的 app/kernel 无法 运行 由于内核的构造而与 app/kernel 的其他实例同时出现，与 MPS 无关。您可能想要验证 the requirements for concurrent kernels, and/or study the concurrentKernels sample app.
这里的大部分信息都是从 test/work 完成的 here 中回收的，尽管这里使用单独应用程序的介绍与那里介绍的 MPI 案例不同。
UPDATE：当 运行 来自多个进程的内核似乎已随着 Pascal 和更新的 GPU 发生变化时，non-MPS 情况下的调度程序行为。以上测试结果对于测试的 GPU（例如开普勒）仍然是正确的，但是当 运行 在 Pascal 或更新的 GPU 上运行上述测试用例时，在 non-MPS 情况下将观察到不同的结果。调度程序被描述为 the latest MPS doc 中的 "time-sliced" 调度程序，似乎正在发生的事情是，调度程序可以根据一些未发布的规则选择等待一个进程的内核完成，而不是pre-empt 一个 运行ning 内核，以便它可以从另一个进程切换到另一个内核。这仍然并不意味着来自不同进程的内核在 CUDA 文档中该词的传统用法是 运行ning "concurrently"，但是上面的代码是 "tricked" time-sliced 调度程序（在 Pascal 和更新版本上），因为它依赖于使用 SM 时钟来设置内核持续时间。 time-sliced 调度程序加上这种 SM 时钟的使用使这个测试用例看起来 运行 "concurrently"。但是，如 MPS 文档中所述，当 A 和 B 源自 non-MPS 情况下的不同进程时，内核 A 的代码与内核 B 的代码不在相同的时钟周期内执行。
使用上述一般方法来证明这一点的另一种方法可能是使用由多个循环设置的内核持续时间，而不是通过读取 SM 时钟设置的内核持续时间，如所述here。在这种情况下必须小心，以避免编译器出现循环 "optimized out"。