1. HPCシステムとしてのクラウドサービスの活用

NVIDIA A100が発表されて半年以上が経過し、少しずつですがA100を用いた性能評価結果が国内でも報告されるようになってきました。ですが、価格的に手を出しづらいという事情と、まだ供給が追いついていないというところもあるかと思います。そこでお勧めしたいのが、クラウドサービスの活用です。

IaaS (Infrastructure as a Service) であるAmazon Web Services, Google Cloud Platform, Microsoft Azure, Oracle Cloud Infrastructureなどが公開または試験的にA100を搭載したGPUインスタンスを提供しています。加えてほとんどのIaaSプラットフォームが、NVIDIA GPU向けの仮想マシンイメージを提供しており、また一部はNVIDIAが提供するNGCカタログに対応していますので、IaaSベンダーまたはNVIDIAが構築したイメージを用いて容易にプログラムの実行・評価が可能です。

また、Amazon Web ServicesとMicrosoft AzureではそれぞれAWS ParallelCluster, Azure CycleCloudという計算クラスター管理ツールを提供しています。どちらもSLURMやPBS ProといったHPCシステムで一般的に用いられるジョブスケジューリングソフトウェアを介したリソース管理を提供しており、ログインノードでジョブを投入すると必要な計算リソースの確保・初期化・ジョブ実行・リソース破棄まで自動化することが可能です。IaaSは計算インフラを確保してから破棄するまでの時間がすべて課金対象になりますので、ジョブ実行時間が課金対象となるスーパーコンピューターと比べ広く設定されています。課金対象時間の最小化だけでなく、ノードの立ち上げや設定、容易なスケール（計算に使用するノードを増やす）とう要望の多くを自動化でき、上記のようなクラスター管理ツールは非常に効果的です。

