A Compiler Framework for Optimizing Dynamic Parallelism on GPUs

[yukarinoki]

https://ieeexplore.ieee.org/abstract/document/9741284

Clangで実装されてる コンパイラはここ！！ https://figshare.com/articles/software/A_Compiler_Framework_for_Optimizing_Dynamic_Parallelism_on_GPUs/17048447/1

[yukarinoki]

GPU上の動的並列性は、GPUスレッドが動的に他のGPUスレッドを起動できることを可能にします。これは特にネストされた並列性の量が不規則で、事前に予測することができないアプリケーションで有用です。しかし、小さいグリッドを大量に起動するとオーバーヘッドが大きいことが知られています。

この問題を解決するために、ネストされた並列性を持つアプリケーションでの動的並列性の使用を最適化するコンパイラフレームワークを提案します。このフレームワークは、スレッショルド、コアセニング、集約という3つの主要な最適化を特徴とします。

• スレッショルドは、子スレッドの数がある閾値を超えた場合にのみ動的にグリッドを起動し、それ以外の場合は親スレッド内で子スレッドをシリアライズすることを含みます。
• コアセニングは、一つのコアセンされたブロックが複数のスレッドブロックの作業を実行し、それらの間で共通の作業を分散することを含みます。
• 集約は、複数の子グリッドを一つの集約されたグリッドに結合することを含みます。
  • 一つの操作で、一気に複数の子グリッドを立ち上げるということだね

[yukarinoki]

ソフトウェア技術の一つのカテゴリーは、親グリッド内のスレッドが動的に起動することなくネストされた並列作業を行うものです[9, 42]。これらの技術は、完全に動的並列性を回避することでそのオーバーヘッドを緩和しますが、親スレッドが彼らが行うべき仕事があるかどうかに関係なくスタンバイしている必要があります。もう一つのカテゴリーの技術は、複数の親スレッドによって起動される子グリッドを単一のグリッドに統合または集約するものです[14, 24, 25, 41]。これらの技術は、起動するグリッドの数を減らすことで、混雑を緩和し、グリッドのサイズを増やすことでハードウェアの利用率を向上させることで、動的並列性のオーバーヘッドを緩和します。

[yukarinoki]

最初の最適化、すなわちスレッショルドは、子スレッドの数が一定の閾値を超える場合にのみ動的にグリッドを起動し、それ以外の場合は親スレッド内で作業をシリアライズすることを含みます。この最適化は、起動するグリッドの数を減らすことで混雑を緩和し、ハードウェアを適切に活用する大規模なグリッドのみが起動されることを確認します。一部の先行研究では、スレッショルドがプログラマによって手動で適用されると仮定しています[24, 25, 34, 41]。しかし、手動でのスレッショルドの適用は、起動コードを複雑にし、コードの複製を必要とし、コードの可読性を損ないます。私たちはコンパイラ内でスレッショルド最適化を自動化することを提案し、その際に関連する課題について議論します。

二つ目の最適化、すなわちコアセニングは、複数の子スレッドブロックを一つに結合することを含みます。この最適化は、スケジュールする必要がある子スレッドブロックの数を減らし、集約最適化と相互作用して、複数の元の子ブロックの作業を通じて集約のオーバーヘッドを分散します。コアセニングはさまざまな文脈で適用される一般的な最適化です[22, 26, 32]が、私たちはこれを動的並列性の文脈で適用し、集約と組み合わせたときの利点を観察することを提案します。

三つ目の最適化である集約は、複数の親スレッドによって起動される子グリッドを単一のグリッドに集約する先行研究[14, 24, 25, 41]と似ています。しかし、先行研究における集約の粒度は、同じワープ、同じスレッドブロック、または全体のグリッドにおける親スレッドによる起動に限られています。私たちは、マルチブロック粒度での新しい集約技術を提案し、集約をさらに強化します。この技術では、起動は単一のスレッドブロックまたは全体のグリッドという先行研究で用いられていた二つの極端な範囲に対して、スレッドブロックのグループ内の親スレッド間で集約されます。

要約すると、以下の貢献を行っています： • 動的並列性に対するスレッショルドを自動化するコンパイラ変換を提案します（セクションIII）。 • 動的並列性の文脈でコアセニングを適用し、そのためのコンパイラ変換を提案します（セクションIV）。 • マルチブロックの粒度で集約を適用し、そのためのコンパイラ変換を提案します（セクションV）。 • スレッショルド、コアセニング、集約を一つのオープンソースコンパイラフレームワークに統合します（セクションVI）。

私たちの評価（セクションVIII）は、私たちのコンパイラフレームワークが、動的並列性を使用するアプリケーションに対して43.0倍、動的並列性を使用しないアプリケーションに対して8.7倍、そして先行研究で提案されたように、単独で集約を使用する動的並列性を使用するアプリケーションに対して3.6倍の幾何平均でネストされた並列性を持つアプリケーションのパフォーマンスを向上させることを示しています。

[yukarinoki]

B. 集約 集約は、複数の親スレッドによって起動された子グリッドを単一のグリッドに統合または集約する最適化です。さまざまな研究[14, 24, 25, 41]では、この最適化を手動で、またはコンパイラ内で適用しています。図1(b)は、図1(a)の例に集約がどのように適用されるかを示しています。この例では、親スレッドは協調して、彼らの子グリッドすべての累積サイズを見つけることで、単一の集約されたグリッドを起動します。先行研究では、協調の範囲が同じワープ、ブロック、またはグリッド内の親スレッド間に渡っています。範囲がワープまたはブロックを越える場合、参加する親スレッドの一つが他の代わりに集約されたグリッドを起動します。範囲がグリッドを越える場合、集約されたグリッドはホストから起動されます。協調の範囲は集約の粒度と呼ばれます。

元のコードでは、各親スレッドがその子グリッドに異なるパラメータと起動設定を提供するかもしれません。しかし、変換されたコードでは、一つのパラメータセットと起動設定のみが提供できます。このため、集約されたグリッドを起動する前に、親スレッドはそれぞれが元のパラメータと起動設定をメモリに保存し、このメモリへのポインタを集約されたグリッドに渡します。その後、子スレッドは正しいパラメータと設定をメモリからロードするために、その元の親スレッドが誰であるかを特定しなければなりません。そのため、各子スレッドブロックは検索操作を実行します。親スレッドが集約されたグリッドのサイズを特定し、それぞれのパラメータと起動設定をメモリに保存するための作業は、集約ロジックと呼ばれます。子スレッドがその元の親スレッドを特定し、元のパラメータと設定をメモリからロードする作業は、分割ロジックと呼ばれます。

[yukarinoki]

III-Bで説明したような変換に適していない子カーネルもあります。特に、2種類の子カーネルを変換しません：(1) __syncthreads()またはワープレベルのプリミティブを介してスレッド間でバリア同期を実行する子カーネル、(2) 共有メモリを使用する子カーネルです。

バリア同期を行う子カーネルについては、そのような同期をサポートしながらGPUスレッドをシリアル化することが文献[12, 18, 20, 21, 33]で行われています。主な戦略は、コードをバリアで区切られた領域に分割し、各領域の周囲にループを挿入することです。しかし、これらの手法は1つのCPUスレッドで複数のGPUスレッドをシリアル化することを目指しています。これらを1つのGPUスレッドで複数のGPUスレッドをシリアル化するように拡張することは2つの理由から実用的ではありません。第一の理由は、これらの手法が全てのローカル変数のスカラー展開を行い、バリアを越えて全てのスレッドの状態を保持することです。GPU上でこのようなスカラー展開を行うと、全てのレジスタアクセスがメモリアクセスに変換され、それが非常に高価になります。バリア同期を実行する子スレッドをシリアル化しない第二の理由は、バリア同期を含むコードは、しばしば並列アルゴリズムを実装しており、シリアル化したときに効率が悪いからです。例えば、並列化された削減操作は削減木を使用し、木のレベル間で同期するためにバリアを利用します。しかし、削減木はシーケンシャルな削減を行うための効率的な方法ではありません。シーケンシャルな削減には単純な削減ループを使用する方が効率的です。この場合、最良のシーケンシャルと並列のアルゴリズムが異なるため、プログラマーが手動で閾値最適化を適用する方が良いです。

共有メモリを使用する子カーネルについては、親スレッドごとに全体の子ブロックと同じくらいの共有メモリが必要になるため、カーネルのシリアルバージョンを構築しません。これは親ブロックの共有メモリの要件を高すぎるものにします。さらに、共有メモリを使用するカーネルは、ほとんどの場合、スレッド間で共有メモリへのアクセスを調整するために__syncthreads()を使用します。したがって、これらのカーネルは前述のバリア同期に関連する理由から、ほとんどの場合変換可能ではないでしょう。

[yukarinoki]

子スレッドの数を特定するために使用するアプローチは、プログラマーが通常、所望のスレッド数をブロックの次元で除した天井関数をグリッドの次元を計算するために使用するという観察に基づいています。天井除算は任意の方法で表現することができるため、静的解析を常に所望のスレッド数を確定的に決定することは可能ではありません。代わりに、我々はプログラマーが天井除算を計算するために使用する最も一般的なパターンを特定し、これらのパターンに基づいた単純な静的解析を行います。

図4は、天井除算を使用してグリッドの次元を計算するためにプログラマが書く一般的な表現を示しています。オプション(a)-(c)は整数演算を使用し、オプション(d)-(e)は浮動小数点に変換し、ceil関数を使用します。オプション(f)は多次元ブロックで使用され、dim3コンストラクタへのオペランドは、オプション(a)-(e)に似た表現になる可能性があります。全てのオプションにおいて、表現は全体として表現されるか、または部分的に表現されてサブ表現が中間変数に格納されるかもしれません。Nとbは任意の表現であることに注意してください。

(a) (N – 1)/b + 1 N: 必要なスレッド数 (b) (N + b – 1)/b b: ブロックの次元 (c) N/b + (N%b == 0)?0:1 (d) ceil((float)N/b) (e) ceil(N/(float)b) (f) dim3(..., ..., ...) // dim3の引数は上記の表現のいずれかになる可能性があります。

[yukarinoki]

粗粒度化という最適化は、多くの異なるコンテキストでプログラマーによって頻繁に適用されています[22]。先行研究でも、コンパイラで粗粒度化を適用していますが[26, 32]、動的並列性のコンテキストではありません。我々は動的並列性のコンテキストで粗粒度化を適用し、その適用をコンパイラ変換によって自動化することを提案します。

Fig.5は、Fig.1(a)の例における子スレッドブロックに粗粒度化がどのように適用されるかを示しています。この例では、変換後のコードの各粗粒度化子スレッドブロックが、元のコードの2つの子スレッドブロックの作業を実行します。動的並列性のコンテキストで粗粒度化を適用する利点は、スケジュールする必要がある子スレッドブロックの数を減らすことです。さらに重要なことは、粗粒度化が集約の前に適用されると、ブロックあたりの作業量が多い子スレッドブロックが提供され、これにより彼らは分割ロジックのオーバーヘッドを償却する能力が高まります。