インテルのみ表示可能 — GUID: mwh1391807501939
Ixiasoft
3.2. ループのアンロール
Intel® FPGA SDK for OpenCL™オフライン・コンパイラー OpenCL™カーネル記述をハードウェア・リソースに変換します。 OpenCLカーネルにループ反復が含まれている場合、ループを展開してパフォーマンスを向上させてください。 ループ展開は、ハードウェア・リソース消費の増加を犠牲にしてオフライン・コンパイラーが実行する反復回数を減少させます。
各Work-Itemが配列内の4つの要素の累積を計算する責任がある並列アプリケーション用のOpenCLコードを検討してください。
__kernel void example ( __global const int * restrict x,
__global int * restrict sum ) {
int accum = 0;
for (size_t i = 0; i < 4; i++) {
accum += x[i + get_global_id(0) * 4];
}
sum[get_global_id(0)] = accum;
}このカーネルでは、以下の3つの主要な動作が行われています。
- 入力xからのロード動作
- 累積
- オペレーションを出力sumに格納する
オフライン・コンパイラーは、OpenCLカーネルコードのデータフローセマンティクスに従って、これらの動作をパイプラインに配置します。たとえば、オフライン・コンパイラーは、ループ終了条件に応じて、パイプラインの最後からパイプラインの先頭に結果を転送することによってループを実装します。
OpenCLカーネルは、クロックサイクルごとに各Work-Itemの1回のループ反復を実行します。十分なハードウェア・リソースがあれば、ループを展開することでカーネルのパフォーマンスを向上させることができ、カーネルが実行する反復回数が減ります。ループをアンロールするには、次のコード例に示すように、 #pragma unrollディレクティブをメインループに追加します。ループをアンロールすると、オフライン・コンパイラーが作成するコンピューティング・ユニットの構造が大幅に変更されることに注意してください。
__kernel void example ( __global const int * restrict x,
__global int * restrict sum ) {
int accum = 0;
#pragma unroll
for (size_t i = 0; i < 4; i++) {
accum += x[i + get_global_id(0) * 4];
}
sum[get_global_id(0)] = accum;
}この例では、 #pragma unrollディレクティブにより、オフライン・コンパイラーはループの4つの繰り返しを完全に展開します。アンロールを達成するために、オフライン・コンパイラーは、加算演算の数を3倍にし、4倍のデータをロードすることによってパイプラインを拡張します。ループが除去されると、コンピューティング・ユニットはフィードフォワード構造を前提します。結果として、コンピューティング・ユニットは、初期ロード動作および加算の完了後にクロックサイクルごとにsum要素を記憶することができる。オフライン・コンパイラーは、4つのロード動作を統合することによってこのカーネルをさらに最適化し、コンピューティング・ユニットが1つのロード動作で結果を計算するために必要なすべての入力データをロードできるようにします。
注意:
ネストされたループ構造は使用しないでください。代わりに、可能であれば#pragma unrollディレクティブを追加して大きな単一ループを実装するか、内部ループを展開してください。
ループをアンロールし、グローバルメモリーからロード動作を統合することにより、ハードウェアによるカーネルのインプリメンテーションでクロックサイクルごとに多くの動作を実行できます。一般に、OpenCLカーネルのパフォーマンスを向上させるために使用する方法は、次の結果を達成する必要があります。
- 並列動作の数を増やす
- 実装のメモリー帯域幅を増やす
- カーネルがハードウェアで実行できるクロックサイクルあたりの動作数を増やす
オフライン・コンパイラーは、以下の状況で完全にループを展開することができない場合があります。
- 非常に多数の反復を伴うデータ依存ループの完全なアンローリングを指定します。したがって、カーネルのハードウェア実装がFPGAに適合しない可能性があります。
- 完全な展開を指定し、ループの境界は定数ではありません。
- ループは、複雑な制御フローで構成されます（たとえば、コンパイル時に不明な複雑な配列インデックスまたは終了条件を含むループ）。
上記の最後の2つのケースでは、オフライン・コンパイラーは次の警告を発行します。
ループのフルアンロールが要求されますが、ループの境界を特定できません。ループはアンロールされません。
これらの状況でループのアンロールをイネーブルするには、 #pragma unroll <N> ディレクティブ(<N>はアンロール係数である）を指定します。アンロール係数は、オフライン・コンパイラーがアンロールする回数を制限します。たとえば、カーネル内のループが展開されないようにするには、そのループに#pragma unroll 1ディレクティブを追加します。
適切に構成されたループを構築するためのヒントについては、Good Design Practices for Single Work-Item Kernelを参照してください。