インテル® グラフィックス・テクノロジーでのオフロードの開始

ターゲットへオフロードする例

unsigned parArrayRHist[256][256],
     parArrayGHist[256][256], parArrayBHist[256][256];

#pragma offload target(gfx) if (do_offload) \
     pin(inputImage: length(imageSize)) \
     out(parArrayRHist, parArrayGHist, parArrayBHist)

     __Cilk_for (int ichunk = 0; ichunk < chunkCount; ichunk++){
          …
     }

上記の例では、CPU コードが生成され、ランタイムは次の処理を行います。

• ターゲットがシステムで利用できるかどうかを判断します。

• ターゲットが利用できない場合、または do_offload の評価に失敗した場合、for ループを CPU で実行します。

• それ以外の場合は、次の処理を行います。

  • ポインター inputImage で参照される imageSize * sizeof(inputImage[0]) バイトを固定して、ターゲットメモリーとの間でデータをコピーせずに、そのメモリーをターゲットと共有します。

  • parArrayRHist、parArrayGHist、および parArrayBHist のターゲットメモリー領域を作成します。

  • for ループの反復空間を N 個のチャンクに分割します (N は chunkCount 以下)。N の特定の値の選択はオフロードランタイムによって行われます。ドキュメントで説明されているように、境界やストライドのような反復空間構成や、環境変数で制御できる最大値のような要因に依存します。

  • それぞれ独自の反復空間チャンクを割り当てた N 個のターゲットスレッドでタスクを作成します。

  • ターゲットで実行するタスクをキューに入れます。

  • ターゲットのタスクの実行が完了するのを待ちます。

  • ターゲットメモリーから CPU メモリーに parArrayRHist、parArrayGHist、および parArrayBHist をコピーして、結果が直ちにすべての CPU スレッドで認識されるようにします。

完全な入れ子構造の _Cilk_for ループを使用してオフロードする例

float (* A)[k] = (float (*)[k])matA;
float (* B)[n] = (float (*)[n])matB;
float (* C)[n] = (float (*)[n])matC;

#pragma offload target(gfx) if (do_offload) \
     pin(A: length(m*k)), pin(B: length(k*n)), pin(C: length(m*n))

     __Cilk_for (int r = 0; r < m; r += TILE_m) {
          __Cilk_for (int c = 0; c < n; c += TILE_n) {
               …
          }
     }

上の例では、次のことが言えます。

• 完全な入れ子構造の __Cilk_for ループを使用することで、コンパイラーが入れ子構造のループを結合できるようになります。オフロードされる入れ子構造のループの反復空間は r ループと c ループの両方を含む 2 次元で、各ターゲットスレッドに並列実行用の 2 次元の反復空間チャンクが割り当てられます。

• A、B および C は配列へのポインターとして定義されますが、length はポインターによって参照された float 型配列の要素で指定されます。

ターゲットへ非同期オフロードする例

この例は、in1 配列の初期化をプロセッサー・グラフィックスへオフロードし、並行してホストで in2 を初期化します。そして、in1 と in2 を使用してホストで out の計算を実行します。

in1 の初期化が完了するのを待機するため、out の計算の前に omp taskwait プラグマを指定しています。

OpenMP* 構文を利用してプロセッサー・グラフィックスへオフロードするため、コンパイル時に、/Qopenmp /Qopenmp-offload=gfx (Windows*) または -qopenmp -qopenmp-offload=gfx (Linux*) オプションを指定する必要があります。

    int* in1 = (int*)malloc(SIZE * sizeof(int));
    int* in2 = (int*)malloc(SIZE * sizeof(int));
    int* out = (int*)malloc(SIZE * sizeof(int));

#pragma omp target map(tofrom: in1[0:SIZE]) nowait
#pragma omp parallel for
    for (int i = 0; i < SIZE; i++) {
        in1[i] = 1;
    }

#pragma omp parallel for
    for (int i = 0; i < SIZE; i++) {
        in2[i] = i;
    }

#pragma omp taskwait

#pragma omp parallel for
    for (int i = 0; i < SIZE; i++) {
        out[i] = in1[i] + in2[i];
    }

この例は、上記の OpenMP* 構文とほぼ同じですが、#pragma offload 構文を使用して記述されています。in1 の初期化が完了するのを待機するため、out の計算の前に offload_wait プラグマを指定しています。

    int* in1 = (int*)malloc(SIZE * sizeof(int));
    int* in2 = (int*)malloc(SIZE * sizeof(int));
    int* out = (int*)malloc(SIZE * sizeof(int));

#pragma offload target(gfx) pin(in1: length(SIZE)) signal(in1)
    _Cilk_for (int i = 0; i < SIZE; i++) {
        in1[i] = SIZE - i;
    }

    _Cilk_for (int i = 0; i < SIZE; i++) {
        in2[i] = i;
    }

#pragma offload_wait target(gfx) wait(in1)

    _Cilk_for (int i = 0; i < SIZE; i++) {
        out[i] = in1[i] + in2[i];
    }