今回は、Amazon Web Services (以下、AWS）が提供するP4dインスタンスを用いたベンチマークの結果を紹介します。
この記事は2020年12月24日に行われたAWS主催イベントでの発表内容について性能評価にフォーカスし加筆修正したものです。

2. AWS P4dインスタンスについて

P4dはAWSが提供する最新のHPC向けインスタンスで、他にV100を搭載したP3、K80を搭載したP2があります。P4dインスタンスのハードウェア詳細や細かな設定についてはAWS Compute Blog: Amazon EC2 P4d instances deep diveに記載されています。

P4dの特性については、NVIDIA DGX A100との比較が最もわかりやすいかと思いますので、以下の表に違いが見える部分を記載します。現在、NVIDIA A100は40 GB/GPUのHBM2版と80 GB/GPUのHBM2e版の2種類が併売されていますが、P4dインスタンスにはHBM2版が搭載されています。

	AWS P4d	NVIDIA DGX A100
CPU	Intel Xeon Platinum 8275CL x2 sockets = 48 physical cores	AMD EPYC 7742 x2 sockets = 128 physical cores
CPU/GPU間の接続	PCIe gen3 x16, 片方向理論値15.75 GB/s	PCIe gen4 x16, 片方向理論値31.51 GB/s
ノード間通信性能	100 Gbps x4 links = 400 Gbps	200 Gbps x8 links = 1600 Gbps

NVIDIA A100はCPUとのインターコネクトとしてPCIe gen4 x16を提供し、双方向で約64 GB/sのデータ転送性能を提供しますが、記事投稿時点ではPCIe gen4が扱える一般的なサーバー向けCPUはAMD EPYCしか選択肢がありません。P4dインスタンスはIntel Cascade Lake-SPアーキテクチャCPUのため、接続がPCIe gen3 x16に制限されます。これは例えばシミュレーションの途中結果をファイルへ書き出すためにGPUからCPUへデータを書き戻すといった、CPU-GPU間の通信が頻繁に発生する場合に不利なことがあります。

以下はCUDA Samplesに含まれるbandwidthTestの実行結果です。P3インスタンスのV100はPCIe gen3 x16を提供しているため仕様通りの性能が得られていますが、P4dインスタンスのA100はCPUによって転送速度が制限されていることがわかります。

NVIDIA V100 (p3dn.24xlarge)
bandwidthTest-H2D-Pinned, Bandwidth = 11.2 GB/s, Time = 0.00286 s, Size = 32000000 bytes, NumDevsUsed = 1
bandwidthTest-D2H-Pinned, Bandwidth = 12.5 GB/s, Time = 0.00255 s, Size = 32000000 bytes, NumDevsUsed = 1

NVIDIA A100 (p4d.24xlarge)
bandwidthTest-H2D-Pinned, Bandwidth = 12.3 GB/s, Time = 0.00260 s, Size = 32000000 bytes, NumDevsUsed = 1
bandwidthTest-D2H-Pinned, Bandwidth = 13.2 GB/s, Time = 0.00243 s, Size = 32000000 bytes, NumDevsUsed = 1

ノード間通信性能は、DGX A100がInfiniBand HDR 8ポート（GPUと対で接続）に対し、P4dインスタンスはElastic Fabric Adapter (EFA) 4ポート（2 GPUと対で接続）が提供されています。したがって、大容量データ通信が頻繁に起きるようなプログラムでは不利な可能性があります。しかしP4dインスタンスおよびDGX A100はノードあたりNVIDIA A100を8台搭載し、倍精度演算 (FP64) で9.7 TFLOPS x8 = 77.6 TFLOPSの演算性能を提供します。1ノードでも十分な性能で計算可能と考えるユーザーも多いと予想され、ノード間通信性能は気にならないかもしれません。ノード内のGPU間通信性能は、DGX A100とP4dインスタンスはNVSwitchによりフルメッシュ接続されているため、同等性能です。

以下はOSU micro-benchmarksが提供する双方向通信性能のベンチマーク実行結果です。ここではCPU間の通信性能を測定していますが、EFA 1リンク分 (100 Gbps) の性能が得られています。P4dインスタンスではEFAを用いたGPU Direct RDMA (GDR) 機能をOpen MPIで利用できますが、設定の問題かEFA GDR通信を有効にするとMPI処理系が内部エラーで停止してしまうため、ノード間でのGPU間通信性能は測定できていません。

最後にCPUですが、プロセッサのメーカー違い（正確には両社が採用する製造プロセスルールや開発遅延）から物理コア数は約2.7倍の差があります。通常、計算はGPUに任せるのでCPUの性能差を気にする必要はありませんが、A100から新たに提供されたMulti-Instance GPU (MIG) 機能を使う場合に若干不利です。

MIGは、1台のA100を最大7つのGPUインスタンスとしてリソースを完全に分割します。MIGで分割されたGPUインスタンスはそれぞれが1台のGPUとして振る舞い、CUDAコアやメモリ帯域等を専有します。P4dインスタンスとDGX A100は8つのA100を搭載しているので、最大56個のGPUインスタンスに分割が可能です。これはP4dインスタンスが持つCPUの物理コア数より多く、56インスタンスを物理CPUコアと一対一で紐付けできず、MPIを用いた分散並列計算には利用できません。従来から提供されているCUDA Multi-Process Services (MPS) も使用できますので、今回はこちらを使用します。

今回はノード間通信性能がボトルネックとなるか、HPCGベンチマークを用いて検証します。評価にはNVIDIA HPC SDK 20.9およびCUDA 11.0を使用し、NVIDIA A100を搭載したP4dとV100を搭載したP3dnの2インスタンスを用いてV100からA100で性能がどれだけ変わるかを紹介します。

3. 姫野ベンチマークの評価

姫野ベンチマークは非圧縮性流体解析コードの性能ベンチマークで、HPCでは特によく知られたベンチマークのひとつです。単精度浮動小数点数 (FP32) でポアソン方程式をJacobi反復法を用いて解き、その求解性能を測定します。

非常に単純なループ構造のため、OpenACCで簡単にGPU化することが可能です。従来、OpenACCはCrayといったスーパーコンピューターのベンダーが各社のシステム専用に提供するコンパイラか、旧PGIコンパイラでしか実運用可能な実装が提供されていませんでしたが、旧PGIコンパイラはNVIDIA HPC SDKのHPCコンパイラとして無償提供されています。

以下はC言語のstatic allocate versionをGPU化した場合のサンプルコードです。5行の#pragma accディレクティブの挿入だけで、測定対象となるループ全体をGPU化することができます。

#pragma acc data copyin(a,b,c,wrk1,bnd) copy(p) create(wrk2)
for(n=0 ; n<nn ; ++n){
  gosa = 0.0;
#pragma acc kernels loop gang
  for(i=1 ; i<imax-1 ; i++)
    for(j=1 ; j<jmax-1 ; j++)
#pragma acc loop vector(256) reduction(+:gosa)
      for(k=1 ; k<kmax-1 ; k++){
        s0 = a[0][i][j][k] * p[i+1][j ][k ] ... ;
        ...;
      }

#pragma acc kernels loop gang
  for(i=1 ; i<imax-1 ; ++i)
    for(j=1 ; j<jmax-1 ; ++j)
#pragma acc loop vector(256)
        for(k=1 ; k<kmax-1 ; ++k)
          p[i][j][k] = wrk2[i][j][k];
}

ヒープ領域に計算領域を確保するdynamic allocate versionの場合、CPU-GPU間のデータ共有を自動化するunified memoryを使用可能です。ただし、姫野ベンチマークは構造体Matrixを静的変数で宣言しているため、以下のように構造体自体はdata構文を用いてGPUに転送する必要がありますが、実際のデータを収めたメンバー変数Matrix::mはunified memoryによりCPU-GPU間のデータ転送を自動化可能です。

#pragma acc data copyin(a,b,c,wrk1,wrk2,bnd,p)
for(n=0 ; n<nn ; ++n){
  gosa = 0.0;

  ...
}

ビルド時にnvcコンパイラに-accオプションを指定するとOpenACCディレクティブを有効にしてビルドできます。-gpuオプションにCompute Capability 7.0, 8.0を明示することで、V100およびA100向けのGPUカーネルバイナリが生成された1つの実行ファイルが作成されます。unified memoryを使用する場合は、-gpuオプションにmanagedを付与します。

nvc -O3 -fast -Kieee -acc=strict -Minfo=accel -gpu=cc70,cc80,fma,fastmath -o xbmt himenoBMTxp.c

C言語版のstatic allocateとdynamic allocate versionの2つについて、シングルGPUで実行した場合の性能が以下の図です。コードはV100とA100で共通なのでGPUの性能差比較と言えますが、V100に対しほぼ2倍近い性能が得られています。姫野ベンチマークはメモリバンド幅に律速されるので、V100とA100のメモリバンド幅の比と計算性能の比が等価であれば妥当と言えますが、メモリバンド幅はV100が900 GB/sに対し、A100は1.555 TB/sです。比は約1.7なので、OpenACCでシンプルに実装した場合でもメモリバンド幅の比以上の性能向上が得られていることになります。L2キャッシュのRead速度がV100から約2.3倍に引き上げられ、この改良点が大きく寄与しているものと考えられます。static/dynamicともにループの並列化は同じ挙動になるように調整していますが、dynamicはstaticに比べ性能に優位性が見られます。どちらの実装も最後にCPUへのデータ転送が完了するまでの間を計測範囲としているため、計算開始から終了までほぼ同じ動作です。この結果は別のコラムとして理由を探していきたいと思います。

次にCUDA Multi-Process Services (MPS) の効果について、MPI版のOpenACC実装を用いた性能評価を以下の図に示します。CUDA MPSは複数プロセスで1 GPUを共有する場合のリソースシェアリングを最適化します。V100より前はソフトウェアによる制御が行われていましたが、V100からハードウェアによる制御が実装されオーバーヘッドが短縮されました。どちらのGPUもMPSを有効にすることで効率的な実行ができていますが、V100では1プロセスの実行と性能が大きく変わらないのに対し、A100は性能低下が見られます。複数プロセス実行の場合はMPI通信が発生しますが、A100の性能が高すぎるためにMPI通信がすべて実行オーバーヘッドになり性能劣化を招いていると考えられます。

最後に、AWS P3dnおよびP4dインスタンスの1ノードが持つ8台のGPUを使った場合の性能を以下の図に示します。V100の場合は1 GPUに複数プロセスを割り当てた方が、専有状態（各GPUに1プロセスを割り当てる場合）よりも高い性能が得られています。しかしA100は専有状態の方が効率よく実行できていて、CUDA MPSの効果がないように見えます。次に示すHPCGベンチマークでA100での効果について考察します。

4. HPCG (High Performance Conjugate Gradient) ベンチマークの評価

HPCGは、世界のスーパーコンピューターのランキングであるTOP500で用いられているベンチマークの1つです。TOP500は従来HPL (High Performance Linpack)のみで計算性能を計測してきましたが、現在の計算機とシミュレーションソフトウェアはHPLではカバーできないデータ転送速度がボトルネックとなることが多く、HPCGはそのメトリックを測定する手段として提案されました。TOP500は依然としてHPLの実行性能でランキングされますが、HPCGの結果も併記されます。

HPCGは疎行列を係数行列に持つ連立一次方程式（3次元熱拡散方程式）について、幾何的マルチグリッド法を前処理に用いたCG法を用いて解き、その求解性能を評価します。疎行列の求解法は一般的にメモリやネットワークなどのデータ転送性能に律速され、密行列の計算性能を求めるHPLよりも現実のソフトウェアに近いベンチマークとされています。しかし測定結果はアルゴリズムの選択や実装方法に依存しており、各システムの結果が対等な性能評価とはなっていないのではないか、という批判もあります。NVIDIA GPU向けの実装がHPCGより公式に提供されていますので、少なくとも同じ実装での評価と比較は問題ないと考えます。

以下にAWS P3dnおよびP4dインスタンス4ノード、合計32 GPUまでの性能評価結果を記載します。問題サイズは、NVIDIA GPUに適切とされるプロセスあたり (256, 256, 256) のWeak scalingです。殆どの場合PPG = 4のときにはメモリ不足で実行できませんが、姫野ベンチマークとは対照的にA100でも複数プロセス実行の実行効率が向上しており、CUDA MPSがA100でも有効と言えます。32 GPUのメモリバンド幅の合計値はV100で28.8 TB/s, A100で49.76 TB/sなので、ノード間通信性能に律速されることなく、合計メモリバンド幅に相当する性能が得られていると言えます。

5. まとめ

AWSのP4dインスタンスについてデータ転送性能が鍵となる2つのベンチマークを評価し、GPU単体では約2倍、複数台では1.5倍以上の性能がハードウェアの進化によって得られました。以前よりV100を用いて実装されていたソフトウェアはA100でも十分な性能を実現可能と考えられ、またこれから実装する場合でもOpenACCが有力な選択肢として提示できるものと思います。AWSのインスタンスとして考えると、前モデルのP3dnと時間単価はほぼ変わりません。1時間あたり32.77ドルですから、物理的に購入するよりも遥かに安くA100をテストすることが可能です。また、OpenACCで書かれたアプリケーションも十分な性能で動くことが示唆できたかと思います。今後はより実用的なアプリケーションで性能評価した結果をご紹介できればと考えております。

HPC WORLDの技術情報ページでは、今後もHPCに関する内容をコラム形式で投稿していく予定です。HPCメーリングリストも合わせてご利用いただければ幸いです。