Before using the Intel^® FPGA SDK for OpenCL™ or the Intel^® FPGA Runtime Environment (RTE) for OpenCL to program your device, familiarize yourself with the respective getting started guides. This document assumes that you have performed the following tasks:

• For developing and deploying OpenCL kernels, download the tar file and run the installers to install the SDK, the Intel^® Quartus^® Prime Pro Edition software, and device support.
• For deployment of OpenCL kernels, download and install the RTE.
• If you want to use the SDK or the RTE to program an Intel^® SoC FPGA, you also have to download and install the Intel^® SoC FPGA Embedded Development Suite (EDS) Pro Edition.
• Install and set up your FPGA board.
• Verify that board installation is successful, and the board functions correctly.

If you have not performed the tasks described above, refer to the SDK's getting starting guides for more information.

The following SDK components work together to program an Intel^® FPGA:

• The host application and the host compiler
• The OpenCL kernel(s) and the offline compiler
• The Custom Platform

The Custom Platform provides the board support package. Typically, the board manufacturer develops the Custom Platform that supports a specific OpenCL board. The offline compiler targets the Custom Platform when compiling an OpenCL kernel to generate a hardware programming image. The host then runs the host application, which usually programs and executes the hardware image onto the FPGA.

In a sequential implementation of a program (for example, on a conventional processor), the program counter controls the sequence of instructions that are executed on the hardware, and the instructions that execute on the hardware across time. In a spatial implementation of a program, such as program implementation within the Intel^® FPGA SDK for OpenCL™ , instructions are executed as soon as the prerequisite data is available. Programs are interpreted as a series of connections representing the data dependencies.

An OpenCL kernel source file (.cl) contains your OpenCL kernel source code that runs on the FPGA. The offline compiler groups one or more kernels into a temporary file and then compiles this file to generate the following files and folders:

• A .aoco object file is an intermediate object file that contains information for later stages of the compilation. It is not presented unless the aoc option -save-temps is specified.
• A.aocx image file is the hardware configuration file and contains information necessary to program the FPGA at runtime.
• The work folder or subdirectory, which contains data necessary to create the .aocx file. By default, the name of the work directory is the name of your .cl file. If you compile multiple kernel source files, the name of the work directory is the name of the last .cl file you list in the aoc command line.

The .aocx file contains data that the host application uses to create program objects, a concept within the OpenCL runtime API, for the target FPGA. The host application first loads these program objects into memory. Then the host runtime uses these program objects to program the target FPGA, as required for kernel launch operations by the host program.

The following figure illustrates the OpenCL kernel design flow that has a single compilation step.

A successful compilation results in the following files and reports:

• A .aocr file
• A .aocx file
• In the <your_kernel_filename>/reports/report.html file, the estimated resource usage summary provides a preliminary assessment of area usage. If you have a single work-item kernel, the optimization report identifies performance bottlenecks.

The following figure outlines the stages in the SDK's design flow. The steps in the design flow serve as checkpoints for identifying functional errors and performance bottlenecks. They allow you to modify your OpenCL kernel code without performing a full compilation on each iteration. You have the option to perform some or all of the compilations steps.

The SDK's design flow includes the following steps:

1. Emulation

   Assess the functionality of your OpenCL kernel by executing it on one or multiple emulation devices on an x86-64 host system. For Linux systems, the Emulator offers symbolic debug support. Symbolic debug allows you to locate the origins of functional errors in your kernel code.

2. Intermediate Compilation

   There are two available intermediate compilation steps. You have the option to include one or both of these compilation steps in your design flow.

   • Compile one or more .cl kernel source files using the -c flag. Doing so instructs the offline compiler to generate .aoco object files that contain the output from the OpenCL parser.
   • Compile one or more .cl kernel source files or .aoco files, but not both, using the -rtl flag. Doing so instructs the offline compiler to perform the following tasks:
     • If the input files are .cl files, the offline compiler generates an intermediate .aoco file for each kernel source file and then links them to generate a .aocr file.
     • If the input files are .aoco files, the offline compiler links them to generate a .aocr file.
     • Creates a <your_kernel_filename> directory.
     The offline compiler uses the .aocr file to generate the final .aocx hardware configuration file.

     Note: If you compile your kernel(s) using the -c flag in an environment where the default board is X, and then you compile your .aoco files using the -rtl flag in an environment where the default board is Y, the offline compiler reads board X from the .aoco files and then pass it on to the subsequent compilation stages.

3. Review HTML Report

   Review the <your_kernel_filename>/reports/report.html file of your OpenCL application to determine whether the estimated kernel performance data is acceptable. The HTML report also provides suggestions on how you can modify your kernel to increase performance.

4. Simulation (Preview)

   Asses the functionality of your OpenCL kernel by running it through simulation. Simulation lets you asses the function correctness and dynamic performance of your kernel without a long compilation time. You can capture and view waveforms for your kernel to help you debug your kernel.

5. Fast Compilation

   Assess the functionality of your OpenCL kernel in hardware. The fast compilation step generates a .aocx file in a fraction of the time required to complete a full compilation. The Intel^® FPGA SDK for OpenCL™ Offline Compiler reduces compilation time by performing only light optimizations.

6. Incremental Compilation

   Assess the functionality of your OpenCL kernel in hardware. The incremental compilation step generates a .aocx file by compiling only the kernels you have modified. The Intel^® FPGA SDK for OpenCL™ Offline Compiler improves your productivity by scaling compilation times with the size of your design changes rather than the size of your overall design.

7. Profiling

   Instruct the Intel^® FPGA SDK for OpenCL™ Offline Compiler to insert performance counters in the FPGA programming image. During execution, the counters collect performance information which you can then review in the Intel^® FPGA Dynamic Profiler for OpenCL™ GUI.

8. Full deployment

   When you are satisfied with the performance of your OpenCL kernel throughout the design flow, perform a full compilation. The resulting .aocx file is suitable for deployment.

For more information on the HTML report and kernel profiling, refer to the Intel^® FPGA SDK for OpenCL™ Pro Edition Best Practices Guide.

Use the disable_loop_pipelining pragma to direct the Intel^® FPGA SDK for OpenCL™ Offline Compiler to disable pipelining of a loop. This pragma applies to single work-item kernels (that is, single-threaded kernels) in which loops are pipelined. Refer to the Single Work-Item Kernel versus NDRange Kernel section of the Intel^® FPGA SDK for OpenCL™ Pro Edition Best Practices Guide for information about loop pipelining and kernel properties that drive the offline compiler's decision about whether to treat a kernel as single-threaded.

Unless otherwise specified, the compiler always attempts to generate a pipelined loop datapath where possible. When generating a pipelined circuit, resources of the loop must be duplicated to execute multiple iterations simultaneously, leading to an increased silicon area utilization. In cases where loop pipelining does not result in an improvement in throughput, avoid the area overhead by applying the disable_loop_pipelining pragma to the loop, as shown in the following code snippet. When you apply this pragma, the offline compiler generates a simple sequential loop datapath.

#pragma 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)
}

In the above example, the offline compiler fails to schedule this loop with a small II due to memory dependency (as reported in the Details pane of the Loops Analysis section of the HTML report). In such cases, loop pipelining is unlikely to be beneficial.

Fusing adjacent loops can help reduce your kernel area use by reducing the overhead required for loop control and increasing the performance of your kernel by executing both original loops concurrently as one (fused) loop.

To specify a block of program code within which the compiler attempts to fuse loops, specify the pragma as follows:

#pragma loop_fuse [clause[[,]clause]...] new-line
    structured_block

where clause is one of the following:

depth(constant-integer-expression)

If a depth clause is present, the constant-integer-expression clause parameter defines the number of nesting depths at which the fusion of adjacent loops is attempted. The depth clause extends the applicability of the loop_fuse construct to all loops nested in top-level loops contained in the construct at nesting depth less-than or equal to the clause parameter, including loops that become adjacent as a result of fusion of their corresponding containing loops. In the absence of a depth clause, only loops at the top-level of the loop_fuse construct are attempted to be fused (that is, loops not contained in other loops defined within the construct). The depth clause with a parameter of 1 is equivalent to the absence of a depth clause.

independent

If an independent clause is present, adjacent loops that are fusion candidates within a loop_fuse construct are assumed to have no negative-distance data access dependencies. That is, for two adjacent loops considered for fusion, iterations of the logically-second loop does not access data elements produced in a later iteration of the logically-first loop. The independent clause overrides the offline compiler's static analysis during loop fusion safety analysis.

If a function call is present in a loop_fuse construct at any of the applicable nesting depths and inlining the function call materializes a loop, then the resulting loop is considered to be a candidate for fusion.

#pragma loop_fuse depth(2) independent
{
  L1: for(...) {}
  L2: for(...) {
    #pragma loop_fuse depth(2)
    {
      L3: for(...) {}
      L4: for(...) {
        L5: for(...) {}
        L6: for(...) {}
      }
    }
  }
}

To specify a loop to not be eligible for fusion, specify the nofusion pragma as follows:

#pragma nofusion 
for (int i = 0; i < N; ++i) {
    ...
}

In the following example, the compiler does not apply the loop fusion transformation to loops L1 and L2.

#pragma nofusion
L1: for (int j=0; j < N; ++j) {
   data[j] += Q;
}
L2: for (int i = 0; i < N; ++l) {
   output[i] = Q * data[i];
}

As an example, consider the loop nest in the following code snippet:

// Loop j is pipelined with ii=1

for (int j = 0; j < M; j++) {

  int a[N];

  // Loop i is pipelined with ii=2

  for (int i = 1; i < N; i++) {
      a[i] = foo(i)
  }
}

In this example, the inner i loop is pipelined with a loop II of 2. Under normal pipelining, this means that the inner loop hardware only achieves 50% utilization since one i iteration is initiated every other cycle. To take advantage of these idle cycles, the compiler interleaves a second invocation of the i loop from the next iteration of the outer j loop. Here, a loop invocation means to start pipelined execution of a loop body. In this example, since the i loop resides inside the j loop, and the j loop has a trip count of M, the i loop is invoked M times. Since the j loop is an outermost loop, it is invoked once. The following table illustrates the difference between normal pipelined execution of the i loop and interleaved execution for this example where N=5:

The table shows the values (j,i) for each inner loop iteration that is initiated at each cycle. At cycle 0, both modes of execution initiate the (0,0)^th iteration of the i loop. Under normal pipelined execution, no i loop iteration is initiated at cycle 1. Under interleaved execution, the (1,0)^th iteration of the innermost loop, that is, the first iteration of the next (j=1) invocation of the i loop is initiated. By cycle 10, interleaved execution has initiated all of the iterations of both the j=0 invocation of the i loop and the j=1 invocation of the i loop. This represents twice the efficiency of the normal pipelined execution.

In some cases, you may decide that the performance benefit from interleaving is not equal to the area cost associated with enabling interleaving. In these cases, you may want to limit or restrict the amount of interleaving to reduce FPGA area utilization. To limit the number of interleaved invocations of an inner loop that can be executed simultaneously, annotate the inner loop with the max_interleaving pragma. The annotated loop must be contained inside another pipelined loop. The required parameter ( n) specifies an upper bound on the degree of interleaving allowed, that is, how many invocations of the containing loop can execute the annotated loop at a given time.

