インテルのみ表示可能 — GUID: ewa1398275007271
Ixiasoft
5.1.1. ループ実行依存関係の削除
最適化レポートからのフィードバックに基づいて、より単純なメモリー・アクセス・パターンを実装することにより、ループで運ばれる依存関係を削除できます。
以下の項目について検討します。
1 #define N 128 2 3 __kernel void unoptimized (__global int * restrict A, 4 __global int * restrict B, 5 __global int* restrict result) 6 { 7 int sum = 0; 8 9 for (unsigned i = 0; i < N; i++) { 10 for (unsigned j = 0; j < N; j++) { 11 sum += A[i*N+j]; 12 } 13 sum += B[i]; 14 } 15 16 * result = sum; 17 }
unoptimizedカーネルの最適化レポートは、次のようになります。
===================================================================================
Kernel: unoptimized
===================================================================================
The kernel is compiled for single work-item execution.
Loop Report:
+ Loop "Block1" (file k.cl line 9)
| Pipelined with successive iterations launched every 2 cycles due to:
|
| Pipeline structure: every terminating loop with subloops has iterations
| launched at least 2 cycles apart.
| Having successive iterations launched every two cycles should still lead to
| good performance if the inner loop is pipelined well and has sufficiently
| high number of iterations.
|
| Iterations executed serially across the region listed below.
| Only a single loop iteration will execute inside the listed region.
| This will cause performance degradation unless the region is pipelined well
| (can process an iteration every cycle).
|
| Loop "Block2" (file k.cl line 10)
| due to:
| Data dependency on variable sum (file k.cl line 7)
|
|
|-+ Loop "Block2" (file k.cl line 10)
Pipelined well. Successive iterations are launched every cycle.
- レポートの最初のロウは、 Intel® FPGA SDK for OpenCL™オフライン・コンパイラーは外部ループのパイプライン実行を正常に推定し、新しいループ反復が他のサイクルごとに起動します。
- due to Pipeline structureというメッセージは、オフライン・コンパイラーが外側ループ反復を2サイクルごとに起動させるパイプライン構造を作成することを示します。この動作は、カーネルコードの構造の結果ではありません。
注: 単一のWork-Itemカーネルを構成する方法の推奨事項については、単一のWork-Itemカーネルのための良いデザイン方法のセクションを参照してください。
- レポートの最初の行の残りのメッセージは、変数sumに対するデータ依存のために、ループがサブループ全体で一度に1回の繰り返しを実行することを示しています。このデータ依存関係は、各外部ループ反復が、内部ループが実行を開始する前に前の反復からのsumの値が必要であるために存在します。
- レポートの2番目のロウは、内部ループがパイプライン形式で実行され、パフォーマンス制限のあるループに依存する依存関係がないことを通知します。
外側のループの反復がサブループ渡って連続的に実行されないように、このカーネルのパフォーマンスを最適化するには、変数sumのデータの依存関係を削除します。 2つのループでsumを含む計算を切り離すには、次のタスクを実行します。
- 内部ループでのみ使用するローカル変数（ sum2など ）を定義します。
- ステップ1のローカル変数を使用して、 A [i * N + j]の累積値を内部ループの繰り返しとして格納します。
- 外部ループでは、 B [i]の累積値とローカル変数に格納されている値を格納する変数sumを格納します。
以下は、再構成されたカーネルoptimizedです。
1 #define N 128 2 3 __kernel void optimized (__global int * restrict A, 4 __global int * restrict B, 5 __global int * restrict result) 6 { 7 int sum = 0; 8 9 for (unsigned i = 0; i < N; i++) { 10 // Step 1: Definition 11 int sum2 = 0; 12 13 // Step 2: Accumulation of array A values for one outer loop iteration 14 for (unsigned j = 0; j < N; j++) { 15 sum2 += A[i*N+j]; 16 } 17 18 // Step 3: Addition of array B value for an outer loop iteration 19 sum += sum2; 20 sum += B[i]; 21 } 22 23 * result = sum; 24 }
以下のような最適化レポートは、変数sumに対するループキャリー依存関係の削除に成功したことを示しています。
===================================================================================
Kernel: optimized
===================================================================================
The kernel is compiled for single work-item execution.
Loop Report:
+ Loop "Block1" (file optimized.cl line 9)
| Pipelined with successive iterations launched every 2 cycles due to:
|
| Pipeline structure: every terminating loop with subloops has iterations
| launched at least 2 cycles apart.
| Having successive iterations launched every two cycles should still lead to
| good performance if the inner loop is pipelined well and has sufficiently
| high number of iterations.
|
|
|-+ Loop "Block2" (file optimized.cl line 14)
Pipelined well. Successive iterations are launched every cycle.
===================================================================================
最適化レポートに次のメッセージだけが表示されたら、ループキャリー依存関係の問題はすべて解決しました。
- Pipelined execution inferred∸最も内側のループ用。
- Pipelined execution inferred. Successive iterations launched every 2 cycles due to: Pipeline structure∸他のすべてのループ用。