6.1.4. 通过推断移位寄存器来松弛循环携带的依赖项

Intel® FPGA SDK for OpenCL™ Pro Edition: 最佳实践实践指南

下载 PDF

在收藏夹中：

Intel® FPGA Development Tools Support Intel® FPGA SDK for OpenCL™ Support Agilex™ 7 FPGAs and SoC FPGAs Support

ID 683521

日期 9/26/2022

版本 22.3

Public

6.1.4. 通过推断移位寄存器来松弛循环携带的依赖项

要启动 Intel® FPGA SDK for OpenCL™ Offline Compiler来处理有效执行双精度浮点运算的单个work-item内核，请通过推断移位寄存器来删除循环携带的依赖项。

请参考如下内核：


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的内核优化报告类似于以下内容：

未优化的内核是一个累加器，它对双精度浮点数组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 }

以下优化报告表明移位寄存器shift_reg[II_CYCLES]的推断成功删除了自变量temp_sum的数据依赖性：

选择您的语言

使用 Intel.com 搜索

快速链接

最近搜索

高级搜索

仅搜索

Intel® FPGA SDK for OpenCL™ Pro Edition: 最佳实践实践指南

6.1.4. 通过推断移位寄存器来松弛循环携带的依赖项