Specify the max_interleaving pragma in one of the following ways:

• #pragma max_interleaving 1

  The compiler restricts the annotated (inner) loop to be invoked only once per outer loop iteration. That is, all iterations of the inner loop travels the pipeline before the next invocation of the inner loop can occur.

• #pragma max_interleaving 0

  The compiler allows the pipeline to contain a number of simultaneous invocations of the inner loop equal to the loop initiation interval (II) of the inner loop. For example, an inner loop with an II of 2 can have iterations from two invocations in the pipeline at a time. This behavior is the default behavior for the compiler if you do not specify the max_interleaving pragma.

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

Implementation of channels decouples data movement between concurrently executing kernels from the host processor.

Consider the following code example:

channel int c0;

__kernel void producer() {
    for (int i = 0; i < 10; i++) {
        write_channel_intel (c0, i);
    }
}

__kernel void consumer (__global uint * restrict dst) {
    for (int i = 0; i < 5; i++) {
        dst[i] = read_channel_intel(c0);
    }
}

The kernel producer writes ten elements ([0, 9]) to the channel. The kernel consumer does not contain any work-item identifier queries; therefore, it receives an implicit reqd_work_group_size attribute of (1,1,1). The implied reqd_work_group_size(1,1,1) attribute means that consumer must be launched as a single work-item kernel. In the example above, consumer reads five elements from the channel per invocation. During the first invocation, the kernel consumer reads values 0 to 4 from the channel. Because the data persists across NDRange invocations, the second time you execute the kernel consumer, it reads values 5 to 9.

For this example, to avoid a deadlock from occurring, you need to invoke the kernel consumer twice for every invocation of the kernel producer. If you call consumer less than twice, producer stalls because the channel becomes full. If you call consumer more than twice, consumer stalls because there is insufficient data in the channel.

Multiple work-item accesses to a channel can be useful in some scenarios. For example, they are useful when data words in the channel are independent, or when the channel is implemented for control logic. The main concern regarding multiple work-item accesses to a channel is the order in which the kernel writes data to and reads data from the channel. If possible, the SDK's channels extension processes work-item read and write operations to the channel in a deterministic order. As such, the read and write operations remain consistent across kernel invocations.

__kernel void ordering (__global int * restrict check,
                        __global int * restrict data) {
  int condition = check[get_global_id(0)];

  if (condition) {
    for (int i = 0; i < N, i++) {
      process(data);
      write_channel_intel (req, data[i]);
    }
  }
  else {
    process(data);
  }
}

When you implement channels in a kernel, the Intel^® FPGA SDK for OpenCL™ Offline Compiler enforces that kernel behavior is equivalent to having at most one work-group in flight within the compute unit at a time. The compiler also ensures that the kernel executes channels in work-item serial execution, where the kernel executes work-items with smaller IDs first. A work-item has the identifier (x, y, z, group), where x, y, z are the local 3D identifiers, and group is the work-group identifier.

The work-item ID (x0, y0, z0, group0) is considered to be smaller than the ID (x1, y1, z1, group1) if one of the following conditions is true:

• group0 < group1
• group0 = group1 and z0 < z1
• group0 = group1 and z0 = z1 and y0 < y1
• group0 = group1 and z0 = z1 and y0 = y1 and x0 < x1

Work-items with incremental IDs execute in a sequential order. For example, the work-item with an ID (x0, y0, z0, group0) executes the write channel call first. Then, the work-item with an ID (x1, y0, z0, group0) executes the call, and so on. Defining this order ensures that the system is verifiable with external models.

__kernel void ordering (__global int * data, int X) {
  int n = 0;
  while (n < X)
  {	
    write_channel_intel (req, data[get_global_id(0)]);
    n++;
  }
}

When you emulate a kernel that has a channel declared with the io attribute, I/O channel input is emulated by reading from a file, and channel output is emulated by writing to a file.

channel uint chanA __attribute__((io("myIOChannel")));

channel uint readChannel __attribute__((io("myIOChannel")));
channel uint writeChannel __attribute__((io("myIOChannel")));

Implementation of pipes decouples kernel execution from the host processor. The foundation of the Intel^® FPGA SDK for OpenCL™ pipes support is the SDK's channels extension. However, the syntax for pipe functions differs from the channels syntax.

For more information on blocking and nonblocking functions, refer to the corresponding documentation on channels.

Consider the following code example:

__kernel void producer (write_only pipe uint __attribute__((blocking)) c0)
{
    for (uint i = 0; i < 10; i++)
    {
        write_pipe (c0, &i);
    }
}

__kernel void consumer (__global uint * restrict dst,
                        read_only pipe uint __attribute__((blocking))
                        __attribute__((depth(10))) c0)
{
    for (int i = 0; i < 5; i++)
    {
        read_pipe (c0, &dst[i]);
    }
}

A read operation to a pipe reads the least recent piece of data written to the pipe first. Pipe data maintains a FIFO ordering within the pipe.

The kernel producer writes ten elements ([0, 9]) to the pipe. The kernel consumer reads five elements from the pipe per NDRange invocation. During the first invocation, the kernel consumer reads values 0 to 4 from the pipe. Because the data persists across NDRange invocations, the second time you execute the kernel consumer, it reads values 5 to 9.

For this example, to avoid a deadlock from occurring, you need to invoke the kernel consumer twice for every invocation of the kernel producer. If you call consumer less than twice, producer stalls because the pipe becomes full. If you call consumer more than twice, consumer stalls because there is insufficient data in the pipe.

