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.