5.1.3. Loop Carried依存関係の簡素化

Intel® FPGA SDK for OpenCL™: ベスト・プラクティス・ガイド

ダウンロード PDF

ID 683521

日付 12/08/2017

バージョン 17.1

Public

5.1.3. Loop Carried依存関係の簡素化

カーネルでループに依存する依存関係を削除したり緩めることができない場合、依存関係を単純化して単一のwork-itemカーネルのパフォーマンスを向上させることができます。

次の式を検討してみましょう。


 1 #define N 128
 2 #define NUM_CH 3
 3 
 4 channel uchar CH_DATA_IN[NUM_CH];
 5 channel uchar CH_DATA_OUT;
 6 
 7 __kernel void unoptimized()
 8 {
 9   unsigned storage = 0;
10   unsigned num_bytes = 0;
11 
12   for (unsigned i = 0; i < N; i++) {
13 
14     #pragma unroll
15     for (unsigned j = 0; j < NUM_CH; j++) {
16       if (num_bytes < NUM_CH) {
17         bool valid = false;
18         uchar data_in = read_channel_nb_intel(CH_DATA_IN[j], &valid);
19         if (valid) {
20           storage <<= 8;
21           storage |= data_in;
22           num_bytes++;
23         }
24       }
25     }
26 
27     if (num_bytes >= 1) {
28       num_bytes -= 1;
29       uchar data_out = storage >> (num_bytes*8);
30       write_channel_intel(CH_DATA_OUT, data_out);
31     }
32   }
33 }

このカーネルは、3つの入力チャンネルから1バイトのデータをノンブロッキング形式で読み込みます。次に、データを一度に1バイトずつ出力チャネルに書き込みます。変数storageを使用して最大4バイトのデータを格納し、変数num_bytesを使用して、格納されているバイト数を追跡します。 storageに使用可能なスペースがある場合、カーネルはチャネルの1つから1バイトのデータを読み出し、storageの最下位バイトに格納します。

次の最適化レポートは、変数num_bytesにループで運ばれる依存関係があることを示しています。

===================================================================================
Kernel: unoptimized
===================================================================================
The kernel is compiled for single work-item execution.

Loop Report:

 + Loop "Block1" (file unoptimized3.cl line 12)
 | Pipelined with successive iterations launched every 7 cycles due to:
 |
 |     Data dependency on variable num_bytes  (file unoptimized3.cl line 10)
 |     Largest Critical Path Contributors:
 |         16%: Integer Compare Operation  (file unoptimized3.cl line 16)
 |         16%: Integer Compare Operation  (file unoptimized3.cl line 16)
 |         16%: Integer Compare Operation  (file unoptimized3.cl line 16)
 |          7%: Integer Compare Operation  (file unoptimized3.cl line 27)
 |          6%: Add Operation  (file unoptimized3.cl line 10, line 22, line 28)
 |          6%: Add Operation  (file unoptimized3.cl line 10, line 22, line 28)
 |          6%: Add Operation  (file unoptimized3.cl line 10, line 22, line 28)
 |          3%: Non-Blocking Channel Read Operation  (file unoptimized3.cl line 18)
 |          3%: Non-Blocking Channel Read Operation  (file unoptimized3.cl line 18)
 |          3%: Non-Blocking Channel Read Operation  (file unoptimized3.cl line 18)
 |
 |
 |-+ Fully unrolled loop (file unoptimized3.cl line 15)
     Loop was fully unrolled due to "#pragma unroll" annotation.

num_bytesの計算パスは次のとおりです。

16行目の比較（ （num_bytes <NUM_CH）の場合 ）
19行目の比較のために、18行目の非ブロッキング・チャネル読み出し操作でvalid変数の計算（ uchar data_in = read_channel_nb_intel（CH_DATA_IN [j]、＆valid） ）。
22行目への追加（ num_bytes ++ ）。
27行目の比較（ if（num_bytes> = 1） ）。
28行目の減算（ num_bytes - = 1 ）。

14行目のunrollプラグマのために、 Intel® FPGA SDK for OpenCL™オフライン・コンパイラーは、ループをアンロールし、ループ本体の比較と追加を3回繰り返します。最適化レポートは、比較がnum_bytesの計算パスで最も高価な演算であり、その後に22行目の加算が続くことを示しています。

num_bytesに対するループキャリーの依存関係を単純化するには、アプリケーションを再構築して次のタスクを実行することを検討してください。

カーネルは、 storageに使用可能な十分なスペースがある場合にのみチャネルから読み出すことを確認し、すべてのチャネルの操作がデータを返す読み出すた場合に（つまり、 storage内の空きスペースの少なくとも3つのバイトがある）。
この条件を設定すると、比較回数を減らすことで変数num_bytesの計算パスが簡単になります。
より簡単に3バイトのスペースしきい値を満たすために、 storageのサイズを4バイトから8バイトに増やします。

Below is the restructured kernel optimized:


 1 #define N 128
 2 #define NUM_CH 3
 3 
 4 channel uchar CH_DATA_IN[NUM_CH];
 5 channel uchar CH_DATA_OUT;
 6 
 7 __kernel void optimized()
 8 {
 9   // Change storage to 64 bits
10   ulong storage = 0;
11   unsigned num_bytes = 0;
12 
13   for (unsigned i = 0; i < N; i++) {
14 
15     // Ensure that we have enough space if we read from ALL channels
16     if (num_bytes <= (8-NUM_CH)) {
17       #pragma unroll
18       for (unsigned j = 0; j < NUM_CH; j++) {
19         bool valid = false;
20         uchar data_in = read_channel_nb_intel(CH_DATA_IN[j], &valid);
21         if (valid) {
22           storage <<= 8;
23           storage |= data_in;
24           num_bytes++;
25         }
26       }
27     }
28 
29     if (num_bytes >= 1) {
30       num_bytes -= 1;
31       uchar data_out = storage >> (num_bytes*8);
32       write_channel_intel(CH_DATA_OUT, data_out);
33     }
34   }
35 }

An optimization report similar to the one below indicates the successful simplification of the loop-carried dependency on the variable num_bytes:

===================================================================================
Kernel: optimized
===================================================================================
The kernel is compiled for single work-item execution.

Loop Report:

 + Loop "Block1" (file optimized3.cl line 13)
 | Pipelined well. Successive iterations are launched every cycle.
 |
 |
 |-+ Fully unrolled loop (file optimized3.cl line 18)
     Loop was fully unrolled due to "#pragma unroll" annotation.

言語の選択

Intel.com サーチを使用

クイックリンク

最近の検索

高度検索

検索のみ

Intel® FPGA SDK for OpenCL™: ベスト・プラクティス・ガイド

5.1.3. Loop Carried依存関係の簡素化