GTC 2014 - OpenACCとMPIやOpenMPを組み合わせて複数GPUを使う(前編)

1つのGPUでは性能が足りない場合は複数のGPUを使うことになる。OpenACCを使うGPUプログラムを書き方は、OpenACCの入門編と上級編でカバーしてきたが、これだけでは複数のGPUを使うプログラムは作れない。GTC 2014において、NVIDIAのPGIコンパイラ部門のMichael Wolfe氏が複数GPUを使う方法を説明した。

MPIを使って複数GPUを使う

複数GPUを使う第1のやり方はMPIとOpenACCを組み合わせることである。MPI(Message Passing Interface)では、それぞれのノードがCPUとメモリを持ち、メモリ空間は独立で、他のノードとの通信は、ネットワークのアドレスを指定してメッセージを送受信することで行われる。

このようなノードのCPUにGPUを接続して、各CPUでは、OpenACCで書かれたGPUに並列度の高い処理をオフロードするプログラムを走らせれば、複数GPUを使うシステムが作れる。

また、各ノードで複数のMPIプロセスを走らせ、それぞれのMPIプロセスがOpenACCで書かれていれば、論理的にはMPIプロセスの数だけのGPUを1つのノードで使うことができる。

例えば、1個のCPUに2台のGPUを接続しMPIプロセスを、2つのコアで走らせれば、1CPU＋2GPUの計算ノードができ、これをInfiniBandやGbit Ethernetで接続してMPIライブラリを載せれば、マルチノードのコンピュータシステムができる。現在の大部分のスパコンは、このような構造になっている。

このようなシステムでは、図1に示すように、まず、MPI_Initでネットワークに接続されているノードを検出し、各ノードでMPIプロセスを起動する。各MPIプロセスにはrankと呼ぶ、一連の番号が付けられる。

そして、MPI_Comm_rankで、各MPIプロセスは自分のrank番号を取得する。図1の例では、OpenACCのacc_get_num_devices関数を呼び、そのCPUに接続されているGPUの個数ngpusを取得する。GPUが接続されておりngpusが0でなければ、acc_set_device_num関数を呼び出してrank%ngpusを使用するdevice番号にセットする。GPUが1台の場合は、使用するdevice番号は0で、2台のGPUが接続されている場合は、そのMPIプロセスのrank番号によって、device番号は0か1になる。

これで、OpenACCでGPUを使う用意ができる。後は、OpenACCで、GPUを使うプログラムを書けばよい。各MPIランクでは同じプログラムが動いているのであるが、メモリ空間は共通ではないので、MPIランク間のデータのやり取りはMPI_SendやMPI_Recvなどを呼び出して、明示的に記述する必要がある。残念ながら、MPIにはOpenACCのデータリージョンのように自動的にデータをコピーしてくれる機能は無い。

OpenMPで複数のGPUを使う

複数GPUを使う第2の方法はOpenMPを使う方法である。OpenMP 4.0では、GPUを含むアクセラレータのサポートが追加されたが、ここでは、4.0以前のCPUだけのマルチスレッド並列化だけを想定している。

この場合は、OpenMPで共通メモリの計算ノード内に複数のスレッドを作って、処理を分担する。そして、それぞれのスレッドがCPUに接続されたGPUを使用する。

図2の例では、まず、OpenMPのomp_get_thread_num関数で自分のスレッド番号tnumを得て、GPUが接続されている場合はtnum%ngpusを、そのスレッドが使用するGPUのデバイス番号としてセットする。複数台のGPUが接続されている場合は、tnumの値で異なったGPU番号が割り当てられるのは、前のMPIの場合と同様である。

そして図3のように、OpenACCのacc_get_num_devices関数を呼んで接続されているGPUの個数numgpusを得て、OpenMPの#pragma omp parallel num_threadsディレクティブでGPUの個数分のスレッドを生成させる。そして、各スレッドでは、自分のスレッド番号を取得してOpenACCのデバイス番号にセットする。この部分の説明は図2と重複する部分があるが、ここではWolfe氏のスライドの記述に合わせて説明している。

そして、図3の例では#pragma acc data copyでデータリージョンを開始し、配列x[0:n]のデータをGPUに転送して処理を行っている。

図4は、OpenMPの#pragma omp parallel num_threadsディレクティブでコア数分のスレッドを生成し、各スレッドは#pragma acc data present_or_copyディレクティブでデータリージョンを開始する例である。そして、その中でOpenACCの#pragma acc parallel loopでGPUでの並列処理を記述している。

この例では各GPUは配列xのtnum*szからsz個の要素の処理を分担しており、処理する範囲は重複していないので問題は発生しない。

しかし、図5のように、0番のGPUは配列xの0からsz+1個の要素を処理し、1番のGPUは配列xのsz要素からsz+1個の要素を処理するというように異なるGPUがオーバラップする領域を処理する場合、GPUのメモリは独立であるので、配列への書き込みがある場合は、意図したように動作しないので注意が必要である。