Description
The XCL_LOOP_TRIPCOUNT attribute can be applied to a loop to manually specify the total number of iterations performed by the loop.
The Vivado High-Level Synthesis (HLS) reports the total latency of each loop, which is the number of clock cycles to execute all iterations of the loop. The loop latency is therefore a function of the number of loop iterations, or tripcount.
The tripcount can be a constant value. It can depend on the value of
variables used in the loop expression (for example, x<y), or depend on control statements used inside the loop. In some cases, the
HLS tool cannot determine the tripcount, and the latency is unknown. This includes cases in
which the variables used to determine the tripcount are:
- Input arguments, or
- Variables calculated by dynamic operation.
In cases where the loop latency is unknown or cannot be calculated, the XCL_LOOP_TRIPCOUNT attribute lets you specify minimum, maximum, and average iterations for a loop. This lets the tool analyze how the loop latency contributes to the total design latency in the reports, and helps you determine appropriate optimizations for the design.
Syntax
Place the attribute in the OpenCL source before the loop declaration.
__attribute__((xcl_loop_tripcount(<min>, <max>, <average>)))Where:
- <min>: Specifies the minimum number of loop iterations.
- <max>: Specifies the maximum number of loop iterations.
- <avg>: Specifies the average number of loop iterations.
Examples
In this example, the WHILE loop in function f is
specified to have a minimum tripcount of 2, a maximum tripcount of 64, and an average
tripcount of 33.
__kernel void f(__global int *a) {
unsigned i = 0;
__attribute__((xcl_loop_tripcount(2, 64, 33)))
while(i < 64) {
a[i] = i;
i++;
}
}