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.
• 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 .aoco file
• 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 a .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 will read 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 will be 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.

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 will receive 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 legalizes two new values in the flags argument of clCreatePipe to make a pipe host accessible, and adds four new API functions (clReadPipeIntelFPGA, clWritePipeIntelFPGA, clMapHostPipeIntelFPGA, 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 will be the host program, and consequently that the pipe will not be 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.

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>

#define ap_int<d> intd_t
#define ap_uint<d> uintd_t

int10_t x_signed;
uint10_t x_unsigned;

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

#pragma OPENCL EXTENSION cl_intel_arbitrary_precision_integers : enable

ap_int<d> intd_t my_signed_integer
ap_uint<d> uintd_t my_unsigned_integer

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.

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 will also terminate along with any tracking activity that it performs. If you restart the host application, a new runtime environment and its associated tracking activities will reinitialize. 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 will be 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.

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 gage the area your design will consume 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 14. Incremental compile Report for a Setup Compilation

Figure 15. Incremental compile Report for an Incremental Compilation

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

• You will experience area, Fmax, 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) Fmax reductions compared to the initial setup compilation. If the Fmax reduction is unacceptable, perform a non-incremental fast compilation instead to reduce the amount of Fmax 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 will print a warning message as a reminder to rerun the setup compilation.

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 fast emulator.

If you did not install the Intel^® FPGA SDK for OpenCL™ with administrator privileges, complete these steps before using the Fast Emulator:

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.

• 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.

• 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.

A few known issues might affect your use of the fast emulator. Review these issues to avoid possible problems when using the fast 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 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 components. To create an OpenCL library, you need to create the following files and components:

Assume each level of operation is one stage in the pipeline. At each stage, the Intel^® FPGA SDK for OpenCL™ Offline Compiler executes all operations in parallel by the thread existing at that stage. For example, thread 2 executes Load A, Load B, and copies the current global ID (via gid) to the next pipeline stage. Similar to the pipelined execution on instructions in reduced instruction set computing (RISC) processors, the SDK's pipeline stages also execute in parallel. The threads will advance to the next pipeline stage only after all the stages have completed execution.

Some operations are capable of stalling the Intel FPGA SDK for OpenCL pipeline. Examples of such operations include variable latency operations like memory load and store operations. To support stalls, ready and valid signals need to propagate throughout the pipeline so that the offline compiler can schedule the pipeline stages. However, ready signals are not necessary if all operations have fixed latency. In these cases, the offline compiler optimizes the pipeline to statically schedule the operations, which significantly reduces the logic necessary for pipeline implementation.

The depicted RTL module has a balanced latency where the threads of the RTL module match the number of pipeline stages. A balanced latency allows the threads of the RTL module to execute without stalling the SDK's pipeline.

Setting the latency of the RTL module in the RTL specification file allows the offline compiler to balance the pipeline latency. RTL supports Avalon™ Streaming ( Avalon™ -ST) interfaces; therefore, the latency of the RTL module can be variable (that is, not fixed). However, the variability in the latency should be small in order to maximize performance. In addition, specify the latency in the <RTL module description file name>.xml specification file so that the RTL module experiences a good approximation of the actual latency in steady state.

A stall-free RTL module might receive an invalid input signal (that is, ivalid is low). In this case, the module ignores the input and produces invalid data on the output. For a stall-free RTL module without an internal state, it might be easier to propagate the invalid input through the module. However, for an RTL module with an internal state, you must handle an ivalid=0 input carefully.

For an RTL module to properly interact with other compiler-generated operations, you must support a simplified Avalon Streaming Interface at both input and output of an RTL module.

The following diagram shows the complete interface of the myMod RTL module shown in Figure 19.

In this diagram, myMod interacts with the upstream module through data signals, A and B, and control signals, ivalid (input) and oready (output). The ivalid control signal equals 1 (ivalid = 1) if and only if data signal A and data signal B contain valid data. When the control signal oready equals 1 (oready = 1), it indicates that the myMod RTL module can process the data signals A and B if they are valid (that is, ivalid = 1). When ivalid = 1 and oready = 0, the upstream module is expected to hold the values of ivalid, A, and B in the next clock cycle.

