6.1.2. 松弛循环携带的依赖性

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

下载 PDF

ID 683521

日期 9/26/2022

版本 22.3

Public

6.1.2. 松弛循环携带的依赖性

基于优化报告的反馈，您可以通过增加依赖项的距离来松弛循环携带的依赖性。通过增加循环携带的值的生成及其使用之间发生的循环迭代的次数，来增加依赖项距离。

请参考如下代码实例：


1 #define N 128
 2 
 3 __kernel void unoptimized (__global float * restrict A,
 4                            __global float * restrict result)
 5 {
 6   float mul = 1.0f;
 7 
 8   for (unsigned i = 0; i < N; i++)
 9     mul *= A[i];
10 
11   * result = mul;
12 }

上述优化报告显式 Intel® FPGA SDK for OpenCL™ Offline Compiler成功推断了循环的流水线执行。但是，变量mul上循环携带的依赖项导致循环迭代每6个周期启动一次。该情况下，第9行上的浮点乘法操作（即，mul *= A[i]）对变量mul的计算贡献最大的延迟。

要松弛循环携带的数据依赖性，除了使用单个变量来存储乘法结果，还可以在变量的M副本上进行操作并每M次迭代使用一个副本：

声明变量mul的多个副本（例如，在名为mul_copies的数组中）。
初始化mul_copies的所有副本。
在乘法运算中使用数组中的最后一个副本。
执行移位操作以将数组的最后一个值传回移位寄存器的开头。
减少mul的副本，并将最终值写入result。

以下是重构的内核：


 1 #define N 128
 2 #define M 8
 3 
 4 __kernel void optimized (__global float * restrict A,
 5                          __global float * restrict result)
 6 {
 7   float mul = 1.0f;
 8 
 9   // Step 1: Declare multiple copies of variable mul
10   float mul_copies[M];
11 
12   // Step 2: Initialize all copies
13   for (unsigned i = 0; i < M; i++)
14     mul_copies[i] = 1.0f;
15 
16   for (unsigned i = 0; i < N; i++) {
17     // Step 3: Perform multiplication on the last copy
18     float cur = mul_copies[M-1] * A[i];
19 
20     // Step 4a: Shift copies
21     #pragma unroll 
22     for (unsigned j = M-1; j > 0; j--)
23       mul_copies[j] = mul_copies[j-1];
24 
25     // Step 4b: Insert updated copy at the beginning
26     mul_copies[0] = cur;
27   }
28 
29   // Step 5: Perform reduction on copies
30   #pragma unroll 
31   for (unsigned i = 0; i < M; i++)
32     mul *= mul_copies[i];
33 
34   * result = mul;
35 }

类似于以下内容的优化报告指示成功松弛了变量mul上的循环携带依赖项：

选择您的语言

使用 Intel.com 搜索

快速链接

最近搜索

高级搜索

仅搜索

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

6.1.2. 松弛循环携带的依赖性