複数GPUを活用するヒント

最終更新日：2018年4月23日

サブシステムBには4基のGPUが搭載されており、最大の性能を得るにはすべてのGPUを活用する必要があります。またGPU間で高速に通信を行うことができるハードウェア（NVLink）や通信ライブラリ（CUDA-Aware MPI）が存在しますが、正しく利用しなければ機能を利用できず、十分な性能を発揮できません。

このページでは複数のGPUを使う方法や GPU間で高速通信を行う方法について説明します。

なお、cuBLASのように複数GPUを利用可能なソフトウェアやライブラリもあり、これらが活用できるような用途（対象アプリケーション）であれば、大きな労力を必要とせずに複数のGPUによる高い性能を得られるかも知れません。

CUDAプログラムにおける対象GPUの指定方法

CUDA（CUDA CおよびCUDA Fortran）では対象GPUを切り替えるためのcudaSetDevice関数が提供されています。 cudaSetDevice関数で対象GPUを変更すると、その後のCUDA APIによる処理やGPUカーネルは指定したGPU上で実行される（指定したGPUに対して処理が行われる）ようになります。

さらに、CUDAが扱う対象GPUは環境変数CUDA_VISIBLE_DEVICESによって制御することも可能です。 1GPUのみを使用するプログラムを複数同時に起動することで複数GPUを使い切りたい場合には、プロセスごとに異なるCUDA_VISIBLE_DEVICES環境変数を与えることで、対象プログラムのソースコードを単一GPU向けのものから変更せずに対応可能です。 CUDA_VISIBLE_DEVICESとcudaSetDeviceを同時に使う場合は、存在しないGPU番号を指定してしまわないよう注意してください。
1ノードで4プロセスを立ち上げ、各プロセスが異なるGPUを使うようにする例：

numactl --cpunodebind=0 --localalloc CUDA_VISIBLE_DEVICES=0 ./a.out &
numactl --cpunodebind=0 --localalloc CUDA_VISIBLE_DEVICES=1 ./a.out &
numactl --cpunodebind=1 --localalloc CUDA_VISIBLE_DEVICES=2 ./a.out &
numactl --cpunodebind=1 --localalloc CUDA_VISIBLE_DEVICES=3 ./a.out

numactlでCPUとメモリの割り当てを指定しつつ、CUDA_VISIBLE_DEVICES環境変数を個別に指定しながら対象プログラムを実行しています。 CPUやメモリの配置とGPUの配置とを適切に指定することが重要です。 numactlとCPU/メモリの配置についてはプロセス・スレッドの割り当て方をご確認ください。（「プロセス・スレッドの割り当て方」のページの図にあるように、サブシステムBはCPU0からGPU0と1、CPU1からGPU2と3が近い構成になっています。）

高速なGPU間通信の使い方

ノード内の異なるGPU間における通信の高速化（GPU間直接通信によるcudaMemcpyの高速化）

1プロセス内における複数GPU間（デバイスメモリ間）のデータ転送はcudaMemcpy関数で簡単に行うことができます。しかし単純にcudaMemcpyを用いてGPU間のデータ転送を行おうとした場合、 GPU上のデバイスメモリにあるデータは一度ホストメモリに転送され、その後に改めて別のGPU上のデバイスメモリに転送されます。受け取る側でも同様にホストメモリとデバイスメモリ間のコピーが必要となるため、通信性能の低下（通信遅延の増大）を招きます。しかし、あらかじめcudaDeviceEnablePeerAccess関数を実行しておくと、ホストメモリを介さないGPU間直接通信を行わせることが可能になり、通信性能を向上させることができます。

コード例

  cudaSetDevice(gid_from);
  cudaMalloc((void**)&d_from, sizeof(double)*N);
  cudaSetDevice(gid_to);
  cudaMalloc((void**)&d_to, sizeof(double)*N);
  cudaDeviceEnablePeerAccess(gid_to, 0); // 違いはこの一行の有無のみ
  cudaMemcpy(d_to, d_from, sizeof(double)*N, cudaMemcpyDefault);

性能例

指定サイズ（横軸Sizeに対応）のcudaMemcpy(cudaMemcpyDefault)を多数繰り返した際の時間を測定して比較しました。（cuda/8.0およびintel/2017環境にて測定。図1-dは縦軸が対数であることに注意。）
図1-a. GPU間通信時間

図1-b. GPU間通信時間（縦軸限定版）

図1-c. GPU間通信速度

図1-d. GPU間通信速度（縦軸限定版）