myMod interacts with the downstream module through data signal C and control signals, ovalid (output) and iready (input). The ovalid control signal equals 1 (ovalid = 1) if and only if data signal C contains valid data. The iready control signal equals 1 (ivalid = 1) indicates that the downstream module is able to process data signal C if it is valid. When ovalid = 1 and iready = 0, the myMod RLT module is expected to hold the valid of the ovalid and C signals in the next clock cycle.

myMod module will assert oready for a single clock cycle to indicate it is ready for an active cycle. Cycles during which myMod module is ready for data are called ready cycles. During ready cycles, the module above myMod module can assert ivalid to send data to myMod.

For a detailed explanation of data transfer under backpressure, refer to " Data Transfer with Backpressure " in Avalon Interface Specifications. Ignore the information pertaining to readyLatency option.

The offline compiler expects the RTL module to support Avalon-ST interface with readyLatency = 0, at both input and output.

The following figure shows the timing diagram for input data transfer with back pressure. For more information about Avalon-ST interfaces, see the " Avalon Streaming Interfaces " section in Avalon Interface Specifications.

For an RTL module with a fixed latency, the output signals (ovalid and oready) can have constant high values, and the input ready signal (iready) can be ignored.

A stall-free RTL module might receive an invalid input signal (ivalid is low). In this case, the module ignores the input and produces invalid data on the output. For a stall-free RTL module without an internal state, it might be easier to propagate the invalid input through the module. However, for an RTL module with an internal state, you must handle an ivalid = 0 input carefully.

Because of the common clock and reset drivers, an RTL module runs in the same clock domain as the OpenCL kernel. The module is reset only when the OpenCL kernel is first loaded onto the FPGA, either via Intel^® FPGA SDK for OpenCL™ program utility or the clCreateProgramwithBinary host function. In particular, if the host restarts a kernel via successive clEnqueueNDRangeKernel or clEnqueueTask invocations, the associated RTL modules will not reset in between these restarts.

The following steps outline the process of setting the kernel clock frequency:

1. The Intel^® Quartus^® Prime software's Fitter applies an aggressive constraint on the kernel clock.
2. The Intel^® Quartus^® Prime software's Timing Analyzer performs static timing analysis to determine the frequency that the Fitter actually achieves.
3. The phase-locked loop (PLL) that drives the kernel clock sets the frequency determined in Step 2 to be the kernel clock frequency.

Optionally, an RTL module can access a system-wide clock that runs at twice the frequency of the OpenCL™ kernel clock. This system-wide clock can be connected to an input signal of the RTL module by including an AVALON element of type clock2x. These two clocks do not have a defined phase-relationship.

Allow your RTL module to interact with external memory only if the interaction is necessary and unavoidable.

The following examples demonstrate how to structure code in an RTL module for easy integration into an OpenCL library:

// my_rtl_fn does:
// out_ptr[idx] = fn(in_ptr[idx])
my_rtl_fn (in_ptr, out_ptr,idx);

int in_value = in_ptr[idx];
// my_rtl_fn now does: out = fn(in)
int out_value = my_rtl_fn (in_value);
out_ptr[idx] = out_value;

The complex RTL module on the left reads a value from external memory, performs a scalar function on the value, and then writes the value back to global memory. Such an RTL module is difficult to describe when you integrate it into an OpenCL library. In addition, this RTL module is harder to verify and causes very conservative pointer analysis in the Intel^® FPGA SDK for OpenCL™ Offline Compiler.

The simplified RTL module on the right provides the same overall functionality as the complex RTL module. However, the simplified RTL module only performs a scalar-to-scalar calculation without connecting to global memory. Integrating this simplified RTL module into the OpenCL library makes it much easier for the offline compiler to analyze the resulting OpenCL kernel.

There are times when an RTL module requires an Avalon^®-MM port to communicate with external memory. This Avalon-MM port connects to the same arbitration network to which all other global load and store units in the OpenCL kernels connect.

If an RTL module receives a memory pointer as an argument, the offline compiler enforces the following memory model:

• If an RTL module writes to a pointer, nothing else in the OpenCL kernel can read from or write to this pointer.
• If an RTL module reads from a pointer, the rest of the OpenCL kernel and other RTL modules may also read from this pointer.

• You may set the access field of the MEM_INPUT attribute to specify how the RTL module uses the memory pointer. Ensure that you set the value for access correctly because there is no way to verify the value.

Example OpenCL C model file for a square root function:

double my_sqrtfd (double a)
{
    return sqrt(a);
}

Intel^® recommends that you emulate your OpenCL system. If you decide not to emulate your OpenCL system, create an empty function with a name that matches the function name you specified in the XML specification file.

Consider a situation where you create and verify your library on a device that does not support Partial Reconfiguration (PR). If a library user then uses the library's RTL module inside a PR region, the module might not function correctly after PR.

For complete PR coding guidelines, refer to Creating a Partial Reconfiguration Design in the Partial Reconfiguration User Guide.

When creating your RTL module, ensure that it operates within the following restrictions:

• An RTL module must use a single input Avalon^®-ST interface. That is, a single pair of ready and valid logic must control all the inputs.

  You have the option to provide the necessary Avalon-ST ports but declare the RTL module as stall-free. In this case, you do not have to implement proper stall behavior because the Intel^® FPGA SDK for OpenCL™ Offline Compiler creates a wrapper for your module. Refer to XML Syntax of an RTL Module and Using an OpenCL Library that Works with Simple Functions (Example 1) for more syntax and usage information, respectively.

  Note: You must handle ivalid signals properly if your RTL module has an internal state. Refer to Stall-Free RTL for more information.
• The RTL module must work correctly regardless of the kernel clock frequency.
• Data input and output sizes must match valid OpenCL data types, from 8 bits for char to 1024 bits for long16.

  For example, if you work with 24-bit values inside an RTL module, declare inputs to be 32 bits and declare function signature in the SDK's library header file to accept the uint data type. Then, inside the RTL module, accept the 32-bit input but discard the top 8 bits.

• RTL modules cannot connect to external I/O signals. All input and output signals must come from an OpenCL kernel.
• An RTL module must have a clock port, a resetn port, and Avalon-ST input and output ports (that is, ivalid, ovalid, iready, oready). Name the ports as specified here.
• RTL modules that communicate with external memory must have Avalon Memory-Mapped (Avalon-MM) port parameters that match the corresponding Custom Platform parameters. The offline compiler does not perform any width or burst adaptation.
• RTL modules that communicate with external memory must behave as follows:
  • They cannot burst across the burst boundary.
  • They cannot make requests every clock cycle and stall the hardware by monopolizing the arbitration logic. An RTL module must pause its requests regularly to allow other load or store units to execute their operations.
• RTL modules cannot act as stand-alone OpenCL kernels. RTL modules can only be helper functions and be integrated into an OpenCL kernel during kernel compilation.
• Every function call that corresponds to RTL module instantiation is completely independent of other instantiations. There is no hardware sharing.
• Do not incorporate kernel code (that is, functions marked as kernel) into a .aoclib library file. Incorporating kernel code into the library file causes the offline compiler to issue an error message. You may incorporate helper functions into the library file.
• An RTL component must receive all its inputs at the same time. A single ivalid input signifies that all inputs contain valid data.
• The SDK does not support I/O RTL modules.
• You can only set RTL module parameters in the <RTL module description file name>.xml specification file, not the OpenCL kernel source file. To use the same RTL module with multiple parameters, create a separate FUNCTION tag for each parameter combination.

Currently, the SDK's RTL module support for the library feature has the following limitations: