2

並列 GPU コード (OpenACC で記述) の強力なスケーリングを調査したいと思います。GPU によるストロング スケーリングの概念は、少なくとも私の知る限りでは、CPU よりも曖昧です。GPU での強力なスケーリングに関して私が見つけた唯一のリソースは、問題のサイズを修正し、GPU の数を増やすことを提案しています。ただし、ストリーミング マルチプロセッサ (Nvidia Kepler アーキテクチャ) のスケーリングなど、GPU内にはある程度の強力なスケーリングがあると思います。

OpenACC と CUDA の目的は、ハードウェアを並列プログラマーに明示的に抽象化し、並列プログラマーをギャング (スレッド ブロック)、ワーカー (ワープ)、ベクトル (スレッドの SIMT グループ) による 3 レベルのプログラミング モデルに制約することです。CUDA モデルは、独立しており、SMX にマップされているスレッド ブロックに関してスケーラビリティを提供することを目的としていると理解しています。したがって、GPU を使用したスト​​ロング スケーリングを調査するには、次の 2 つの方法があります。

  1. 問題のサイズを修正し、スレッド ブロック サイズとブロックあたりのスレッド数を任意の定数に設定します。スレッド ブロックの数 (グリッド サイズ) をスケーリングします。
  2. 基礎となるハードウェアに関する追加の知識 (例: CUDA 計算機能、最大ワープ/マルチプロセッサ、最大スレッド ブロック/マルチプロセッサなど) を考慮して、ブロックが全体および単一の SMX を占有するように、スレッド ブロック サイズとブロックあたりのスレッド数を設定します。したがって、スレッド ブロックのスケーリングは、SMX のスケーリングと同じです。

私の質問は次のとおりです: GPU での強力なスケーリングに関する私の一連の考えは正しいですか? もしそうなら、OpenACC 内で上記の #2 を行う方法はありますか?

4

2 に答える 2

5

GPU は強力なスケーリングを行いますが、必ずしもあなたが考えている方法ではありません。そのため、複数の GPU への強力なスケーリングに関する情報しか見つけることができませんでした。マルチコア CPU を使用すると、実行する CPU コアの数を簡単に決定できるため、作業を修正し、コア全体のスレッドの程度を調整できます。GPU を使用すると、SM 間の割り当ては自動的に処理され、完全に制御できなくなります。これは設計によるものです。適切に作成された GPU コードは、プログラマーやユーザーの介入なしに、スローされた GPU (または GPU) を埋めるために強力にスケーリングされることを意味するためです。

少数の OpenACC ギャング/CUDA スレッドブロックで実行し、14 のギャングが 14 の異なる SM で実行されると仮定することもできますが、これにはいくつかの問題があります。まず、1 つのギャング/スレッドブロックが 1 つのケプラー SMX を飽和させることはありません。スレッドの数や占有率に関係なく、ハードウェアを十分に活用するには、SM ごとにより多くのブロックが必要です。第二に、ハードウェアがそのようにブロックをスケジュールすることを選択するという保証はありません。最後に、お使いのデバイスで SM ごとのブロックまたはギャングの最適な数を見つけたとしても、他のデバイスには拡張できません。GPU の秘訣は、SM が 1 のデバイスから、SM が 100 のデバイス (存在する場合)、または複数のデバイスにスケーリングできるように、できるだけ多くの並列処理を公開することです。

一定量の作業に対して OpenACC ギャングの数を変えることがパフォーマンスにどのように影響するかを実験したい場合は、句 (リージョンを使用している場合) または句 (リージョンnum_gangsを使用している場合) を使用してそれを行います。ループの特定のマッピングを強制しようとしているので、より規範的なディレクティブであるため、 を使用する方が本当に良いです。やりたいことは、次のようなものです。parallelgangkernelsparallel

#pragma acc parallel loop gang vector num_gangs(vary this number) vector_length(fix this number)
for(i=0; i<N; i++)
  do something

これにより、指定されたベクトル長を使用してループをベクトル化し、OpenACC ギャング間でループを分割するようコンパイラーに指示します。私が期待しているのは、ギャングを追加すると、SM の数の倍数までパフォーマンスが向上し、その時点でパフォーマンスがほぼ横ばいになることです (もちろん異常値はあります)。上で述べたように、最適なパフォーマンスが得られるポイントでギャングの数を修正することは、これが唯一のデバイスである場合を除き、必ずしも最良のアイデアではありません。代わりに、ループを分解する方法をコンパイラーに決定させることによって、 、これにより、コンパイラは、構築するように指示したアーキテクチャに基づいてスマートな決定を下すことができます。または、できるだけ多くのギャングを公開することにより、より大きな GPU または複数の GPU に強力にスケーリングする追加の並列処理が提供されます。

于 2014-12-18T17:01:06.763 に答える
0

完全な SMX を占有するには、共有メモリを占有の制限リソースとして使用することをお勧めします。SMX が別のブロックのリソースを使い果たしているため、32kB の共有メモリをすべて消費するカーネルを作成すると、ブロックが SMX 全体を占有します。ブロックを 1 から 13 (K20c の場合) にスケールアップすると、スケジューラは (うまくいけば) 各ブロックを異なる SMX にスケジュールします。最初にブロックあたりの therads を 192 にスケールアップして、各 CUDA コアをビジー状態にし、さらにワープ スケジューラを満足させることができます。GPU は、レイテンシの隠蔽を通じてパフォーマンスを提供します。したがって、SMX を占有する 1 ブロックから N ブロックに移動する必要があります。これは、使用する共有メモリを減らすことで実現できます。そして再びワープをスケールアップして、レイテンシーの隠蔽をカバーします。

私は OpenACC に触れたことはありません。実験的なコードを完全に制御したい場合は、OpenACC の代わりに CUDA を使用してください。OpenACC コンパイラの内部と、コードで使用されているプラ​​グマで何をしているのかを確認することはできません。

于 2014-11-14T09:17:54.870 に答える