Multiple work-item accesses to a pipe can be useful in some scenarios. For example, they are useful when data words in the pipe are independent, or when the pipe is implemented for control logic. The main concern regarding multiple work-item accesses to a pipe is the order in which the kernel writes data to and reads data from the pipe. If possible, the OpenCL pipes process work-items read and write operations to a pipe in a deterministic order. As such, the read and write operations remain consistent across kernel invocations.

__kernel void
ordering (__global int * check, global int * data,
          write_only pipe int __attribute__((blocking)) req)
{
    int condition = check[get_global_id(0)];

    if (condition)
    {
        for (int i = 0; i < N; i++)
        {
            process(data);
            write_pipe (req, &data[i]);
        }
    }
	else
    {
        process(data);
    }
}

When you implement pipes in a kernel, the Intel^® FPGA SDK for OpenCL™ Offline Compiler enforces that kernel behavior is equivalent to having at most one work-group in flight. The offline compiler also ensures that the kernel executes pipes in work-item serial execution, where the kernel executes work items with smaller IDs first. A work-item has the identifier (x, y, z, group), where x, y, z are the local 3D identifiers, and group is the work-group identifier.

The work-item ID (x0, y0, z0, group0) is considered to be smaller than the ID (x1, y1, z1, group1) if one of the following conditions is true:

• group0 < group1
• group0 = group1 and z0 < z1
• group0 = group1 and z0 = z1 and y0 < y1
• group0 = group1 and z0 = z1 and y0 = y1 and x0 < x1

Work items with incremental IDs execute in a sequential order. For example, the work-item with an ID (x0, y0, z0, group0) executes the write channel call first. Then, the work-item with an ID (x1, y0, z0, group0) executes the call, and so on. Defining this order ensures that the system is verifiable with external models.

__kernel void ordering (__global int * data,
                        write_only pipe int __attribute__((blocking)) req)
{
    write_pipe (req, &data[get_global_id(0)]);
}

The extension provides two new values in the flags argument of clCreatePipe to make a pipe host accessible, and adds four new API functions (clReadPipeIntelFPGA, clWritePipeIntelFPGA, clMapHostPipeIntelFPGA, and clUnmapHostPipeIntelFPGA) to allow the host to read from and write to a pipe that was created with host access enabled. A new optional kernel argument attribute is added to specify in the kernel language that the opposing end of a pipe kernel argument is the host program, and consequently that the pipe is not connected to another kernel. A pipe kernel argument is specialized in the kernel definition to connect to either a host pipe or another kernel, and cannot dynamically switch between the two at runtime.

When a pipe kernel argument is marked for host accessibility, the kernel language pipe accessors are restricted to a subset of the 2.x functions (reservations are not supported), and memory consistency or visibility guarantees are made beyond OpenCL™ synchronization points.

__attribute__((intel_host_accessible))

• clReadPipeIntelFPGA and clWritePipeIntelFPGA functions operate on single words of the pipe’s width.
• clMapHostPipeIntelFPGA function is an advanced mechanism to reduce latency and overhead when performing many word reads or writes on a host pipe.
• clUnmapHostPipeIntelFPGA function allows the host program to signal to the OpenCL runtime that it has written to or read from either a portion of or the entire mapped region that was created through a previous clMapHostPipeIntelFPGA function call.

The following sections describe the API functions for bound cl_mem objects:

cl_int clReadPipeIntelFPGA (cl_mem pipe, 
                            gentype *ptr);

cl_int clWritePipeIntelFPGA (cl_mem pipe, 
                             gentype *ptr);

void * clMapHostPipeIntelFPGA (cl_mem pipe,
                               cl_map_flags map_flags,
                               size_t requested_size,
                               size_t * mapped_size,
                               cl_int * errcode_ret);

cl_int clUnmapHostPipeIntelFPGA (cl_mem pipe, 
                                 void * mapped_ptr,
                                 size_t size_to_unmap,
                                 size_t * unmapped_size);

To enable host access (reading or writing) to pipes, the cl_intel_fpga_host_pipe extension legalizes the following two flags values to clCreatePipe:

• CL_MEM_HOST_READ_ONLY
• CL_MEM_HOST_WRITE_ONLY

When one of these flags is passed to the clCreatePipe function, the corresponding cl_mem object can be passed as the first argument to clReadPipeIntelFPGA and clWritePipeIntelFPGA functions. Throughout the remainder of the cl_intel_fpga_host_pipe extension, such a pipe is referred to as a host pipe.

Use the Intel^® FPGA SDK for OpenCL™ arbitrary precision integer extension to define integers with a custom bit-width. You can define integer custom bit-widths up to and including 64 bits.

#include "ihc_apint.h"

aoc <other command options> -I $INTELFPGAOCLSDKROOT/include/kernel_headers <my_kernel_file>

The header defines signed and unsigned arbitrary precision integer type definitions of the form intd_t and uintd_t, where d is the bit width of the type.

int10_t x_signed;
uint10_t x_unsigned;

You can declare arbitrary precision integers with widths up to 64 bits.

If you do operations where the bit width of the result is larger than the bit widths of the arguments, you must explicitly cast one of the arguments to the resulting bit width.

int10_t a;
int10_t b;
int20_t res;

res = a * b;

In the example, the compiler attempts to instantiate a multiplier that multiplies two 10-bit integers and put the results into another 10-bit integer. The result is then sign extended or zero extended up to 20-bits.

res = ((int20_t)a) * b

To match member data types, use the cl_ version of the data type in your host application that corresponds to the data type in the kernel code. The cl_ version of the data type is available in the opencl.h header file. For example, if you have a data member of type float4 in your kernel code, the corresponding data member you declare in the host application is cl_float4.

Align the structures and align the struct data members between the host and kernel applications. Manage the alignments carefully because of the variability among different host compilers.

For example, if you have float4 OpenCL data types in the struct, the alignments of these data items must satisfy the OpenCL specification (that is, 16-byte alignment for float4).

The following rules apply when the Intel^® FPGA SDK for OpenCL™ Offline Compiler compiles your OpenCL kernels:

1. Alignment of built-in scalar and vector types follow the rules outlined in Section 6.1.5 of the OpenCL Specification version 1.0.

   The offline compiler usually aligns a data type based on its size. However, the compiler aligns a value of a three-element vector the same way it aligns a four-element vector.

2. An array has the same alignment as one of its elements.
3. A struct (or a union) has the same alignment as the maximum alignment necessary for any of its data members.

   Consider the following example:

   struct my_struct
   {
       char data[3];
       float4 f4;
       int index;
   };

   The offline compiler aligns the struct elements above at 16-byte boundaries because of the float4 data type. As a result, both data and index also have 16-byte alignment boundaries.

4. The offline compiler does not reorder data members of a struct.
5. Normally, the offline compiler inserts a minimum amount of data structure padding between data members of a struct to satisfy the alignment requirements for each data member.
   1. In your OpenCL kernel code, you may specify data packing (that is, no insertion of data structure padding) by applying the packed attribute to the struct declaration. If you impose data packing, ensure that the alignment of data members satisfies the OpenCL alignment requirements. The Intel^® FPGA SDK for OpenCL™ does not enforce these alignment requirements. Ensure that your host compiler respects the kernel attribute and sets the appropriate alignments.
   2. In your OpenCL kernel code, you may specify the amount of data structure padding by applying the aligned(N) attribute to a data member, where N is the amount of padding. The SDK does not enforce these alignment requirements. Ensure that your host compiler respects the kernel attribute and sets the appropriate alignments.

      For Windows systems, some versions of the Microsoft Visual Studio compiler pack structure data types by default. If you do not want to apply data packing, specify an amount of data structure padding as shown below:

      struct my_struct
      {
          __declspec(align(16)) char data[3];
      
          /*Note that cl_float4 is the only known float4 definition on the host*/
          __declspec(align(16)) cl_float4 f4;
          
          __declspec(align(16)) int index;
      };
      

      Tip: An alternative way of adding data structure padding is to insert dummy struct members of type char or array of char.

The offline compiler infers private arrays as registers either as single values or in a piecewise fashion. Piecewise implementation results in very efficient hardware; however, the offline compiler must be able to determine data accesses statically. To facilitate piecewise implementation, hardcode the access points into the array. You can also facilitate register inference by unrolling loops that access the array.

If array accesses are not inferable statically, the offline compiler might infer the array as registers. However, the offline compiler limits the size of these arrays to 64 bytes in length for single work-item kernels. There is effectively no size limit for kernels with multiple work-items.

Consider the following code example:

int array[SIZE];
for (int j = 0; j < N; ++j)
{
    for (int i = 0; i < SIZE - 1; ++i)
    {
        array[i] = array[i + 1];
    }
}

The indexing into array[i] is not inferable statically because the loop is not unrolled. If the size of array[SIZE] is less than or equal to 64 bytes for single work-item kernels, the offline compiler implements array[SIZE] into registers as a single value. If the size of array[SIZE] is greater than 64 bytes for single work-item kernels, the offline compiler implements the entire array in block RAMs. For multiple work-item kernels, the offline compiler implements array[SIZE] into registers as a single value as long as its size is less than 1 kilobyte (KB).

Consider the following code example:

channel int in, out;

#define SIZE 512
//Shift register size must be statically determinable

__kernel void foo()
{
    int shift_reg[SIZE];
	   //The key is that the array size is a compile time constant

    // Initialization loop
    #pragma unroll
	for (int i=0; i < SIZE; i++)
	{
        //All elements of the array should be initialized to the same value
	    shift_reg[i] = 0;
    }
	
	while(1)
    {
        // Fully unrolling the shifting loop produces constant accesses
        #pragma unroll
        for (int j=0; j < SIZE–1; j++)
	    {
            shift_reg[j] = shift_reg[j + 1];
        } 
        shift_reg[SIZE – 1] = read_channel_intel(in);

        // Using fixed access points of the shift register
        int res = (shift_reg[0] + shift_reg[1]) / 2;

        // ‘out’ channel will have running average of the input channel
        write_channel_intel(out, res);
	}
}

In each clock cycle, the kernel shifts a new value into the array. By placing this shift register into a block RAM, the Intel^® FPGA SDK for OpenCL™ Offline Compiler can efficiently handle multiple access points into the array. The shift register design pattern is ideal for implementing filters (for example, image filters like a Sobel filter or time-delay filters like a finite impulse response (FIR) filter).

When implementing a shift register in your kernel code, keep in mind the following key points:

1. Unroll the shifting loop so that it can access every element of the array.
2. All access points must have constant data accesses. For example, if you write a calculation in nested loops using multiple access points, unroll these loops to establish the constant access points.
3. Initialize all elements of the array to the same value. Alternatively, you may leave the elements uninitialized if you do not require a specific initial value.
4. If some accesses to a large array are not inferable statically, they force the offline compiler to create inefficient hardware. If these accesses are necessary, use __local memory instead of __private memory.
5. Do not shift a large shift register conditionally. The shifting must occur in very loop iteration that contains the shifting code to avoid creating inefficient hardware.

Before declaring any double precision floating-point data type in your OpenCL kernel, include the following OPENCL EXTENSION pragma in your kernel code:

