Developer Guide

FPGA Optimization Guide for Intel® oneAPI Toolkits

ID 767853
Date 7/13/2023
Public

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

Document Table of Contents

FPGA Loop Directives

The following table summarizes loop directives:

FPGA Loop Directives

Directive (Pragma, Attribute, or Function)

Description

Example

disable_loop_pipelining

Directs the Intel® oneAPI DPC++/C++ Compiler to disable pipelining of a loop.

[[intel::disable_loop_pipelining]]
for (int i = 1; i < N; i++) {
  int j = a[i-1];
  // Memory dependency induces a high-latency loop feedback path
  a[i] = foo(j)
}
initiation_interval

Forces a loop to have a loop initialization interval (II) of a specified value.

// ii set to 5
[[intel::initiation_interval(5)]]
  for (int i = 0; i < N; ++i){
}
ivdep

Ignores memory dependencies between iterations of this loop

// ivdep loop
[[intel::ivdep]] for (…) {}
//ivdep safelen 
[[intel::ivdep(safelen)]] for (;;) {} 
// ivdep accessor
[[intel::ivdep(accessorA)]] for (;;) {}
//ivdep array safelen
[[intel::ivdep(accessorA, safelen)]]
for (;;){}
loop_coalesce

Coalesces nested loops into a single loop without affecting the loop functionality.

[[intel::loop_coalesce(2)]] 
for (int i = 0; i < N; i++)
  for (int j = 0; j < M; j++)
    sum[i][j] += i+j;
max_concurrency

Limits the number of iterations of a loop that can simultaneously execute at any time.

//max concurrency set to 1
[[intel::max_concurrency(1)]] 
  for (int i = 0; i < c; ++i){
}
max_interleaving

Maximizes the throughput and hardware resource occupancy of pipelined inner loops in a loop nest.

// Loop j is pipelined with ii=1
for (int j = 0; j < M; j++) {
  int a[N];
  // Loop i is pipelined with ii=2 
  [[intel::max_interleaving(1)]]
  for (int i = 1; i < N; i++) {
    a[i] = foo(i)
  }
  …
}
speculated_iterations

Improves the performance of pipelined loops.

[[intel::speculated_iterations(1)]]
  while (m*m*m < N) {
    m += 1;
  }
  dst[0] = m;
unroll

Unrolls a loop in the kernel code.

// unroll factor N set to 2
#pragma unroll 2
for(size_t k = 0; k < 4; k++){
  mac += data_in[(gid * 4) + k] * coeff[k];
}
nofusion Prevents the compiler from fusing the annotated loop with any of the adjacent loops.
for (int x = 0; x < N; x++) { 
  a1_acc[x] = x; 
}

[[intel::nofusion]] 
for (int x = 0; x < N; x++) { 
  a2_acc[x] = x; 
}
sycl::ext::intel::fpga_loop_fuse<v>(f) Fuses loops within the function f up to a depth of v >= 1, where v = 1 by default.
[=]() [[intel::kernel_args_restrict]] { 
  sycl::ext::intel::fpga_loop_fuse<v>{
    for (int x = 0; x < N; x++) {
      for (int y = 0; y < N; y++) {
        for (int z = 0; z < N; z++) {
          a1_acc[x][y][z] = 0;
        }
      }
    }
    for (int x = 0; x < N + 1; x++) {
      for (int y = 0; y < N + 1; y++) {
        for (int z = 0; z < N + 1; z++) {
          a2_acc[x][y][z] = 0;
        }
      }
    }
  }
}
sycl::ext::intel::fpga_loop_fuse<v><v>(f) Fuses loops within the function f up to a depth v >= 1 while overriding fusion-safety checks. Here, v = 1 by default.
[=]() { //Kernel
  sycl::ext::intel::fpga_loop_fuse_independent([&] {
    for(int x = 0; x < N; x++){
      a3_acc[x] = x;
    }
    for(int x = 0; x < N + 1; x++){
      a4_acc[x] = x;
    }
  });
}