凡例の数字（0,1,2,3）はGPU番号に対応し、 0-1であればGPU0からGPU1への通信であることを意味します。いくつか重なってしまっているグラフがあるためわかりにくいですが、全体としては「peer 0-1」「peerの0-2と0-3」「defaultの3種」の3グループに分類できます。
特に転送サイズが大きな場合に注目すると、「peer 0-1」は同一CPUソケットにつながった2つのGPUであり、最大で37GB/sec程度の転送性能が得られています。これはNVLinkを2リンク使った直接通信でなくては得られない性能です。次点の「peer 0-2と0-3」はそれぞれ18.5GB/sec程度の転送性能が得られており、 NVLinkを1リンク使った性能と考えると妥当です。最後に「defaultの3種」は最大で10GB/sec程度であり、PCE-Expressバス経由の転送性能であると考えるのが妥当です。
転送サイズが小さい場合に注目すると、「peer 0-1」と「peerの0-2と0-3」は差が見られず、「defaultの3種」のみが遅いという結果が得られており、 NVLinkを使ったGPU間直接通信の優位性がよくわかります。 peerの3種類の性能がどれも同程度であることから、 NVLinkのリンク数が影響してくるのは転送サイズが16KB程度以上の場合であることが確認できます。
以上のように、GPU間直接通信はとても簡単に使える割にとても効果が高いため、 1プロセスで複数のGPUを使用しGPU間の直接通信が存在するようなプログラムにおいては是非活用してください。

GPU間におけるMPI通信の高速化（CUDA-Aware MPIの活用）

通常のMPIライブラリ（を使った複数プロセスプログラム）は、メインメモリ上のデータをプロセス間で送受信します。一方、GPUプログラムにおいてMPI通信を行うことを考えてみると、 GPU上で行った計算の結果はGPU上のメモリ（デバイスメモリ）に格納されているため、一度ホストメモリに転送してから通信せねばなりません。これでは（転送速度は確保できるかも知れませんが）転送に遅延が生じるため、より高速・低遅延にデータを転送するための仕組みが用意されています。 CUDA-Aware MPIを使えば、メインメモリとデバイスメモリ間の不要なデータコピーを削減し、 GPUプログラムの通信性能を高めることができます。
ITOにて提供されているCUDA-Aware MPIとしては、 OpenMPI (exp-openmpi/3.0.0-intel)と MVAPICH (exp-mvapich2/2.2-intel)があります。各MPIは性能の傾向に違いがあるため、対象プログラムの通信パターン等を考慮して利用するMPIを選択してください。

使用方法

IntelコンパイラとOpenMPIを使う場合

export MODULEPATH=$MODULEPATH:/home/exp/modulefiles
module load intel/2017 cuda/8.0 exp-openmpi/3.0.0-intel

IntelコンパイラとMVAPICHを使う場合

export MODULEPATH=$MODULEPATH:/home/exp/modulefiles
module load intel/2017 cuda/8.0 exp-mvapich2/2.2-intel

PGIコンパイラ（と同梱のOpenMPI）を使う場合

module load pgi/17.7 pgi/2017_openmpi

CUDA-Aware MPIを使うとデバイスメモリを指す変数（配列）を直接MPI通信時の送受信対象として指定できます。プログラムの記述しやすさと通信性能の両方のメリットがあるため、 GPUのみを計算に利用するプログラムでは積極的に利用してください。

プログラム例

MPIランク0のGPUにあるデータを、MPIランク1のGPUで処理して、MPIランク0のGPUに戻す例を示します。

CUDA-Aware MPIを使わないGPU間通信の場合

#include <cuda_runtime.h>

// メインメモリの確保
data = (double*)malloc(sizeof(double)*N);
// デバイスメモリの確保、初期データの準備
cudaSetDevice(gpuid);
cudaMalloc((void*)&d_data, sizeof(double)*N);
cudaMemcpy(d_data,data,sizeof(double)*N,cudaMemcpyHostToDevice);
if(myrank==0){
  // デバイスメモリからホストメモリへ転送し、MPI送信
  cudaMemcpy(data, d_data, sizeof(double)*N, cudaMemcpyDeviceToHost);
  ierr = MPI_Send(data, N, MPI_DOUBLE, 1, 0, MPI_COMM_WORLD);
  // MPI受信し、ホストメモリからデバイスメモリへ転送
  ierr = MPI_Recv(data, N, MPI_DOUBLE, 1, 0, MPI_COMM_WORLD, &status);
  cudaMemcpy(d_data, data, sizeof(double)*N, cudaMemcpyHostToDevice);
}else if(myrank==1){
  // MPI受信し、ホストメモリからデバイスメモリへ転送
  ierr = MPI_Recv(data, N, MPI_DOUBLE, 0, 0, MPI_COMM_WORLD, &status);
  cudaMemcpy(d_data, data, sizeof(double)*N, cudaMemcpyHostToDevice);
  // GPU上で処理
  gpukernel<<<...>>>(...);
  // デバイスメモリからホストメモリへ転送し、MPI送信
  cudaMemcpy(data, d_data, sizeof(double)*N, cudaMemcpyDeviceToHost);
  ierr = MPI_Send(data, N, MPI_DOUBLE, 0, 0, MPI_COMM_WORLD);
}