#pragma OPENCL EXTENSION cl_khr_fp64 : enable

The offline compiler supports an accumulator that adds or subtracts a value. To leverage this feature, describe the accumulation in a way that allows the offline compiler to infer the accumulator.

Below are examples of a description that results in the correct inference of the accumulator by the offline compiler.

channel float4 RANDOM_STREAM;

__kernel void acc_test(__global float *a, int k) {
    // Simplest example of an accumulator.
    // In this loop, the accumulator acc is incremented by 5.
    int i;
    float acc = 0.0f;
    for (i = 0; i < k; i++) {
        acc+=5;
    }
    a[0] = acc;
}

__kernel void acc_test2(__global float *a, int k) {
    // Extended example showing that an accumulator can be
    // conditionally incremented. The key here is to describe the increment
    // as conditional, not the accumulation itself.
    int i;
    float acc = 0.0f;
    for (i = 0; i < k; i++) {
        acc += ((i < 30) ? 5 : 0);
    }
    a[0] = acc;
}

__kernel void acc_test3(__global float *a, int k) {
    // A more complex case where the accumulator is fed
    // by a dot product.
    int i;
    float acc = 0.0f;
    for (i = 0; i < k; i++ ){
        float4 v = read_channel_intel(RANDOM_STREAM);
        float x1 = v.x;
        float x2 = v.y;
        float y1 = v.z;
        float y2 = v.w;
 
        acc += (x1*y1+x2*y2);
    }
    a[0] = acc;
}

__kernel void loader(__global float *a, int k) {
    int i;
    float4 my_val = 0;
    for(i = 0; i < k; i++) {
        if ((i%4) == 0)
            write_channel_intel(RANDOM_STREAM, my_val);
        if ((i%4) == 0) my_val.x = a[i];
        if ((i%4) == 1) my_val.y = a[i];
        if ((i%4) == 2) my_val.z = a[i];
        if ((i%4) == 3) my_val.w = a[i];
    }
}

