OpenMP* オフロードの最良の事例

2022.04.23
インテル® oneAPI
GPUOpenMPドキュメント最適化

この記事は、インテル® デベロッパー・ゾーンに公開されている『oneAPI GPU Optimization Guide』の「OpenMP Offload Best Practices」からの抜粋の日本語参考訳です。原文は更新される可能性があります。原文と翻訳文の内容が異なる場合は原文を優先してください。

OpenMP* オフロードの最良の事例

ここでは、GPU にオフロードするアプリケーションのパフォーマンスを向上させる最良の事例を紹介します。事例を以下のカテゴリーに分類して、以降のセクションで説明します。

  • 多くの GPU リソースを使用する
  • データ転送とメモリー割り当てを最小化する
  • OpenMP* 構造を有効活用する
  • メモリー割り当て
  • インテル® oneMKL の計算を GPU へオフロードする
  • 節: is_device_ptr、use_device_ptr、has_device_addr、use_device_addr

注: OpenMP* のパフォーマンスを収集するため次の構成を使用しました。

  • 内部バージョンのインテル® コンパイラーと GPU ドライバー
  • GPU: ATS-P✝ B0、2 タイル
  • L0-plugin
  • 起動時のオーバーヘッドを計測しないようにダミーの target 構造を挿入しています。
  • ジャストインタイム (JIT) コンパイルモードを使用します。
  • 1 タイルのみを使用します (暗黙と明示的なスケーリングはありません)

多くの GPU リソースを使用する

並列に実行できる work-item 数を多くして、より多くの GPU リソースを利用する (GPU を処理で一杯にする) ことで、オフロードされたコードのパフォーマンスを向上できます。

collapse 節

入れ子になったループの並列性を高める方法として、collapse 節を使用して入れ子になった 2 階層以上のループを 1 つのループに融合する方法があります。1 つのループに融合することで、並列に実行できる反復回数が増加し、GPU でより多くの work-item を処理できます。

次の例では、4 つの入れ子になったループを GPU にオフロードしています。parallel for ディレクティブは、一番外側のループ (47 行目) を並列実行することを指示しています。ループの反復数は BLOCKS であり 8 が設定されています。

// clang-format off
#include <stdio.h>
#include <stdlib.h>
#include <time.h>

#include <math.h>
#include <omp.h>

#define P 16
#define BLOCKS 8
#define SIZE (BLOCKS * P * P * P)

#define MAX 100
#define scaled_rand() ((rand() % MAX) / (1.0 * MAX))

#define IDX2(i, j) (i * P + j)
#define IDX4(b, i, j, k) (b * P * P * P + i * P * P + j * P + k)

int main(int argc, char *argv[]) {
  double w[SIZE];            /* output */
  double u[SIZE], dx[P * P]; /* input */
  int b, i, j, k, l;         /* loop counters */
  double start, end;         /* timers */

  omp_set_default_device(0);

  /* 起動時のオーバーヘッドを計測しないよう、ダミーの target 領域を追加しています */
  #pragma omp target
  { ; }

  /* 乱数値で入力を初期化します */
  srand(0);
  for (int i = 0; i < SIZE; i++)
    u[i] = scaled_rand();

  for (int i = 0; i < P * P; i++)
    dx[i] = scaled_rand();

  /* デバイスへデータをマップ */
  #pragma omp target enter data map(to: u[0:SIZE], dx[0:P * P])

  start = omp_get_wtime();

  /* collapse 節なしでカーネルをオフロードします */
  #pragma omp target teams distribute parallel for \
    private(b, i, j, k, l)
  for (b = 0; b < BLOCKS; b++) {
    for (i = 0; i < P; i++) {
      for (j = 0; j < P; j++) {
        for (k = 0; k < P; k++) {
          double ur = 0.;
          double us = 0.;
          double ut = 0.;

          for (l = 0; l < P; l++) {
            ur += dx[IDX2(i, l)] * u[IDX4(b, l, j, k)];
            us += dx[IDX2(k, l)] * u[IDX4(b, i, l, k)];
            ut += dx[IDX2(j, l)] * u[IDX4(b, i, j, l)];
          }

          w[IDX4(b, i, j, k)] = ur * us * ut;
        }
      }
    }
  }

  end = omp_get_wtime();

  #pragma omp target exit data map(from: w[0:SIZE])

  /* 結果を出力します */
  printf("no-collapse-clause: w[0]=%lf time=%lf\n", w[0], end - start);

  return 0;
}

