cuda - Nvidia マルチプロセスサービス (MPS) を使用して複数の非 MPI CUDA アプリケーションを実行するにはどうすればよいですか?

Question

MPS を使用する NVIDIA Kepler GPU で非 MPI CUDA アプリケーションを同時に実行できますか? 私のアプリケーションは GPU を十分に活用できないため、これを行いたいので、それらを一緒に実行したいと考えています。これを行うコード例はありますか?

score 37 · Accepted Answer

必要な手順は、MPS サービスのドキュメントに含まれています。これらの命令は実際には MPI に依存したり、MPI を呼び出したりするわけではないので、MPI 固有のものは何もないことに注意してください。

これがウォークスルー/例です。

さまざまな要件と制限については、上記のリンク先のドキュメントのセクション 2.3 を参照してください。これには、CUDA 7、7.5、またはそれ以降を使用することをお勧めします。CUDA MPS の以前のバージョンとの構成の違いがいくつかありましたが、ここでは説明しません。また、単一のサーバー/単一の GPU を使用したデモも行います。テストに使用しているマシンは、CUDA 7.0 を搭載した K40c (cc3.5/Kepler) GPU を使用する 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
$

MPS サーバーを起動するには、bash スクリプトを使用します。

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

そして、テストアプリの 2 つのコピーを「同時に」起動する bash スクリプト:
```
$ 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 サーバーを実際に有効にしないと、アプリの 1 つのインスタンスが予想される ~5 秒かかるのに対し、もう 1 つのインスタンスはその約 2 倍 (~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 によりカーネルが同時に実行されているため、両方のアプリの実行にかかる時間が同じであることがわかります。
お好きなように実験してみてください。このシーケンスが正しく機能しているように見えても、独自のアプリケーションを実行しても期待どおりの結果が得られない場合、考えられる理由の 1 つは、アプリケーション/カーネルがアプリケーション/カーネルの他のインスタンスと同時に実行できない可能性があります。 MPS とは何の関係もありません。同時実行カーネルの要件を確認したり、 concurrentKernel サンプルアプリを調べたりすることをお勧めします。
ここでの情報の多くは、ここで行われたテスト/作業から再利用されましたが、別のアプリでのここでのプレゼンテーションは、そこでの MPI ケースとは異なります。

更新: 複数のプロセスからカーネルを実行する場合の非 MPS ケースでのスケジューラの動作は、Pascal および新しい GPU で変更されたようです。上記のテスト結果は、テストされた GPU (Kepler など) では正しいですが、Pascal またはそれ以降の GPU で上記のテストケースを実行すると、非 MPS ケースでは異なる結果が観察されます。スケジューラは、最新の MPS docで「タイムスライス」スケジューラとして説明されています発生しているように見えるのは、あるプロセスからのカーネルが完了するのを待つのではなく、スケジューラーが、未公開のルールに従って、実行中のカーネルを先取りして、別のプロセスから別のカーネルに切り替えることができるようにすることです。これは、別のプロセスからのカーネルが、CUDA ドキュメントでのその単語の従来の使用法で「同時に」実行されていることを意味するものではありませんが、上記のコードは、依存するため、(Pascal 以降の) タイムスライススケジューラによって「だまされて」います。 SM クロックを使用してカーネル期間を設定する。タイムスライススケジューラと SM クロックのこの使用の組み合わせにより、このテストケースは「同時に」実行されているように見えます。ただし、MPS ドキュメントで説明されているように、カーネル A のコードはカーネル B のコードと同じクロックサイクルで実行されていません。

上記の一般的なアプローチを使用してこれを実証する別の方法は、ここで説明されているように、SM クロックを読み取ることによって設定されるカーネル期間ではなく、ループの数によって設定されるカーネル期間を使用することです。その場合、コンパイラによってループが「最適化」されないように注意する必要があります。

cuda - Nvidia マルチプロセス サービス (MPS) を使用して複数の非 MPI CUDA アプリケーションを実行するにはどうすればよいですか?

1 に答える 1

Related

Reference

cuda - Nvidia マルチプロセスサービス (MPS) を使用して複数の非 MPI CUDA アプリケーションを実行するにはどうすればよいですか?