• If both operands are of standard integer type (for example char or short), integers are promoted following the C/C++ standard. That is, the operation is carried out in the data type and size of the largest operand, with at least 32 bits. The expression returns the result in that larger data type.
• If both operands are intX_t data types, operations are carried out in the largest intX_t data type even if that data type is smaller than 32 bits. The expression returns the result in that type.
• If the expression has one standard data type and one intX_t data type, the rules for intX_t data type promotion apply. The resulting expression type is always an intX_t data type. For example, if the largest data type is a standard integer type short, the resulting data type is an int16_t.
• In C/C++, literals are by default an int data type, so when you use a literal without any casting, the expression type is always at least 32 bits. For example, if you have code as shown in the following snippet, the comparison is carried out in 32 bits:

  int5_t ap;
  ...
  if (ap < 4) {
  ...

• If operands are of different signage and the unsigned type is at least as large as the other type, the operation is carried out as an unsigned operation. Otherwise, the unsigned operand is converted to a signed value.

For example, if you have code as shown in the following snippet, -1 expands to a 32-bit negative number (0xffffffff) while the uint3_t ends up as the positive 32-bit number 7 (0x00000007), which are not equal.

uint3_t x = 7;
if (x != -1) {
   // FAIL
}

The host machine must support the following components:

• The host application and operating system.
• The working set for the host application.
• The maximum amount of OpenCL™ memory buffers that can be allocated at once. Every device-side cl_mem buffer is associated with a corresponding storage area in the host process. Therefore, the amount of host memory necessary might be as large as the amount of external memory supported by the FPGA.

All OpenCL APIs are thread safe except the clSetKernelArg function.

It is safe to call clSetKernelArg from any host thread or in a reentrant way as long as concurrent calls to clSetKernelArg operate on different cl_kernel objects.

A single in-order command queue can only dispatch a single operation for execution. Subsequent operations are not dispatched until the previous operation is fully complete. Thus, to execute kernels within the same OpenCL program object concurrently, instantiate a separate command queue for each kernel you want to run concurrently.

Similarly, multiple in-order command queues are also required to concurrently execute different transfers (clEnqueueReadBuffer or clEnqueueWriteBuffer), including executing transfers concurrently with kernels. For example, to execute kernel A concurrently with kernel B as well as concurrently with a clEnqueueWriteBuffer transfer, you must create three command queues and enqueue each of the operations in a separate queue. Achieving this in a steady state leads to maximum utilization of the FPGA device.

Out-of-order command queues may also be used to launch buffer writes, reads, and kernel execution concurrently. Enqueueing events that do not have any dependencies on one another into an out-of-order queue schedules them to execute concurrently, with the exception that these events are not blocked by their other dependencies.

The following clCreatePipe host API creates a pipe object:

cl_mem clCreatePipe(cl_context context,
                    cl_mem_flags flags,
                    cl_uint pipe_packet_size,
                    cl_uint pipe_max_packets,
                    const cl_pipe_properties *properties,
                    cl_int *errcode_ret)

For more information on the clCreatePipe host API function, refer to section 5.4.1 of the OpenCL Specification version 2.0.

Below is an example syntax of the clCreatePipe host API function:

cl_int status;
cl_mem c0_pipe = clCreatePipe(context,
                              0,
                              sizeof(int),
                              1,
                              NULL,
                              &status);
status = clSetKernelArg(kernel, 1, sizeof(cl_mem), &c0_pipe);

Following is the code snippet of the clGetProfileDataDeviceIntelFPGA host library call:

cl_int clGetProfileDataDeviceIntelFPGA (cl_device_id device_id,
                                        cl_program program,
                                        cl_bool read_enqueue_kernels,
                                        cl_bool read_auto_enqueued,
                                        cl_bool clear_counters_after_readback,
                                        size_t param_value_size,
                                        void *param_value,
                                        size_t *param_value_size_ret,
                                        cl_int *errcode_ret);

where,

• read_enqueue_kernels parameter profiles enqueued kernels. In this release, this parameter has no effect.
• read_auto_enqueued parameter profiles autorun kernels.
• Following are the placeholder parameters for the future releases:
  • clear_counters_after_readback
  • param_value_size
  • param_value
  • param_value_size_ret
  • errcode_ret

The clGetProfileDataDeviceIntelFPGA host library call returns CL_SUCCESS on success. Else, it returns one of the following errors:

• CL_INVALID_DEVICE if the device is not a valid device.
• CL_INVALID_PROGRAM if the program is not a valid program.

Pausing data acquisition is not synchronized exactly across all kernels. The skew between halting profile data acquisition across kernels is dependent on the communication link with the device, driver overhead, and congestion on communication buses. Exact synchronized snapshotting of profile data between kernels should not be relied upon.

You can present up to 128 FPGA devices to your system in the following manner:

• Multiple FPGA accelerator boards, each consisting of a single FPGA.
• Multiple FPGAs on a single accelerator board that connects to the host system via a PCIe^® switch.
• Combinations of the above.

The host runtime can load kernels onto each and every one of the FPGA devices. The FPGA devices can then operate in a parallel fashion.

The runtime environment is a library that is compiled as part of the host application. When the host application terminates, the runtime environment also terminates along with any tracking activity that it performs. If you restart the host application, a new runtime environment and its associated tracking activities reinitializes. The initialization functions reset the kernel's hardware state.

In same cases, unexpected termination of the host application causes the configuration of certain hardware (for example, PCIe^® hard IP) to be incomplete. To restore the configuration of these hardware, the host needs to reprogram the FPGA.

If you use a Custom Platform that implements customized hardware blocks, be aware that restarting the host application and resetting these blocks might have design implications:

• When the host application calls the clGetPlatformIDs function, all kernels and channels are reset for all available devices.
• When the host application calls the clGetPlatformIDs function, it resets FIFO buffers and channels as it resets the device.
• The host application initializes memory buffers via the clCreateBuffer and clEnqueueWriteBuffer function calls. You cannot access the contents of buffers from a previous host execution within a new host execution.

You can then pass these devices to clCreateContext function where the devices are locked by that OpenCL program until it either calls clReleaseContext function or terminates.

Figure 14. Sharing Multiple Devices Across Multiple Host Programs

Multiple processes or multiple trusted users can arbitrate between devices in a multi-device system using this locking mechanism. In case users have decided ahead of time which device (by name) each person will use, they can use clGetDeviceInfo function to select the cl_device_id with the name that a user is assigned to. To arbitrate more dynamically in case each of the N users want Di devices, use the following scheme:

1. Each user queries the clGetDeviceIDs function to obtain a list of devices.
2. Each user chooses Di devices (ideally randomly to minimize collisions) and passes those to the clCreateContext function.

It is possible that during step 2, another user may have already called clCreateContext function with that same device, in which case, the clCreateContext function call fails. The user should then repeat steps 2 (optionally changing the device selection) until it succeeds.

Consider the following example code snippet:

do {
    for(i = num_devices - 1; i >= 0; i--) {
      context = clCreateContext(0, 1, &(device_ids[i]), NULL, NULL, &status);

      if (status != CL_SUCCESS) {
        printf("Failed to get context with %d (error: %d), waiting\n", i, status
        sleep(1);
      } else {
        device_id = device_ids[i];
        break;
      }
    }
  } while (status != CL_SUCCESS);
// Exit this loop only when we’ve succeeded in creating a context

The Incremental compile report provides the following metrics on your OpenCL design:

• The <%> of design not preserved metric at the bottom of the report provides a quick summary of the overall changes to your design. It is the best predictor of compilation time.

  Note:

  The FPGA resources listed in the Incremental compile report are calculated based on the estimated area models that the Intel^® FPGA SDK for OpenCL™ Offline Compiler produces. The area numbers represent an estimate of the area usage in a standard (that is, non-incremental) compilation. You can use these numbers to gauge the area your design consumes in a standard compilation.

  The FPGA resource information might not fully match the final area in the Intel^® Quartus^® Prime Pro Edition software compilation reports.

Figure 15. Incremental compile Report for a Setup Compilation

Figure 16. Incremental compile Report for an Incremental Compilation

In addition to device support, the incremental compilation has the following limitations:

• You will experience area, F_max, and power degradations when you enable the incremental compilation feature ( -incremental ) or the fast compilation feature ( -fast-compile ), or both.
• In congested designs, incremental compilations can experience severe (that is, 25% or more) F_max reductions compared to the initial setup compilation. If the F_max reduction is unacceptable, perform a non-incremental fast compilation instead to reduce the amount of F_max degradation while preserving some of the savings in compilation time.

• The offline compiler does not detect changes in RTL libraries that you have included by invoking the -l <library_name>.aoclib offline compiler command option. After you modify an RTL library, you must perform a setup compilation again.

  The offline compiler prints a warning message as a reminder to rerun the setup compilation.

• When compiling an OpenCL kernel containing calls to HLS tasks, incremental compile may trigger recompilation for unaffected kernels. However, this is not a functional bug. It may result in a more conservative incremental compile.

The Intel^® FPGA SDK for OpenCL™ Emulator generates a .aocx file that executes on x86-64 Windows or Linux host. This feature allows you to emulate the functionality of your kernel and iterate on your design without executing it on the actual FPGA each time. For Linux platform, you can also use the Emulator to perform functional debug.

The emulator supports 64-bit Windows and Linux operating systems. On Linux systems, the GNU C Library (glibc) version 2.15 or later is required. If you Linux system does not meet this requirements, you might be able to use the Legacy Emulator. For details, see Using the Legacy Emulator.

If you installed the Intel^® FPGA SDK for OpenCL™ Pro edition with administrator privileges, no additional setup is needed. If you did not install the Intel^® FPGA SDK for OpenCL™ with administrator privileges, you must perform some additional steps to enable the emulator.

env LD_LIBRARY_PATH=<path to sufficiently new libstdc++.so>:$LD_LIBRARY_PATH <host> [host arguments]

You may emulate a kernel that passes a channel or pipe by value, as shown in the following example:

channel uint my_ch;

void my_function (channel uint ch,
                  __global uint * dst, int i)
{
    dst[i] = read_channel_intel(ch);
}

__kernel void
consumer (__global uint * restrict dst)
{
    for (int i=0;i<5;i++)
    {
        my_function(my_ch, dst, i );
    }
}

When you compile your OpenCL* kernel for emulation, the default channel depth is different from the default channel depth generated when your kernel is compiled for hardware. You can change this behavior when you compile your kernel for emulation with the CL_CONFIG_CHANNEL_DEPTH_EMULATION_MODE environment variable.

• Execution model

  The Emulator supports the same compilation modes as the FPGA variant. As a result, you must call the clCreateProgramBinary function to create cl_program objects for emulation.

• Concurrent execution

  Modeling of concurrent kernel executions has limitations. During execution, the Emulator does not actually run interacting work-items in parallel. Therefore, some concurrent execution behaviors, such as different kernels accessing global memory without a barrier for synchronization, might generate inconsistent emulation results between executions.

• Same address space execution

  The Emulator executes the host runtime and the kernels in the same address space. Certain pointer or array usages in your host application might cause the kernel program to fail, and vice versa. Example usages include indexing external allocated memory and writing to random pointers. You may use memory leak detection tools such as Valgrind to analyze your program. However, the host might encounter a fatal error caused by out-of-bounds write operations in your kernel, and vice versa.

• Conditional channel operations

  Emulation of channel behavior has limitations, especially for conditional channel operations where the kernel does not call the channel operation in every loop iteration. In these cases, the Emulator might execute channel operations in a different order than on the hardware.

• GCC version

  Emulator host programs on Linux* must be run with a version of libstdc++.so from GCC 7.2.0 or later. This can be achieved either by installing GCC 7.2.0 or later on your system, or setting LD_LIBRARY_PATH variable such that a sufficiently new libstdc++.so is identified.

A few known issues might affect your use of the emulator. Review these issues to avoid possible problems when using the emulator.

Many of the recommendations outlined in Emulating and Debugging Your OpenCL Kernel apply to both emulators. However, the Legacy Emulator has a few differences that are outlined in the sections that follow.

When you compile your OpenCL* kernel for emulation, the default channel depth is different from the default channel depth generated when your kernel is compiled for hardware. You can change this behavior when you compile your kernel for emulation with the -emulator-channel-depth-model option.

The Intel^® Code Builder for OpenCL™ provides a set of Microsoft Visual Studio and Eclipse plug-ins that enable capabilities for creating, building, debugging, and analyzing Windows and Linux applications accelerated with OpenCL.

You can configure a session by right-clicking the session in the Code Builder Session Explorer and selecting Session Options. Alternatively, you can also open the Session Settings dialog box by selecting Code-Builder > OpenCL Kernel Development > Session Options.

The Session Settings dialog box allows you to configure:

• Device options such as target machine, OpenCL platform, and OpenCL device.
• Build options such as offline compiler flags and build architecture.
• Build artifacts such as .aocx and .aoco files, and static reports.
• General options such as job architecture and network settings.

In the Device Options tab, ensure to select Intel^® FPGA SDK for OpenCL™ in the OpenCL platform drop-down list.

Under the Build Options tab, in the OpenCL Build Options section, enter the Intel^® FPGA SDK for OpenCL™ Offline Compiler flags manually.

For more information about configuring a session and variable management, refer to the Developer Guide for Intel SDK for OpenCL Applications.

You can create an OpenCL library in OpenCL or register transfer level (RTL). You can then include this library file and use the functions inside your OpenCL kernels or in HLS components. For information about HLS libraries, refer to Intel High Level Synthesis Compiler: Reference Manual .

You may use a previously-created library or create your own library. To use an OpenCL library, you do not require in-depth knowledge in hardware design or in the implementation of library primitives. To create an OpenCL library, you need to create the following files:

Creating a OpenCL library is a two-step process:

1. Each object file is created from input source files using the fpga_crossgen command. The required input source files depend on the type of source code you are creating the object from.
   • An object is effectively an intermediate representation of your source code with both a CPU representation and an FPGA representation of your code.
   • An object can be targeted for use with only one Intel^® high-level design product. If you want to target more than one high-level design product, you must generate a separate object for each target product.
2. Object files are combined into a library file using the fpga_libtool command.

   Objects created from different types of source code can be combined into a library, provided all objects target the same high-level design product.

   A library is assigned a version number, and can be used only with the targeted high-level design product with the same version number (for example, Intel^® FPGA SDK for OpenCL™ Pro Edition version 20.2).

You can create a library from object files from your OpenCL source code. An OpenCL-based object file includes code for CPU as well as hardware execution (CPU-capturing testbench and emulation use).

A library can contain multiple object files. You can create object files for use in different Intel high-level design tools from the same OpenCL source code. Depending on the target high-level design tool, your source code might require adjustments to support tool-specific data types or constructs.