libomptarget.so のデバッグ情報 (環境変数 LIBOMPTARGET_DEBUG=1 の場合に実行時に出力される) には、ループ反復の ND-range パーティション化と collapse 節を使用した際の並列性の向上が示されています。出力の Lb と Ub は、パーティション化された各次元の並列ループの上限と下限を示します。

collapse がないと、LIBOMPTARGET_DEBUG=1 による出力は、45 行目 target 領域に対し次のような情報を示します。

Libomptarget --> Launching target execution __omp_offloading_802_b85fb2__Z4main_l45 with pointer 0x0000000000ff1b48 (index=1).
Libomptarget --> Manifesting used target pointers:
Target LEVEL0 RTL --> Executing a kernel 0x0000000000ff1b48...
Target LEVEL0 RTL --> Assumed kernel SIMD width is 32
Target LEVEL0 RTL --> Preferred group size is multiple of 64
Target LEVEL0 RTL --> Level 0: Lb = 0, Ub = 7, Stride = 1
Target LEVEL0 RTL --> Group sizes = {8, 1, 1}
Target LEVEL0 RTL --> Group counts = {1, 1, 1}

collapse 節がない場合、最も外側のループの上限は (BLOCKS) = 8 であるため、並列ループの反復回数は 8 になります。この場合、1 つの work-group が 8 つの work-item を持ちます (work-group の合計は 1 x 1 x 1 = 1 で、各 work-group のサイズは 8 x 1 x 1 = 8 work-item)。カーネルは SIMD32 でベクトル化されており、32 個の work-item が 1 つの sub-group にまとめられています。work-item が 8 つしかないため、すべての SIMD レーンがアクティブではない sub-group は 1 つしかないことになります。

parallel for ディレクティブに collapse 節を追加することで、並列性を高め GPU で実行する work-item の数を増やすことができます。最初に、次の例に示すように collapse(2) 節を追加してみます。

  /* collapse 節を指定してカーネルをオフロードします */
  #pragma omp target teams distribute parallel for collapse(2) \
    private(b, i, j, k, l)
  for (b = 0; b < BLOCKS; b++) {
    for (i = 0; i < P; i++) {
      for (j = 0; j < P; j++) {
        for (k = 0; k < P; k++) {
          double ur = 0.;
          double us = 0.;
          double ut = 0.;

          for (l = 0; l < P; l++) {
            ur += dx[IDX2(i, l)] * u[IDX4(b, l, j, k)];
            us += dx[IDX2(k, l)] * u[IDX4(b, i, l, k)];
            ut += dx[IDX2(j, l)] * u[IDX4(b, i, j, l)];
          }

          w[IDX4(b, i, j, k)] = ur * us * ut;
        }
      }
    }
  }

LIBOMPTARGET_DEBUG=1 による出力は、collapse(2) が指定された場合のパーティション化の情報を次のように示します。

Libomptarget --> Launching target execution __omp_offloading_802_b85fb3__Z4main_l45 with pointer 0x0000000001dffc98 (index=1).
Libomptarget --> Manifesting used target pointers:
Target LEVEL0 RTL --> Executing a kernel 0x0000000001dffc98...
Target LEVEL0 RTL --> Assumed kernel SIMD width is 16
Target LEVEL0 RTL --> Preferred group size is multiple of 32
Target LEVEL0 RTL --> Level 0: Lb = 0, Ub = 15, Stride = 1
Target LEVEL0 RTL --> Level 1: Lb = 0, Ub = 7, Stride = 1
Target LEVEL0 RTL --> Group sizes = {16, 1, 1}
Target LEVEL0 RTL --> Group counts = {1, 8, 1}

collapse(2) の場合、並列ループの反復回数は BLOCKS x P = 8 x 16 = 128 となります。この場合、最終的に work-group は 8 つになり、各 work-group には 16 個の work-item があります (work-group の合計は 1 x 8 x 1 = 8 で、各 work-group のサイズは 16 x 1 x 1 = 16 work-item です)。カーネルは SIMD16 でベクトル化されており、16 個の work-item が 1 つの sub-group にまとめられています。つまり、各 work-group は 1 つの sub-group を持つことになります。

