OpenMP* を意識した最適化

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

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

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 はリマークが有効である場合、グローバル化された呼び出しが残っていることを開発者に通知します。これは、プログラムの問題として扱うべきです。

OpenMPOpt

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

グローバル化

関連情報