これは「標準並列プログラミング」シリーズの 3 回目の投稿です。このシリーズでは、標準言語による並列化をアクセラレーテッド コンピューティングで使用することの利点を取り上げています。
パート 1 では、次を説明しました。
- C++ 並列プログラミングの基礎
- 格子ボルツマン法 (LBM)
- 標準の C++ を使用して GPU で効率的に実行するために、Palabos ライブラリをリファクタリングするための最初の手順を行いました。
この投稿では、引き続き ISO C++ アルゴリズムのパフォーマンスを最適化し、MPI を使用して複数の GPU にアプリケーションを拡張します。
最適なパフォーマンスを目指す
自分のコードを CPU から GPU に移植するとき、パフォーマンスが専用 HPC コードより低くなると予想するのは当然かもしれません。結局のところ、ソフトウェア アーキテクチャの制約や、確立している API、ユーザー ベースから求められる高性能な追加機能を考慮する必要性によって、制限を受けます。それだけでなく、C++ 標準並列処理の単純なプログラミング モデルでは、CUDA のような専用言語よりも手動による微調整が少なくなります。
実際には、このパフォーマンスの損失を無視できる程度まで制御し、制限できることが多々あります。重要なのは、個々のコード部分のパフォーマンス指標を分析し、ソフトウェア フレームワークの実際のニーズを反映していないパフォーマンスのボトルネックを解消することです。
数値計算アルゴリズムの中心的コンポーネントに対して、別の原理証明 (proof-of-principle) コードを整備することが推奨されます。この手法のパフォーマンスはもっと自由に最適化することができて、完全で複雑なソフトウェア フレームワーク (Palabos の場合は STLBM ライブラリなど) に匹敵します。また、nvprof など、GPU 対応プロファイラを利用すると、パフォーマンスのボトルネックの原因を効率的に発見できます。
一般的なパフォーマンスの問題とその解決策を次の推奨事項でまとめました。
- CPU 上のデータに触れない
- アルゴリズムを知る
- パフォーマンス モデルを確立する
CPU 上のデータに触れない
パフォーマンスの低下原因として多いのが CPU と GPU メモリの間での隠れたデータ転送です。これが非常に遅くなることがあります。CUDA ユニファイド メモリ モデルでは、このタイプの転送が CPU から GPU データにアクセスするたびに行われます。1 バイトのデータにアクセスすると、メモリ ページ全体が 1 回転送されるため、壊滅的なパフォーマンス低下につながる可能性があります。
明白な解決策は、可能な限り GPU だけでデータを操作することです。そのためには、コードを注意深く見て、データへのアクセスをすべて調べ、並列アルゴリズム呼び出しにまとめる必要があります。やや荒っぽい方法ですが、どんな単純な操作にもこのプロセスは必要になります。
特に調べるべき場所は、データ統計値の後処理操作や中間評価です。もうひとつの典型的なパフォーマンス ボトルネックは、MPI 通信層で見つかります。データのパッキングとアンパッキングは GPU で実行しなければならないからです。
for_each と transform_reduce のフォーマリズムは、ほとんどの場合、均等に構造化されたメモリ アクセスに最適であるため、GPU でアルゴリズムを表現することは、「言うは易く行うは難し」です。
不規則なデータ構造の場合、競合状態を避け、これら 2 つのアルゴリズムでコアレスメモリ アクセスを保証することは困難でしょう。このような場合、次の推奨事項に従い、C++ STL で提供される並列化アルゴリズム群を理解する必要があります。
アルゴリズムを知る
ここまでのところ、並列 STL は、並列 for ループを大仰な関数構文で表すことと大差ないように思えます。実際のところ、STL は for_each と transform_reduce 以上の大規模なアルゴリズム セットを提供します。それは、並べ替えアルゴリズムや検索アルゴリズムなど、数値的手法を表現するのに便利なものです。
exclusive_scan アルゴリズムは累積和を計算するものであり、一般的に、非構造化データのインデックス再作成に便利なことから、特に言及に値します。たとえば、MPI 通信のパッキングアルゴリズムについて考えてみます。この通信では、各グリッド ノードから通信バッファーに与えられる変数の数は、事前に知る術がありません。この場合、すべてのグリッド ノードがバッファーに書き込むインデックスを決定するには、スレッド間のグローバル通信が必要です。
次のコード例は、並列アルゴリズムを使用し、GPU で良い並列効率で、この種の問題を解決する方法を示しています。
// Step 1: compute the number of variables contributed by every node.
int* numValuesPtr = allocateMemory(numberOfCells);
for_each(execution::par_unseq, numValuesPtr,
numValuesPtrl + numberOfCells, [=](int& numValues)
{
int i = &numValues - numValuesPtr;
// Compute number of variables contributed by current node.
numValues = computeNumValues(i);
} );
// 2. Compute the buffer index for every node.
int* indexPtr = allocateMemory(numberOfCells);
exclusive_scan(execution::par_unseq, numValuesPtr,
numValuesPtr + numberOfCells, indexPtr, 0);
// 3. Pack the data into the buffer.
for_each(execution::par_unseq, indexPtr,
indexPtr + numberOfCells, [=](int& index)
{
int i = &index - indexPtr;
packCellData(i, index);
} );この例では、GPU プログラミングのアルゴリズムベース手法の表現力を楽しむことができます。このコードでは、同期ディレクティブやその他の低レベルのコンストラクタを必要としません。
パフォーマンス モデルを確立する
パフォーマンス モデルは、ボトルネック分析を介してアルゴリズムのパフォーマンスの上限を確立するものです。このモデルでは通常、プロセッサのピーク パフォーマンス (FLOPS で測定) とピーク メモリ帯域幅を、制限を与える主要なハードウェア特性と見なします。
前の投稿の「例: 格子ボルツマン ソフトウェアと Palabos」セクションで説明したように、LBM コードはメモリ アクセスに対する演算の比率が低く、最新式の GPU では完全にメモリに制約されます。これは、少なくとも、単精度演算や倍精度演算向けに最適化された GPU を使用する場合です。
ピーク パフォーマンスは単純に、GPU のメモリ帯域幅とコードで実行されるメモリ アクセス数の比率で表されます。直接的な結果として、LBM コードを倍精度演算から単精度演算に切り替えると、パフォーマンスが 2 倍になります。
図 1 は、NVIDIA A100 (40 GB) GPU で単精度浮動小数点と倍精度浮動小数点に対して得られた Palabos の GPU 移植のパフォーマンスを示しています。
実行されたテスト ケースは、乱流領域におけるリッド駆動キャビティ流れで、形状は単純な立方体です。しかしながら、このケースには境界条件が含まれ、複雑な流れパターンを示します。パフォーマンスは MLUPS (Million Lattice-node Updates Per Second) で計測され、GPU メモリをピーク値で使用するという仮定の下で得られる理論ピーク値と比較されます。
このコードは倍精度でピーク パフォーマンスの 73% に到達し、単精度で 74% に到達します。そのようなパフォーマンス指標は、使用される言語やライブラリに関係なく、格子ボルツマン モデルの最先端の実装で共通です。
一部の実装では割合が少し上がり、80% に近い値に達する可能性がありますが、パフォーマンス モデルによって暗示されるハード リミットに近づいていることは明らかです。大きな視点から見ると、コードのシングル GPU パフォーマンスは、可能な限り良好なものです。
既存の MPI バックエンドを再利用してマルチ GPU コードを得る
C++ 並列アルゴリズムは既存のソフトウェア プロジェクトにシームレスに統合され、コードの重要な部分を高速化するので、プロジェクトの通信バックエンドを再利用してマルチ GPU パフォーマンスを達成することを妨げるものは何もありません。ただし、通信バッファーを監視し、CPU メモリに迂回しないようにします。迂回すると、結果的にコストのかかるページ フォールトが発生します。
複数の GPU で GPU 移植版の Palabos を初めて実行してみたところ、結果は技術的に正しいものでしたが、許容範囲のパフォーマンスが示されませんでした。1 つの GPU から 2 つの GPU へ切り替えたところ、スピードアップどころか、スピードが 1 桁下がりました。この問題の原因として、通信データのパッキングとアンパッキングが考えられます。元のバックエンドでは、これを CPU 上で行い、また通信バッファーのサイズ変更など、CPU メモリに不要なデータ アクセスが行われました。
このような問題はプロファイラを利用して発見できます。プロファイラは、ユニファイド メモリのページ フォールトの発生をすべて見つけます。ページ フォールトは、該当するコード部分を並列アルゴリズムに移行することで修正されます。「アルゴリズムを知る」セクションでは、データが不規則なパターンに従う場合に通信バッファーをパッキングし、アンパッキングする方法を説明しました。
この時点で、MPI 以外の拡張機能を使用しない標準 C++ を使用することで、シングル GPU で最先端のパフォーマンスを実現し、マルチ GPU で安定した並列パフォーマンスを実現するというハイブリッド CPU/GPU ソフトウェア プロジェクトが得られます。
残念ながら、言語仕様とそれに対応する GPU 実装の現在の制限に起因し、マルチ GPU のパフォーマンスは期待を下回っています。C++ 標準並列処理はまだ新しいテクノロジであり、将来の改良が待たれます。この投稿で提供する回避策の一部は、C++ 標準以外の手法に基づいています。
マルチ CPU とマルチ GPU のコード実行を調整する
この投稿では、CPU と GPU のハイブリッド プログラミングに焦点を当てていますが、CPU 部分のハイブリッド並列 (MPI やマルチスレッド) の問題にも対処しなければなりません。
たとえば、元のバージョンの Palabos は非ハイブリッドであり、MPI 通信層を使用して CPU のコア間だけでなくネットワーク全体で作業を分散します。GPU に移植した後、結果として得られるマルチ CPU とマルチ GPU のコードは、MPI タスクごとに 1 つの CPU コアとGPU でグループを作り、CPU が相対的にパワー不足になります。
そのため、計算負荷の高いタスクを CPU で実行する必要があるとき、そのほうが都合の良いとき、パフォーマンスのボトルネックが発生します。流体力学では、幾何処理やメッシュ生成など、前処理の段階でこのボトルネックが頻繁に発生します。
明白な解決策としては、たとえば、マルチスレッドを使用し、MPI タスク内から複数の CPU コアにアクセスします。マルチスレッドの共有メモリ空間を、CUDA ユニファイド メモリ形式により、GPU と直接共有できます。
ただし、C++ 並列アルゴリズムは、GPU とマルチコア CPU の実行という両方の目的のために利用することはできません。これは、C++ では、言語内から並列アルゴリズムのターゲット プラットフォームを選択できないためです。
C++ ではネイティブでこの問題を解決する方法が提供されますが、OpenMP から最も便利で邪魔にならない解決策が提供されていることを見つけました。この場合、現在の MPI タスクに割り当てられているグリッド部分を複数のスレッドに分配する作業は、for loop の OpenMP 注釈で十分でした。
ピン メモリを介した通信
現在のバージョンの HPC SDK では、CUDA ユニファイド メモリ モデルは MPI との組み合わせで別のパフォーマンス問題を示します。
MPI 通信層では、固定ハードウェア アドレス (いわゆる ピン メモリ) を持つデータが求められるため、マネージド メモリ領域に存在するすべてのバッファーは最初に、ホスト CPU のピン メモリ バッファーに暗黙的にコピーされます。GPU と CPU の間の転送に起因し、この操作はいくぶんコストがかかることがあります。
したがって、通信バッファーは GPU メモリ アドレスに明示的に固定する必要があります。nvc++ コンパイラを利用すると、cudaMalloc で通信バッファーを割り当てることでこれが達成されます。
// Allocate the communication buffer
// vector<double> buffer(N);
// double* buffer = buffer.data();
double* buffer; cudaMalloc((void**)&buffer, N * sizeof(double));
for_each(buffer, buffer + N, … // Proceed with data packingもうひとつの解決策は、STL の vector を Thrust ライブラリの thrust::d evice_vector に置き換えることです。これは GPU ピン メモリをデフォルトで使用します。
近い将来、HPC SDK はこのようなケースをもっと効率的かつ自動的に処理することでしょう。cudaMalloc や thrust::device_vector に頼る必要がなくなります。そのときまでお待ちください。
この投稿に記載しているさまざまな改善の後、Palabos ライブラリは、4 つの GPU を搭載した DGX A100 (40 GB) ワークステーションでテストし、リッド駆動キャビティのベンチマーク ケースでもテストしました。得られたパフォーマンスを図 2 に示し、48 コア Xeon Gold 6240R CPU で達成したパフォーマンスと比較します。
Xeon Gold の場合、Palabos のオリジナル実装はもっと効率的であることが判明し、48 個の MPI タスクで実行しました。一方、1 GPU と 4 GPU の実行では、nvc++ でコンパイルされた並列アルゴリズム バックエンドが使用されました。
パフォーマンス値は、1 GPU 実行と比較して 4 GPU 実行ではスピードが 3.27 倍にアップしたことを示しています。両方の実行でドメインサイズが等しいストロング スケーリングで、82% という極めて満足できる並列効率となります。4 GPU 実行の問題サイズを 4 倍にするウィーク スケーリングでは、スピードアップは 3.72 に上昇します (93% の効率)。
図 2 からは、MPI 通信バッファーが cudaMalloc で割り当てられないなど、ピン留めしない通信バッファーを使用するとき、並列効率が 82% から 61% に落ちることもわかります。
最後に、4 GPU DGX ワークステーションは Xeon Gold CPU の 55 倍の速さで動作します。2 つのマシンのスコープが異なるため、直接的な比較は公平ではないかもしれませんが、コードを GPU に移植することで得られる高速化がどんなものかはわかります。DGX は一般的な電源プラグに接続されるデスクトップ ワークステーションですが、CPU クラスターでは数千個の CPU コアでしか得られなかったパフォーマンスを与えます。
結論
Palabos のようなライブラリを GPU に移植する目的で C++ 標準言語による並列処理を利用するとき、コードのパフォーマンスが驚くほど改善されることがわかりました。
- Palabos ライブラリのエンド ユーザーにとっては、このパフォーマンスの向上は、コードを 1 行変更し、CPU から GPU バックエンドに切り替えることで得られます。
- Palabos ライブラリ開発者にとっては、対応する GPU バックエンドを開発するためにいくつかの作業が必要でした。
ただし、この作業では、ドメイン固有の言語を新たに学習する必要がなく、GPU アーキテクチャの詳細な知識も不要でした。
2 部からなるこの投稿では、自分のコードで同様の結果を得るために応用できるガイドラインを提供しました。詳細については、次のリソースを確認することをお勧めします。
- コンパイラ サポートについては、HPC SDK ページをご覧ください。
- 投稿記事の「標準言語による並列処理を利用したコードの高速化」をお読みください。
- HPC SDK は無料でダウンロードできます。
- C++ 標準並列処理と MPI を使用する方法は、2D 熱拡散方程式のコードを読むことで学習できます。
- LBM を GPU に実装する方法は、簡単な自己完結型のサンプルで学習できます。あるいは、完全版の STLBM ライブラリをダウンロードできます。
- Palabos GPU 移植のプロジェクト ページをご覧ください。
- Palabos をダウンロードする
