6.1.1. Removing Loop-Carried Dependency

Intel® FPGA SDK for OpenCL™ Pro Edition: Best Practices Guide

Download PDF

ID 683521

Date 3/28/2022

Version

Public

A newer version of this document is available. Customers should click here to go to the newest version.

6.1.1. Removing Loop-Carried Dependency

Based on the feedback from the optimization report, you can remove a loop-carried dependency by implementing a simpler memory access pattern.

Consider the following kernel:

 1 #define N 128
 2 
 3 __kernel void unoptimized (__global int * restrict A,
 4                            __global int * restrict B,
 5                            __global int* restrict result)
 6 {
 7   int sum = 0;
 8 
 9   for (unsigned i = 0; i < N; i++) {
10     for (unsigned j = 0; j < N; j++) {
11       sum += A[i*N+j];
12     }
13     sum += B[i];
14   }
15 
16   * result = sum;
17 }

The optimization report for kernel unoptimized resembles the following:

The first row of the report indicates that the Intel® FPGA SDK for OpenCL™ Offline Compiler successfully infers pipelined execution for the outer loop, and a new loop iteration launches every other cycle.
The message due to Pipeline structure indicates that the offline compiler creates a pipeline structure that causes an outer loop iteration to launch every two cycles. The behavior is not a result of how you structure your kernel code.
Note: For recommendations on how to structure your single work-item kernel, refer to the Good Design Practices for Single Work-Item Kernel section.
The remaining messages in the first row of report indicate that the loop executes a single iteration at a time across the subloop because of data dependency on the variable sum. This data dependency exists because each outer loop iteration requires the value of sum from the previous iteration to return before the inner loop can start executing.
The second row of the report notifies you that the inner loop executes in a pipelined fashion with no performance-limiting loop-carried dependencies.

To optimize the performance of this kernel, remove the data dependency on variable sum so that the outer loop iterations do not execute serially across the subloop. Perform the following tasks to decouple the computations involving sum in the two loops:

Define a local variable (for example, sum2) for use in the inner loop only.
Use the local variable from Step 1 to store the cumulative values of A[i*N + j] as the inner loop iterates.
In the outer loop, store the variable sum to store the cumulative values of B[i] and the value stored in the local variable.

Below is the restructured kernel optimized:

 1 #define N 128
 2 
 3 __kernel void optimized (__global int * restrict A,
 4                          __global int * restrict B,
 5                          __global int * restrict result)
 6 {
 7   int sum = 0;
 8 
 9   for (unsigned i = 0; i < N; i++) {
10     // Step 1: Definition
11     int sum2 = 0;
12 
13     // Step 2: Accumulation of array A values for one outer loop iteration
14     for (unsigned j = 0; j < N; j++) {
15       sum2 += A[i*N+j];
16     }
17 
18     // Step 3: Addition of array B value for an outer loop iteration
19     sum += sum2;
20     sum += B[i];
21   }
22 
23   * result = sum;
24 }

An optimization report similar to the one below indicates the successful removal of the loop-carried dependency on the variable sum:

You have addressed all the loop-carried dependence issues successfully when you see only the following messages in the optimization report:

Pipelined execution inferred for innermost loops.
Pipelined execution inferred. Successive iterations launched every 2 cycles due to: Pipeline structure for all other loops.

Select Your Language

Using Intel.com Search

Quick Links

Recent Searches

Advanced Search

Only search in

Intel® FPGA SDK for OpenCL™ Pro Edition: Best Practices Guide

6.1.1. Removing Loop-Carried Dependency