CUDA-Aware MPIを使ったGPU間通信の場合

#include <cuda_runtime.h>

// メインメモリの確保
data = (double*)malloc(sizeof(double)*N);
// デバイスメモリの確保、初期データの準備
cudaSetDevice(gpuid);
cudaMalloc((void*)&d_data, sizeof(double)*N);
cudaMemcpy(d_data,data,sizeof(double)*N,cudaMemcpyHostToDevice);
if(myrank==0){
  // デバイスメモリを直接MPI送信
  ierr = MPI_Send(d_data, N, MPI_DOUBLE, 1, 0, MPI_COMM_WORLD);
  // デバイスメモリへ直接MPI受信
  ierr = MPI_Recv(d_data, N, MPI_DOUBLE, 1, 0, MPI_COMM_WORLD, &status);
}else if(myrank==1){
  // デバイスメモリへ直接MPI受信
  ierr = MPI_Recv(d_data, N, MPI_DOUBLE, 0, 0, MPI_COMM_WORLD, &status);
  // GPU上で処理
  gpukernel<<<...>>>(...);
  // デバイスメモリを直接MPI送信
  ierr = MPI_Send(d_data, N, MPI_DOUBLE, 0, 0, MPI_COMM_WORLD);
}

コンパイル例

IntelコンパイラとOpenMPIを使う場合

mpicc -O3 -qopenmp -I/usr/local/cuda-8.0/include -L/usr/local/cuda-8.0/lib64 ./test.c

IntelコンパイラとMVAPICHを使う場合

mpicc -O3 -qopenmp -I/usr/local/cuda-8.0/include -L/usr/local/cuda-8.0/lib64 ./test.c

PGIコンパイラ（と同梱のOpenMPI）を使う場合

mpicc -fast -mp -tp=haswell -Mcuda=cc60,cuda8.0 ./test.c

実行例

※利用するGPU番号が実行時引数で与えられるプログラムになっていると仮定し、 numactlのパラメタや実行時引数をある程度自由にいじるためrun.shというスクリプトを介していますが、プログラムの作りによってはmpirunから直接実行ファイルを呼び出しても構いません

IntelコンパイラとOpenMPIの場合

#PJM -L "vnode=2"
#PJM -L "vnode-core=36"
#PJM -L "rscunit=ito-b"
#PJM -L "rscgrp=ito-g-16-dbg"
#PJM -L "elapse=10:00"
#PJM -S
export MODULEPATH=$MODULEPATH:/home/exp/modulefiles
module load intel/2017 cuda/8.0 exp-openmpi/3.0.0-intel

# 1ノード内GPU-GPU通信(同一CPUソケット)の例
FILE=./run.sh
cat<<EOF>${FILE}
#!/bin/bash
ID=\${OMPI_COMM_WORLD_RANK}
case \${ID} in
[0])
  numactl -N 0 --localalloc ./gpu2gpu 0
  ;;
[1])
  numactl -N 0 --localalloc ./gpu2gpu 1
  ;;
esac
EOF
chmod +x ${FILE}
mpirun -n 2 -display-devel-map -map-by ppr:2:socket --mca plm_rsh_agent /bin/pjrsh \
-machinefile ${PJM_O_NODEINF} --mca btl_openib_want_cuda_gdr 1 ${FILE}

# 1ノード内GPU-GPU間通信(異なるCPUソケット)の例
FILE=./run.sh
cat<<EOF>${FILE}
#!/bin/bash
ID=\${OMPI_COMM_WORLD_RANK}
case \${ID} in
[0])
  numactl -N 0 --localalloc ./gpu2gpu 0
  ;;
[1])
  numactl -N 1 --localalloc ./gpu2gpu 3
  ;;
esac
EOF
chmod +x ${FILE}
mpirun -n 2 -display-devel-map -map-by ppr:1:socket --mca plm_rsh_agent /bin/pjrsh \
-machinefile ${PJM_O_NODEINF} --mca btl_openib_want_cuda_gdr 1 ${FILE}

# 2ノード間GPU-GPU間通信の例
FILE=./run.sh
cat<<EOF>${FILE}
#!/bin/bash
ID=\${OMPI_COMM_WORLD_RANK}

case \${ID} in
[0])
  numactl -N 1 --localalloc ./gpu2gpu 2
  ;;
[1])
  numactl -N 1 --localalloc ./gpu2gpu 3
  ;;
esac
EOF
chmod +x ${FILE}
mpirun -n 2 -display-devel-map -map-by ppr:1:node --mca plm_rsh_agent /bin/pjrsh \
-machinefile ${PJM_O_NODEINF} --mca btl_openib_want_cuda_gdr 1 ${FILE}