一方、collapse(3) 節にすると、LIBOMPTARGET_DEBUG=1 の出力は次のようなパーティション化の情報を示します。

Libomptarget --> Launching target execution __omp_offloading_802_b85fb4__Z4main_l45 with pointer 0x0000000000a2b9b8 (index=1).
Libomptarget --> Manifesting used target pointers:
Target LEVEL0 RTL --> Executing a kernel 0x0000000000a2b9b8...
Target LEVEL0 RTL --> Assumed kernel SIMD width is 16
Target LEVEL0 RTL --> Preferred group size is multiple of 32
Target LEVEL0 RTL --> Level 0: Lb = 0, Ub = 15, Stride = 1
Target LEVEL0 RTL --> Level 1: Lb = 0, Ub = 15, Stride = 1
Target LEVEL0 RTL --> Level 2: Lb = 0, Ub = 7, Stride = 1
Target LEVEL0 RTL --> Group sizes = {16, 1, 1}
Target LEVEL0 RTL --> Group counts = {1, 16, 8}

collapse(3) の場合、並列ループの反復数は BLOCKS x P x P = 8 x 16 x 16 = 2048 となります。この場合、最終的に work-group は 128 になり、各 work-group には 16 個の work-item があります (work-group の合計は 1 x 16 x 8 = 128 で、各 work-group のサイズは 16 x 1 x 1 = 16 work-item です)。カーネルは SIMD16 でベクトル化されており、16 個の work-item が 1 つの sub-group にまとめられています。つまり、各 work-group は 1 つの sub-group を持つことになります。

collapse(3) 節の代わりに collapse(4) 節を使用すると、LIBOMPTARGET_DEBUG=1 の出力は次のようなパーティション化情報を示します。

Libomptarget --> Launching target execution __omp_offloading_802_b85fb5__Z4main_l45 with pointer 0x0000000000aeec98 (index=1).
Libomptarget --> Manifesting used target pointers:
Target LEVEL0 RTL --> Executing a kernel 0x0000000000aeec98...
Target LEVEL0 RTL --> Assumed kernel SIMD width is 16
Target LEVEL0 RTL --> Preferred group size is multiple of 32
Target LEVEL0 RTL --> Level 0: Lb = 0, Ub = 32767, Stride = 1
Target LEVEL0 RTL --> Group sizes = {32, 1, 1}
Target LEVEL0 RTL --> Group counts = {1024, 1, 1}

collapse(4) の場合、並列ループの反復回数は BLOCKS x P x P x P = 8 x 16 x 16 x 16= 32768 となります。この場合、コンパイラーとランタイムは 1 次元にパーティション化し、1024 の work-group を持ち、各 work-group は 32 個の work-item になります (work-group の合計は 1024 x 1 x 1 = 1024 で、各 work-group のサイズは 32 x 1 x 1 = 32 work-item です)。カーネルは SIMD16 でベクトル化されており、16 個の work-item が 1 つの sub-group にまとめられています。つまり、各 work-group は 2 つの sub-group を持つことになります。

collapse 節を使用すると、入れ子になったループの実行時間が大幅に短縮されます。ここで使用する ATS✝ GPU (1 タイルのみ) で実行した各バージョンのパフォーマンスは次のようになります。

no collapse version : 0.028665 seconds
collapse(2) version : 0.003309 seconds
collapse(3) version : 0.002016 seconds
collapse(4) version : 0.002016 seconds

上記の結果から、collapse(3) 節または collapse(4) 節を追加すると、パフォーマンスがおよそ 14 倍向上したことがわかります (0.002016 秒 vs 0.028665 秒)。

GPU 上では collapse 節で実際のループを融合できないかもしれませんが、この節はコンパイラーとランタイムに入れ子になったループの並列性の度合いを伝え、ND-range のパーティション化に利用されることに留意してください。

ベクトルロードとストアの利点を活用するため、入れ子になったループの最も内側のループは、ベクトル化できるように collapse 節で融合しないことを推奨します。最も内側のループがユニットストライドで、反復数が SIMD 幅と同じになるよう十分に大きければ、最良のパフォーマンスを得ることができます。

✝開発コード名

製品および性能に関する情報

1 性能は、使用状況、構成、その他の要因によって異なります。詳細については、http://www.intel.com/PerformanceIndex/ を参照してください。

タイトルとURLをコピーしました