2.8.4. Nested Loops
Intel® FPGA SDK for OpenCL™オフライン・コンパイラーループ反復の順序付けのためにパイプライン実行を推測しません。その結果、内部ループの反復回数は、異なるループ反復で異なる可能性があるため、外部ループ反復は、その後の内部ループに対して順序が乱れる可能性があります。
アウトオブオーダーのアウターループ反復の問題を解決するには、アウターループ反復間で変化しない上限と下限を持つ内側ループをデザインします。
Single Work-Itemの実行
FPGA上でのハイスループットのWork-Itemベースのカーネル実行を確実にするために、 Intel® FPGA SDK for OpenCL™オフライン・コンパイラーはある時点で複数のパイプライン・ステージを並列に処理する必要があります。この並列性は、ループの反復をパイプライン化することによって実現されます。
Single Work-Itemでの累積を示す次の簡単なコード例を検討してください。
1 kernel void accum_swg (global int* a,
global int* c,
int size,
int k_size) {
2 int sum[1024];
3 for (int k = 0; k < k_size; ++k) {
4 for (int i = 0; i < size; ++i) {
5 int j = k * size + i;
6 sum[k] += a[j];
7 }
8 }
9 for (int k = 0; k < k_size; ++k) {
10 c[k] = sum[k];
11 }
12 }
次の図は、iの各反復がどのようにブロックに入るかを示しています。
外部ループを観測すると、IIの値が1であるということは、スレッドの各反復がすべてのクロックサイクルに入ることもできることを意味します。この例では、 k_sizeが20でsizeが4であると見なされます。これは、最初の8クロック・サイクルでは真であり、外側のループ反復0〜7は、それをストールさせることなく下流に入ることができるからです。スレッド0が内部ループに入ると、それは4回の反復で終了します。スレッド1〜8は内部ループに入ることができず、スレッド0によって4サイクル停止します。スレッド0の反復が完了すると、スレッド1は内部ループに入ります。その結果、スレッド9は、クロックサイクル13で外側ループに入る。スレッド9から20は、 sizeの値である4クロックサイクルごとにループに入る。この例では、外部ループの動的開始間隔が静的に予測される開始間隔1よりも大きく、内部ループのトリップ数の関数であることがわかります。
非線形実行
ループ構造は線形実行をサポートしていません。次のコード例は、外部ループiに2つの分岐する内部ループが含まれていることを示しています。外側ループの各反復は、1つの内側ループまたは非直線的実行パターンである他方を実行することができます。
__kernel void structure (__global unsigned* restrict output1,
__global unsigned* restrict output2,
int N) {
for (unsigned i = 0; i < N; i++) {
if ((i & 3) == 0) {
for (unsigned j = 0; j < N; j++) {
output1[i+j] = i * j;
}
}
else
{
for (unsigned j = 0; j < N; j++) {
output2[i+j] = i * j;
}
}
}
}アウトオブオーダーのループ反復
内側ループの反復回数は、外側ループの反復ごとに異なります。次のコード例を検討してください。
__kernel void order( __global unsigned* restrict input,
__global unsigned* restrict output
int N ) {
unsigned sum = 0;
for (unsigned i = 0; i < N; i++) {
for (unsigned j = 0; j < i; j++) {
sum += input[i+j];
}
}
output[0] = sum;
}この例は、 i = 0の場合、内側のループjがゼロ回反復することを示しています。 i = 1の場合、 jは 1回反復します。内部ループの反復回数が変わるため、オフライン・コンパイラーはパイプライン処理を推測できません。
シリアルリージョン
内部ループアクセスが外部ループ依存関係を引き起こすと、ネストされたループ内で直列エリアが発生することがあります。内部ループは、データまたはメモリーの依存性のために、外部ループの反復で直列エリアになります。
定常状態では、外側ループのII =内側ループのII *内側ループのトリップ数となります。 IIが1より大きい内部ループと、直列実行エリアがない外部ループの場合、スレッドを外部ループからインターリーブすることは可能です。
次の式を検討してみましょう。
kernel void serially_execute (global int * restrict A,
global int * restrict B,
global int * restrict result,
unsigned N) {
int sum = 0;
for (unsigned i = 0; i < N; i++) {
int res;
for (int j = 0; j < N; j++) {
sum += A[i*N+j];
}
sum += B[i];
}
*result = sum;
}
