Intel® FPGA SDK for OpenCL™ Standard Edition: Programming Guide

Download PDF
ID 683342
Date 4/22/2019
Public
Document Table of Contents
Document Table of Contents x
1. Intel® FPGA SDK for OpenCL™ Standard Edition Overview 2. Intel® FPGA SDK for OpenCL™ Offline Compiler Kernel Compilation Flows 3. Obtaining General Information on Software, Compiler, and Custom Platform 4. Managing an FPGA Board 5. Structuring Your OpenCL Kernel 6. Designing Your Host Application 7. Compiling Your OpenCL Kernel 8. Emulating and Debugging Your OpenCL Kernel 9. Reviewing Your Kernel's report.html File 10. Profiling Your OpenCL Kernel 11. Developing OpenCL™ Applications Using Intel® Code Builder for OpenCL™ 12. Intel® FPGA SDK for OpenCL™ Standard Edition Advanced Features A. Support Statuses of OpenCL Features B. Document Revision History of the Intel® FPGA SDK for OpenCL™ Standard Edition Programming Guide
1. Intel® FPGA SDK for OpenCL™ Standard Edition Overview x
1.1. Intel® FPGA SDK for OpenCL™ Standard Edition Programming Guide Prerequisites 1.2. Intel® FPGA SDK for OpenCL™ Standard Edition FPGA Programming Flow
2. Intel® FPGA SDK for OpenCL™ Offline Compiler Kernel Compilation Flows x
2.1. One-Step Compilation for Simple Kernels 2.2. Multistep Intel® FPGA SDK for OpenCL™ Standard Edition Design Flow
3. Obtaining General Information on Software, Compiler, and Custom Platform x
3.1. Displaying the Software Version (version) 3.2. Displaying the Compiler Version (-version) 3.3. Listing the Intel® FPGA SDK for OpenCL™ Standard Edition Utility Command Options (help) 3.4. Listing the Intel® FPGA SDK for OpenCL™ Offline Compiler Command Options (no argument, -help, or -h) 3.5. Listing the Available FPGA Boards in Your Custom Platform (-list-boards) 3.6. Displaying the Compilation Environment of an OpenCL Binary (env)
3.3. Listing the Intel® FPGA SDK for OpenCL™ Standard Edition Utility Command Options (help) x
3.3.1. Displaying Information on an Intel® FPGA SDK for OpenCL™ Utility Command Option (help <command_option>)
4. Managing an FPGA Board x
4.1. Installing an FPGA Board (install) 4.2. Uninstalling the FPGA Board (uninstall) 4.3. Querying the Device Name of Your FPGA Board (diagnose) 4.4. Running a Board Diagnostic Test (diagnose <device_name>) 4.5. Programming the FPGA Offline or without a Host (program <device_name>) 4.6. Programming the Flash Memory (flash <device_name>)
5. Structuring Your OpenCL Kernel x
5.1. Guidelines for Naming the Kernel 5.2. Programming Strategies for Optimizing Data Processing Efficiency 5.3. Programming Strategies for Optimizing Pointer-to-Local Memory Size 5.4. Implementing the Intel® FPGA SDK for OpenCL™ Standard Edition Channels Extension 5.5. Implementing OpenCL Pipes 5.6. Implementing Arbitrary Precision Integers 5.7. Using Predefined Preprocessor Macros in Conditional Compilation 5.8. Declaring __constant Address Space Qualifiers 5.9. Including Structure Data Types as Arguments in OpenCL Kernels 5.10. Inferring a Register 5.11. Enabling Double Precision Floating-Point Operations 5.12. Single-Cycle Floating-Point Accumulator for Single Work-Item Kernels
5.2. Programming Strategies for Optimizing Data Processing Efficiency x
5.2.1. Unrolling a Loop 5.2.2. Coalescing Nested Loops 5.2.3. Specifying a Loop Initiation interval (II) 5.2.4. Loop Concurrency (max_concurrency Pragma) 5.2.5. Specifying Work-Group Sizes 5.2.6. Specifying Number of Compute Units 5.2.7. Specifying Number of SIMD Work-Items
5.4. Implementing the Intel® FPGA SDK for OpenCL™ Standard Edition Channels Extension x
5.4.1. Overview of the Intel® FPGA SDK for OpenCL™ Standard Edition Channels Extension 5.4.2. Channel Data Behavior 5.4.3. Multiple Work-Item Ordering for Channels 5.4.4. Restrictions in the Implementation of Intel® FPGA SDK for OpenCL™ Standard Edition Channels Extension 5.4.5. Enabling the Intel® FPGA SDK for OpenCL™ Standard Edition Channels for OpenCL Kernel
5.4.3. Multiple Work-Item Ordering for Channels x
5.4.3.1. Work-Item Serial Execution of Channels
5.4.5. Enabling the Intel® FPGA SDK for OpenCL™ Standard Edition Channels for OpenCL Kernel x
5.4.5.1. Declaring the Channel Handle 5.4.5.2. Implementing Blocking Channel Writes 5.4.5.3. Implementing Blocking Channel Reads 5.4.5.4. Implementing I/O Channels Using the io Channels Attribute 5.4.5.5. Emulating I/O Channels 5.4.5.6. Use Models of Intel® FPGA SDK for OpenCL™ Standard Edition Channels Implementation 5.4.5.7. Implementing Buffered Channels Using the depth Channels Attribute 5.4.5.8. Enforcing the Order of Channel Calls
5.4.5.2. Implementing Blocking Channel Writes x
5.4.5.2.1. Implementing Nonblocking Channel Writes
5.4.5.3. Implementing Blocking Channel Reads x
5.4.5.3.1. Implementing Nonblocking Channel Reads
5.4.5.8. Enforcing the Order of Channel Calls x
5.4.5.8.1. Defining Memory Consistency Across Kernels When Using Channels
5.5. Implementing OpenCL Pipes x
5.5.1. Overview of the OpenCL Pipe Functions 5.5.2. Pipe Data Behavior 5.5.3. Multiple Work-Item Ordering for Pipes 5.5.4. Restrictions in OpenCL Pipes Implementation 5.5.5. Enabling OpenCL Pipes for Kernels 5.5.6. Direct Communication with Kernels via Host Pipes
5.5.3. Multiple Work-Item Ordering for Pipes x
5.5.3.1. Work-Item Serial Execution of Pipes
5.5.5. Enabling OpenCL Pipes for Kernels x
5.5.5.1. Ensuring Compatibility with Other OpenCL SDKs 5.5.5.2. Declaring the Pipe Handle 5.5.5.3. Implementing Pipe Writes 5.5.5.4. Implementing Pipe Reads 5.5.5.5. Implementing Buffered Pipes Using the depth Attribute 5.5.5.6. Implementing I/O Pipes Using the io Attribute 5.5.5.7. Enforcing the Order of Pipe Calls
5.5.5.7. Enforcing the Order of Pipe Calls x
5.5.5.7.1. Defining Memory Consistency Across Kernels When Using Pipes
5.5.6. Direct Communication with Kernels via Host Pipes x
5.5.6.1. Optional intel_host_accessible Kernel Argument Attribute 5.5.6.2. API Functions for Interacting with cl_mem Pipe Objects Bound to Host-Accessible Pipe Kernel Arguments 5.5.6.3. Creating a Host Accessible Pipe 5.5.6.4. Example Use of the cl_intel_fpga_host_pipe Extension
5.9. Including Structure Data Types as Arguments in OpenCL Kernels x
5.9.1. Matching Data Layouts of Host and Kernel Structure Data Types 5.9.2. Disabling Insertion of Data Structure Padding 5.9.3. Specifying the Alignment of a Struct
5.10. Inferring a Register x
5.10.1. Inferring a Shift Register
5.12. Single-Cycle Floating-Point Accumulator for Single Work-Item Kernels x
5.12.1. Programming Strategies for Inferring the Accumulator
6. Designing Your Host Application x
6.1. Host Programming Requirements 6.2. Allocating OpenCL Buffers for Manual Partitioning of Global Memory 6.3. Collecting Profile Data During Kernel Execution 6.4. Accessing Custom Platform-Specific Functions 6.5. Modifying Host Program for Structure Parameter Conversion 6.6. Managing Host Application 6.7. Allocating Shared Memory for OpenCL Kernels Targeting SoCs 6.8. Debugging Your OpenCL System That is Gradually Slowing Down
6.1. Host Programming Requirements x
6.1.1. Host Machine Memory Requirements 6.1.2. Host Binary Requirement 6.1.3. Multiple Host Threads 6.1.4. Out-of-Order Command Queues 6.1.5. Requirement for Multiple Command Queues to Execute Kernels Concurrently
6.2. Allocating OpenCL Buffers for Manual Partitioning of Global Memory x
6.2.1. Partitioning Buffers Across Multiple Interfaces of the Same Memory Type 6.2.2. Partitioning Buffers Across Different Memory Types (Heterogeneous Memory) 6.2.3. Creating a Pipe Object in Your Host Application
6.3. Collecting Profile Data During Kernel Execution x
6.3.1. Profiling Enqueued and Autorun Kernels 6.3.2. Profile Data Acquisition 6.3.3. Multiple Autorun Profiling Calls
6.6. Managing Host Application x
6.6.1. Displaying Example Makefile Fragments (example-makefile or makefile) 6.6.2. Compiling and Linking Your Host Application 6.6.3. Linking Your Host Application to the Khronos ICD Loader Library 6.6.4. Programming an FPGA via the Host 6.6.5. Termination of the Runtime Environment and Error Recovery
6.6.2. Compiling and Linking Your Host Application x
6.6.2.1. Displaying Flags for Compiling Host Application (compile-config) 6.6.2.2. Displaying Paths to OpenCL Host Runtime and MMD Libraries (ldflags) 6.6.2.3. Listing OpenCL Host Runtime and MMD Libraries (ldlibs) 6.6.2.4. Displaying Information on OpenCL Host Runtime and MMD Libraries (link-config or linkflags)
6.6.3. Linking Your Host Application to the Khronos ICD Loader Library x
6.6.3.1. Linking to the ICD Loader Library on Windows 6.6.3.2. Linking to the ICD Loader Library on Linux
6.6.4. Programming an FPGA via the Host x
6.6.4.1. Programming Multiple FPGA Devices
6.6.4.1. Programming Multiple FPGA Devices x
6.6.4.1.1. Probing the OpenCL FPGA Devices 6.6.4.1.2. Querying Device Information 6.6.4.1.3. Loading Kernels for Multiple FPGA Devices
7. Compiling Your OpenCL Kernel x
7.1. Compiling Your Kernel to Create Hardware Configuration File 7.2. Compiling Your Kernel without Building Hardware (-c) 7.3. Specifying the Location of Header Files (-I=<directory>) 7.4. Specifying the Name of an Intel® FPGA SDK for OpenCL™ Offline Compiler Output File (-o=<filename>) 7.5. Compiling a Kernel for a Specific FPGA Board (-board=<board_name>) 7.6. Resolving Hardware Generation Fitting Errors during Kernel Compilation (-high-effort) 7.7. Defining Preprocessor Macros to Specify Kernel Parameters (-D<macro_name>) 7.8. Generating Compilation Progress Report (-v) 7.9. Displaying the Estimated Resource Usage Summary On-Screen (-report) 7.10. Suppressing Warning Messages from the Intel® FPGA SDK for OpenCL™ Offline Compiler (-W) 7.11. Converting Warning Messages from the Intel® FPGA SDK for OpenCL™ Offline Compiler into Error Messages (-Werror) 7.12. Removing Debug Data from Compiler Reports and Source Code from the .aocx File (-g0) 7.13. Disabling Burst-Interleaving of Global Memory (-no-interleaving=<global_memory_type>) 7.14. Configuring Constant Memory Cache Size (-const-cache-bytes=<N>) 7.15. Relaxing the Order of Floating-Point Operations (-fp-relaxed) 7.16. Reducing Floating-Point Rounding Operations (-fpc)
8. Emulating and Debugging Your OpenCL Kernel x
8.1. Modifying Channels Kernel Code for Emulation 8.2. Compiling a Kernel for Emulation (-march=emulator) 8.3. Emulating Your OpenCL Kernel 8.4. Debugging Your OpenCL Kernel on Linux 8.5. Limitations of the Intel® FPGA SDK for OpenCL™ Standard Edition Emulator 8.6. Discrepancies in Hardware and Emulator Results
8.1. Modifying Channels Kernel Code for Emulation x
8.1.1. Emulating a Kernel that Passes Pipes or Channels by Reference 8.1.2. Emulating Channel Depth
10. Profiling Your OpenCL Kernel x
10.1. Instrumenting the Kernel Pipeline with Performance Counters (-profile) 10.2. Launching the Intel® FPGA Dynamic Profiler for OpenCL™ GUI (report) 10.3. Profiling Autorun Kernels
11. Developing OpenCL™ Applications Using Intel® Code Builder for OpenCL™ x
11.1. Configuring the Intel® Code Builder for OpenCL™ Offline Compiler Plug-in for Microsoft Visual Studio 11.2. Configuring the Intel® Code Builder for OpenCL™ Offline Compiler Plug-in for Eclipse 11.3. Creating a Session in the Intel® Code Builder for OpenCL™ 11.4. Configuring a Session
12. Intel® FPGA SDK for OpenCL™ Standard Edition Advanced Features x
12.1. OpenCL Library 12.2. Kernel Attributes for Configuring Local and Private Memory Systems 12.3. Kernel Attributes for Reducing the Overhead on Hardware Usage 12.4. Kernel Replication Using the num_compute_units(X,Y,Z) Attribute
12.1. OpenCL Library x
12.1.1. Understanding RTL Modules and the OpenCL Pipeline 12.1.2. Packaging an OpenCL Helper Function File for an OpenCL Library 12.1.3. Packaging an RTL Component for an OpenCL Library 12.1.4. Verifying the RTL Modules 12.1.5. Packaging Multiple Object Files into a Library File 12.1.6. Specifying an OpenCL Library when Compiling an OpenCL Kernel 12.1.7. Using an OpenCL Library that Works with Simple Functions (Example 1) 12.1.8. Using an OpenCL Library that Works with External Memory (Example 2) 12.1.9. OpenCL Library Command-Line Options
12.1.1. Understanding RTL Modules and the OpenCL Pipeline x
12.1.1.1. Overview: Intel FPGA SDK for OpenCL Pipeline Approach 12.1.1.2. Integration of an RTL Module into the Intel FPGA SDK for OpenCL Pipeline 12.1.1.3. Stall-Free RTL 12.1.1.4. RTL Module Interfaces 12.1.1.5. Avalon Streaming (Avalon-ST) Interface 12.1.1.6. RTL Reset and Clock Signals 12.1.1.7. XML Syntax of an RTL Module 12.1.1.8. Interaction between RTL Module and External Memory 12.1.1.9. Order of Threads Entering an RTL Module 12.1.1.10. OpenCL C Model of an RTL Module 12.1.1.11. Potential Incompatibility between RTL Modules and Partial Reconfiguration
12.1.1.7. XML Syntax of an RTL Module x
12.1.1.7.1. XML Elements for ATTRIBUTES 12.1.1.7.2. XML Elements for INTERFACE 12.1.1.7.3. XML Elements for RESOURCES
12.1.3. Packaging an RTL Component for an OpenCL Library x
12.1.3.1. Restrictions and Limitations in RTL Support for the Intel® FPGA SDK for OpenCL™ Standard Edition Library Feature
12.2. Kernel Attributes for Configuring Local and Private Memory Systems x
12.2.1. Restrictions on the Usage of Variable-Specific Attributes
12.3. Kernel Attributes for Reducing the Overhead on Hardware Usage x
12.3.1. Hardware for Kernel Interface
12.3.1. Hardware for Kernel Interface x
12.3.1.1. Omit Hardware that Generates and Dispatches Kernel IDs 12.3.1.2. Omit Communication Hardware between the Host and the Kernel
12.4. Kernel Replication Using the num_compute_units(X,Y,Z) Attribute x
12.4.1. Customization of Replicated Kernels Using the get_compute_id() Function 12.4.2. Using Channels with Kernel Copies
A. Support Statuses of OpenCL Features x
A.1. Support Statuses of OpenCL 1.0 Features A.2. Support Statuses of OpenCL 1.2 Features A.3. Support Statuses of OpenCL 2.0 Features A.4. Intel® FPGA SDK for OpenCL™ Standard Edition Allocation Limits
A.1. Support Statuses of OpenCL 1.0 Features x
A.1.1. OpenCL1.0 C Programming Language Implementation A.1.2. OpenCL C Programming Language Restrictions A.1.3. Argument Types for Built-in Geometric Functions A.1.4. Numerical Compliance Implementation A.1.5. Image Addressing and Filtering Implementation A.1.6. Atomic Functions A.1.7. Embedded Profile Implementation
A.2. Support Statuses of OpenCL 1.2 Features x
A.2.1. OpenCL 1.2 Runtime Implementation A.2.2. OpenCL 1.2 C Programming Language Implementation
A.3. Support Statuses of OpenCL 2.0 Features x
A.3.1. OpenCL 2.0 Headers A.3.2. OpenCL 2.0 Runtime Implementation A.3.3. OpenCL 2.0 C Programming Language Restrictions for Pipes
1. Intel® FPGA SDK for OpenCL™ Standard Edition Overview
2. Intel® FPGA SDK for OpenCL™ Offline Compiler Kernel Compilation Flows
3. Obtaining General Information on Software, Compiler, and Custom Platform
4. Managing an FPGA Board
5. Structuring Your OpenCL Kernel
6. Designing Your Host Application
7. Compiling Your OpenCL Kernel
8. Emulating and Debugging Your OpenCL Kernel
9. Reviewing Your Kernel's report.html File
10. Profiling Your OpenCL Kernel
11. Developing OpenCL™ Applications Using Intel® Code Builder for OpenCL™
12. Intel® FPGA SDK for OpenCL™ Standard Edition Advanced Features
A. Support Statuses of OpenCL Features
B. Document Revision History of the Intel® FPGA SDK for OpenCL™ Standard Edition Programming Guide

5.4.5.8. Enforcing the Order of Channel Calls

To enforce the order of channel calls, introduce memory fence or barrier functions in your kernel program to control memory accesses. A memory fence function is necessary to create a control flow dependence between the channel accesses before and after the fence.

When the Intel® FPGA SDK for OpenCL™ Offline Compiler generates a compute unit, it does not always create instruction-level parallelism on all instructions that are independent of each other. As a result, channel read and write operations might not execute independently of each other even if there is no control or data dependence between them. When channel calls interact with each other, or when channels write data to external devices, deadlocks might occur.

For example, the code snippet below consists of a producer kernel and a consumer kernel. Channels c0 and c1 are unbuffered channels. The schedule of the channel read operations from c0 and c1 might occur in the reversed order as the channel write operations to c0 and c1. That is, the producer kernel writes to c0 but the consumer kernel might read from c1 first. This rescheduling of channel calls might cause a deadlock because the consumer kernel is reading from an empty channel. 

__kernel void producer (__global const uint * src,
                        const uint iterations)
{
    for (int i = 0; i < iterations; i++)
    {
        write_channel_intel(c0, src[2*i]);
        write_channel_intel(c1, src[2*i+1]);
    }
}

__kernel void consumer (__global uint * dst,
                        const uint iterations)
{
    for (int i = 0; i < iterations; i++)
    {
        /*During compilation, the AOC might reorder the way the consumer kernel 
        writes to memory to optimize memory access. Therefore, c1 might be read
        before c0, which is the reverse of what appears in code.*/  

        dst[2*i+1] = read_channel_intel(c0);
        dst[2*i] = read_channel_intel(c1);
    }
}
To prevent deadlocks from occurring by enforcing the order of channel calls, include memory fence functions (mem_fence) in your kernel.
Inserting the mem_fence call with each kernel's channel flag forces the sequential ordering of the write and read channel calls. The code snippet below shows the modified producer and consumer kernels: 
channel uint c0 __attribute__((depth(0)));
channel uint c1 __attribute__((depth(0)));

__kernel void producer (__global const uint * src,
                        const uint iterations)
{
    for (int i = 0; i < iterations; i++)
    {
        write_channel_intel(c0, src[2*i]);
        mem_fence(CLK_CHANNEL_MEM_FENCE);
        write_channel_intel(c1, src[2*i+1]);
    }
}

__kernel void consumer (__global uint * dst;
                        const uint iterations)
{
    for (int i = 0; i < iterations; i++)
    {
        dst[2*i+1] = read_channel_intel(c0);
        mem_fence(CLK_CHANNEL_MEM_FENCE);
        dst[2*i] = read_channel_intel(c1);
    }
}

In this example, mem_fence in the producer kernel ensures that the channel write operation to c0 occurs before that to c1. Similarly, mem_fence in the consumer kernel ensures that the channel read operation from c0 occurs before that from c1.


Section Content
Defining Memory Consistency Across Kernels When Using Channels

Level Two Title