5.1.5. シフトレジスターの推測によるループキャリー依存関係の削除

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

ダウンロード PDF

ID 683521

日付 12/08/2017

バージョン 17.1

Public

5.1.5. シフトレジスターの推測によるループキャリー依存関係の削除

Intel® FPGA SDK for OpenCL™オフライン・コンパイラーで倍精度浮動小数点演算を効率的に実行する1つの作業項目カーネルを処理できるようにするには、シフトレジスターを推論してループ実行依存関係を削除します。

以下の項目について検討します。


 1 __kernel void double_add_1 (__global double *arr,
 2                             int N,
 3                             __global double *result)
 4 {
 5   double temp_sum = 0;
 6
 7   for (int i = 0; i < N; ++i)
 8   {
 9       temp_sum += arr[i];
10   }
11 
12   *result = temp_sum;
13 }

unoptimizedカーネルのOptimizationレポートは、次のようになります。

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

Loop Report:

 + Loop "Block1" (file unoptimized5.cl line 7)
   Pipelined with successive iterations launched every 11 cycles due to:

       Data dependency on variable temp_sum  (file unoptimized5.cl line 9)
       Largest Critical Path Contributor:
           97%: Fadd Operation  (file unoptimized5.cl line 9)

最適化されていないカーネルは、倍精度浮動小数点配列arr [i]の要素を合計する累算器です。各ループ反復に対して、オフライン・コンパイラーは加算の結果を計算するために11サイクルを要し、それを変数temp_sumに格納します。各ループの反復では、以前のループ反復からのtemp_sumの値が必要です。これにより、 temp_sumにデータ依存関係が作成されます。

データの依存関係を削除するには、配列arr [i]をシフトレジスターとして推定します。

以下は、再構成されたカーネルoptimizedです：


 1 //Shift register size must be statically determinable
 2 #define II_CYCLES 12
 3
 4 __kernel void double_add_2 (__global double *arr,
 5                             int N,
 6                             __global double *result)
 7 {
 8     //Create shift register with II_CYCLE+1 elements
 9     double shift_reg[II_CYCLES+1];
10	
11     //Initialize all elements of the register to 0
12     for (int i = 0; i < II_CYCLES + 1; i++)
13     {
14         shift_reg[i] = 0;
15     }
16    
17     //Iterate through every element of input array
18     for(int i = 0; i < N; ++i)
19     {
20         //Load ith element into end of shift register
21         //if N > II_CYCLE, add to shift_reg[0] to preserve values
22         shift_reg[II_CYCLES] = shift_reg[0] + arr[i];
23 
24         #pragma unroll
25         //Shift every element of shift register
26         for(int j = 0; j < II_CYCLES; ++j)
27         {
28             shift_reg[j] = shift_reg[j + 1];
29         }
30     }
31 
32     //Sum every element of shift register
33     double temp_sum = 0;
34 	
35     #pragma unroll 
36     for(int i = 0; i < II_CYCLES; ++i)
37     {
38         temp_sum += shift_reg[i];
39     }
40
41     *result = temp_sum;
42 }

次のOptimizationレポートは、シフトレジスターshift_reg [II_CYCLES]の推論が変数temp_sumのデータ依存性を正常に削除することを示しています。

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

Loop Report:

 + Fully unrolled loop (file optimized5.cl line 12)
   Loop was automatically and fully unrolled.
   Add "#pragma unroll 1" to prevent automatic unrolling.


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


 + Fully unrolled loop (file optimized5.cl line 36)
   Loop was fully unrolled due to "#pragma unroll" annotation.

言語の選択

Intel.com サーチを使用

クイックリンク

最近の検索

高度検索

検索のみ

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

5.1.5. シフトレジスターの推測によるループキャリー依存関係の削除