IntelコンパイラとMVAPICHの場合

#PJM -L "vnode=2"
#PJM -L "vnode-core=36"
#PJM -L "rscunit=ito-b"
#PJM -L "rscgrp=ito-g-16-dbg"
#PJM -L "elapse=10:00"
#PJM -S

export MODULEPATH=$MODULEPATH:/home/exp/modulefiles
module load intel/2017 cuda/8.0 exp-mvapich2/2.2-intel
export MV2_SHOW_CPU_BINDING=1
# プロセス配置をnumactlに任せる場合に指定
export MV2_ENABLE_AFFINITY=0
export MV2_USE_CUDA=1
# GDRDRVを用いたデータ転送を行いたい場合に指定が必要
export MV2_GPUDIRECT_GDRCOPY_LIB=/usr/local/lib64/libgdrapi.so.1.2
export LD_PRELOAD=/home/usr0/m70000a/opt/mvapich2-2.2-gdr/lib64/libmpi.so.12.0.5

# 1ノード内GPU-GPU間通信の例
# CPUソケットへの配置はnumactlに任せている
FILE=./run.sh
cat<<EOF>${FILE}
#!/bin/bash
ID=\${MV2_COMM_WORLD_RANK}

case \${ID} in
[0])
  numactl -N 1 --localalloc ./gpu2gpu 2
  ;;
[1])
  numactl -N 1 --localalloc ./gpu2gpu 3
  ;;
esac
EOF
chmod +x ${FILE}
mpirun -n 2 -ppn 2 -launcher-exec /bin/pjrsh -machinefile ${PJM_O_NODEINF} ${FILE}

# 2ノード間GPU-GPU間通信の例
FILE=./run.sh
cat<<EOF>${FILE}
#!/bin/bash
ID=\${MV2_COMM_WORLD_RANK}

case \${ID} in
[0])
  numactl -N 1 --localalloc ./gpu2gpu 2
  ;;
[1])
  numactl -N 1 --localalloc ./gpu2gpu 3
  ;;
esac
EOF
chmod +x ${FILE}
mpirun -n 2 -ppn 1 -launcher-exec /bin/pjrsh -machinefile ${PJM_O_NODEINF} ${FILE}

PGIコンパイラ（と同梱のOpenMPI）の場合

module load の行以外はIntelコンパイラとOpenMPIの場合と同じです
（もちろん、PGIコンパイラ向けの環境変数等を使いたい場合はそれに準じます）

性能例

gccgはGPUからCPUへcudamemcpyによってデータを転送し、CPUプロセス間でMPI通信を行い、CPUからGPUへデータを転送した場合の性能です。 g2gはGPU-GPU間で直接MPI通信を行った場合の性能です。 0-1などの数字は通信に用いたGPUの番号を意味します。いずれもSizeに示されたサイズのデータを往復させる時間を100回程度測定した場合の平均実行時間です。

OpenMPIを用いた場合の性能はIntelコンパイラとPGIコンパイラで同様の傾向です。 GPU間の直接通信はノード内ノード間に関わらずCPUを経由した場合よりも高速です。 GPU間通信は、ノード内では同一のCPUソケットに接続されたGPU同士（0-1）がわずかに高速ですが、ノード間の時間についてはほぼ性能差に見分けが付きません。
なおこのバージョンのIntelコンパイラとOpenMPIの組み合わせでは 32KB以上のデータ転送がうまく行かない（制御が戻ってこない）という問題が発生しています。原因は調査中です。（一度だけの通信で送受信できるサイズと複数回に分けて送受信される必要があるサイズの境界であることが影響しているとは考えられます。）

MVAPICHについては全体的にOpenMPIより高速ですが、環境変数設定によりGDRDRVを有効化しておく必要がある点には注意が必要です。またGPU間データ転送時間については16KBのノード内GPU-GPU間転送時間が遅めであったりノード間GPU-GPU間転送時間がややばらついています。環境変数設定などで安定化させることができるかもしれませんが、現時点では確認できていません。

小さい転送サイズにおけるノード間GPU-GPU間転送時間をさらに細かく見てみると、 Intel+OpenMPIはPGI+OpenMPIよりわずかに高速、 Intel+MVAPICHは更に高速だが性能の上下が大きいようです。

NCCLによるNVLinkを用いた高速通信の活用

GPU間の直接通信を簡単に行うための通信ライブラリNCCLがNVIDIA社によって提供されています。 NCCLが提供しているのは一部のMPI集団通信に相当する機能のみですが、マルチノードマルチGPU環境において高速なGPU間通信を容易に行うことができます。

参考資料：NVIDIA Collective Communications Library (NCCL) | NVIDIA Developer