仅对英特尔可见 — GUID: ewa1425568230374
Ixiasoft
产品终止通知
1. Intel® FPGA SDK for OpenCL™ Pro Edition最佳实践指南介绍
2. 查看您Kernel的report.html文件
3. OpenCL内核设计概念
4. OpenCL内核设计最佳实践
5. 分析(Profiling)您的内核来识别性能瓶颈
6. 提高单个Work-Item内核性能的策略
7. 提高NDRange内核数据处理效率的策略
8. 提高存储器访问效率的策略
9. 优化FPGA面积使用的策略
10. 优化英特尔 Stratix 10 OpenCL设计的策略
11. 提高主机应用程序性能的策略
12. Intel® FPGA SDK for OpenCL™ Pro版最佳实践指南存档
A. Intel® FPGA SDK for OpenCL™ Pro版最佳实践指南修订历史
仅对英特尔可见 — GUID: ewa1425568230374
Ixiasoft
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的数据依赖性: