OpenMP* を意識した最適化

同カテゴリーの次の記事

インテル® oneAPI 2022.3 リリース

この記事は、GitHub* の LLVM で公開されている 2022年10月13日現在の「OpenMP-Aware Optimizations」を、インテル社の許可を得て iSUS (IA Software User Society) が翻訳した日本語参考訳です。原文は更新される可能性があります。原文と翻訳文の内容が異なる場合は原文を優先してください。


LLVM は、バージョン 11 (2020年10月12日) (英語) 以降、OpenMP* を意識した最適化パスをサポートしています。この最適化パスでは、OpenMP* に特化した分野の見地からモジュールの最適化を試みます。このパスは、OpenMP* サポートを有効にしてコンパイルする際に、高度な最適化レベル (O2 と O3) を指定するとデフォルトで有効になります。

OpenMPOpt

OpenMPOpt には、いくつかの OpenMP* を考慮した最適化が含まれます。この最適化パスは、モジュール全体に対し早期に適用され、コールグラフ全体に対してはそれ以降に適用されます。OpenMPOpt による最適化のほとんどは、リマークをサポートします。次のフラグを使用してコンパイルすることで、最適化のリマークを有効にできます。

$ clang -Rpass=openmp-opt -Rpass-missed=openmp-opt -Rpass-analysis=openmp-opt

OpenMP* ランタイム呼び出しの重複排除

OpenMP* ランタイム・ライブラリーでは、OpenMP* 仕様の機能を実装する際に使用されるいくつかの関数が用意されています。ランタイム呼び出しのいくつかは、並列領域内で一定です。一般的な最適化は、不変的なコードを単一の参照に置き換えることですが、その場合コンパイラーはランタイム・ライブラリーへの不透過な呼び出しのみを認識します。これを回避するため、OpenMPOpt は一定である OpenMP* ランタイム関数のリストを保持して、手動で重複を排除します。

グローバル化

OpenMP* 仕様では、異なるスレッド間でデータが共有できることを要求されています。この要件は、GPU アクセラレーターにオフロードを行う際に課題となります。GPU スレッド間では、デフォルトでデータを共有できないため、グローバルメモリーまたは共有メモリーを介してデータを共有する必要があります。正しい OpenMP* プログラムを作成するには、変数が共有される可能性があるたびにこれを行わなければなりません。残念ながら、これはパフォーマンスに重大な影響を与えるため、ほとんどの場合使用されません。次の例では、Clang がオフロード領域のコードを生成する際に、変数 x がエスケープされ、共有される可能性があることを認識することができます。この場合、変数をグローバル化する必要があり、デバイス上のレジスターに常駐することはできません。

void use(void *) { }

void foo() {
  int x;
  use(&x);
}

int main() {
#pragma omp target parallel
  foo();
}

すべてのケースにおいてこの変換が必要ではありませんが、かなりのパフォーマンス・ペナルティーが伴います。そのため、OpenMPOpt はプロシージャー間の最適化によって、グローバル変数が使用されているかどうかをスキャンして、さらにほかのスレッドが取得することで共有される可能性を判断できます。実際に取得されていない場合、安全に高速レジスターに値を戻すことができます。

もう1 つ考えられるのは、スレッド間で共有されているメモリーが、1 つのスレッドからほかのすべてのスレッドに共有されるケースです。そのような変数は、ランタイム・ライブラリーを介することなく、コンパイル時に共有メモリーに移動できます。これにより、開発者はカスタム OpenMP* アロケーターを使用したり、ランタイムに依存することなくデバイス上で共有メモリーを確保できます。

void share(void *);

static void foo() {
  int x[64];
#pragma omp parallel
  share(x);
}

int main() {
  #pragma omp target
  foo();
}

これらの最適化は、パフォーマンスに重大な影響を与える可能性があります。また、これらの最適化はプロシージャー間の解析にかなり依存します。そのため、オフロード・アプリケーションは単一のコンパイル単位にとどめ、可能な限り関数の外部参照を避けることが理想的です。OpenMPOpt はリマークが有効である場合、グローバル化された呼び出しが残っていることを開発者に通知します。これは、プログラムの問題として扱うべきです。

関連情報

  • •2021 OpenMP* ウェビナー: 「コンパイラーから見た OpenMP*」 https://youtu.be/eIMpgez61r4 (英語)
  • 2020 LLVM 開発者ミーティング: 「(OpenMP*) 並列化を意識した最適化」 https://youtu.be/gtxWkeLCxmU (英語)
  • 2019 EuroLLVM 開発者ミーティング: 「コンパイラーによる GPU への (OpenMP*) target オフロード向けの最適化」 https://youtu.be/3AbS82C3X30 (英語)

関連記事

  • インテル Parallel Universe 47 号日本語版の公開インテル Parallel Universe 47 号日本語版の公開 インテル Parallel Universe マガジンの最新号 (英語) が公開されました。 注目記事: インテルの CPU と GPU 向けの LLVM と GCC のベクトル化 掲載記事 OpenMP* を使用した効率良いヘテロジニアス並列プログラミング ArrayFire* と […]
  • DPC++ への Codeplay の貢献により NVIDIA* GPU の SYCL* サポートを提供 この記事は 2020 年 2 月 3 日に Codeplay のウェブサイトで公開された「Codeplay contribution to DPC++ brings SYCL support for NVIDIA GPUs」を Codeplay の許可を得て日本語訳したものです。 Codeplay は当初から SYCL* […]
  • インテル® oneAPI 2021.4 リリースインテル® oneAPI 2021.4 リリース この記事は、インテル® デベロッパー・ゾーンに公開されている「Intel® oneAPI 2021.4 update available」の日本語参考訳です。 公開日: 2021 年 10 月 1 日 インテル® oneAPI ツールキットの最新のアップデート (2021.4) […]
  • OpenMP* オフロードの最良の事例OpenMP* オフロードの最良の事例 この記事は、インテル® デベロッパー・ゾーンに公開されている『oneAPI GPU Optimization Guide』の「OpenMP Offload Best Practices」からの抜粋の日本語参考訳です。原文は更新される可能性があります。原文と翻訳文の内容が異なる場合は原文を優先してください。 OpenMP* […]
  • とらえどころのないアルゴリズム – 並列スキャン (追記)とらえどころのないアルゴリズム – 並列スキャン (追記) この記事は、インテル® デベロッパー・ゾーンに公開されている「Elusive Algorithms - Parallel Scan」の日本語参考訳です。 先月、IDZ の MIC フォーラムでの問い合わせ "C 言語でインテル® Cilk™ Plus を使用して包括的なスキャンを行うには?" […]