Developer Guide

FPGA Optimization Guide for Intel® oneAPI Toolkits

Download PDF
ID 767853
Date 12/16/2022
Public

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

Document Table of Contents
Document Table of Contents x
FPGA Optimization Guide for Intel® oneAPI Toolkits
FPGA Optimization Guide for Intel® oneAPI Toolkits x
Introduction To FPGA Design Concepts Analyze Your Design Optimize Your Design FPGA Optimization Flags, Attributes, Pragmas, and Extensions Quick Reference Additional Information Document Revision History for the FPGA Optimization Guide for Intel® oneAPI Toolkits Notices and Disclaimers
Introduction To FPGA Design Concepts x
FPGA Architecture Overview Concepts of FPGA Hardware Design Methods of Hardware Design How Source Code Becomes a Custom Hardware Datapath Scheduling Mapping Parallelism Models to FPGA Hardware Memory Types
FPGA Architecture Overview x
Adaptive Logic Module (ALM) Lookup Table (LUT) Register Digital Signal Processing (DSP) Block Random Access Memory (RAM) Blocks
Concepts of FPGA Hardware Design x
Maximum Frequency (fMAX) Latency Pipelining Throughput Datapath Control Path Occupancy
How Source Code Becomes a Custom Hardware Datapath x
Mapping Source Code Instructions to Hardware Mapping Arrays and Their Accesses to Hardware
Scheduling x
Dynamic Scheduling Clustering the Datapath Handshaking Between Clusters
Mapping Parallelism Models to FPGA Hardware x
Data Parallelism Task Parallelism
Data Parallelism x
Executing Independent Operations Simultaneously Pipelining
Memory Types x
Kernel Memory Global Memory
Analyze Your Design x
Analyze the FPGA Early Image Analyze the FPGA Image
Analyze the FPGA Early Image x
Review the FPGA Optimization Report Access HLD FPGA Reports in JSON Format
Review the FPGA Optimization Report x
Loop Analysis Bottlenecks Viewer Area Estimates System Viewer Kernel Memory Viewer Schedule Viewer
Analyze the FPGA Image x
Quartus (Static) Summary Intel® FPGA Dynamic Profiler for DPC++ System-level Profiling Using the Intercept Layer for OpenCL* Applications
Quartus (Static) Summary x
Timing Failures
Intel® FPGA Dynamic Profiler for DPC++ x
Measure Kernel Performance Instrument the Kernel Pipeline with Performance Counters (<span class='codeph'>-Xsprofile</span>) Obtain Profiling Data During Runtime Reduce Area Resource Use While Profiling Profiler Analyses of Example SYCL* Design Scenarios Limitations
Obtain Profiling Data During Runtime x
Invoke the Profiler Runtime Wrapper to Obtain Profiling Data Use Intel® VTune™ Profiler
Use Intel® VTune™ Profiler x
Interpret Performance Counter Data
System-level Profiling Using the Intercept Layer for OpenCL* Applications x
Set Up the Intercept Layer for OpenCL* Applications
Optimize Your Design x
Throughput Resource Use
Throughput x
Single Work-item Kernels NDRange Kernels Memory Accesses Pipes Host
Single Work-item Kernels x
Single Work-item Kernel Design Guidelines Loops Single-Cycle Floating-Point Accumulator for Single Work-Item Kernels
Loops x
Refactor the Loop-Carried Data Dependency Relax Loop-Carried Dependency Transfer Loop-Carried Dependency to Local Memory Minimize the Memory Dependencies for Loop Pipelining Unroll Loops Fuse Loops to Reduce Overhead and Improve Performance Optimize Loops With Loop Speculation Remove Loop Bottlenecks Shannonization to Improve FMAX/II Optimize Inner Loop Throughput Improve Loop Performance by Caching On-Chip Memory
Single-Cycle Floating-Point Accumulator for Single Work-Item Kernels x
Strategies for Inferring the Accumulator
Memory Accesses x
Load-Store Units Global Memory Accesses Optimization Perform Kernel Computations Using Local or Private Memory Local and Private Memory Accesses Optimization Annotating Unified Shared Memory Pointers Zero-Copy Memory Access Additional Recommendations
Load-Store Units x
Load-Store Unit Styles Load-Store Unit Modifiers Load-Store Unit Controls
Global Memory Accesses Optimization x
Global Memory Bandwidth Use Calculation Manual Partition of Global Memory Partitioning Buffers Across Different Memory Types (Heterogeneous Memory) Partitioning Buffers Across Memory Channels of the Same Memory Type Ignoring Dependencies Between Accessor Arguments Contiguous Memory Accesses Static Memory Coalescing
Host x
Multi-Threaded Host Application Utilizing Hardware Kernel Invocation Queue Double Buffering Host Utilizing Kernel Invocation Queue N-Way Buffering to Overlap Kernel Execution Prepinning Memory Simple Host-Device Streaming Buffered Host-Device Streaming
Double Buffering Host Utilizing Kernel Invocation Queue x
Applying Double-Buffering Using the Intercept Layer for OpenCL* Applications
Resource Use x
Data Types and Operations Kernel Variable Accesses
Data Types and Operations x
Optimize Floating-point Operation Avoid Expensive Functions Variable-Precision Integer and Floating-Point Support
Variable-Precision Integer and Floating-Point Support x
Advantages and Limitations of Arbitrary Precision Data Types Declare and Use the AC Data Types
Declare and Use the AC Data Types x
Declare the <span class='codeph'>ac_int</span> Data Type Declare the <span class='codeph'>ac_fixed</span> Data Type Declare the <span class='codeph'>ac_complex</span> Data Type Declare the <span class='codeph'>ap_float</span> Data Type
Declare the <span class='codeph'>ap_float</span> Data Type x
Conversion Rules for <span class='codeph'>ap_float</span> Operations with Explicit Precision Controls Comparison Operators Additional <span class='codeph'>ap_float</span> Functions Additional Data Types Provided by the <span class='codeph'>ap_float.hpp</span> Header File Quality of Results and the ap_float Data Type
FPGA Optimization Flags, Attributes, Pragmas, and Extensions x
Optimization Flags Kernel Attributes Kernel Controls Kernel Variables Memory Attributes Loop Directives Floating-Point Pragmas Latency Controls (Beta) System of Tasks Extension (<span class='codeph'>task_sequence</span>)
Optimization Flags x
Specify Schedule FMAX Target for Kernels (<span class='codeph'>-Xsclock=<clock target>) Disable Burst-Interleaving of Global Memory (<span class='codeph'>-Xsno-interleaving=<global_memory_type></span>) Force Ring Interconnect for Global Memory (<span class='codeph'>-Xsglobal-ring</span>) Force a Single Store Ring to Reduce Area (<span class='codeph'>-Xsforce-single-store-ring</span>) Force Fewer Read Data Reorder Units to Reduce Area (<span class='codeph'>-Xsnum-reorder</span>) Disable Hardware Kernel Invocation Queue (<span class='codeph'>-Xsno-hardware-kernel-invocation-queue</span>) Modify the Handshaking Protocol Between Clusters (<span class='codeph'>-Xshyper-optimized-handshaking</span>) Disable Automatic Fusion of Loops (<span class='codeph'>-Xsdisable-auto-loop-fusion</span>) Fuse Adjacent Loops With Unequal Trip Counts (<span class='codeph'>-Xsenable-unequal-tc-fusion</span>) Pipeline Loops in Non-task Kernels (<span class='codeph'>-Xsauto-pipeline</span>) Control Semantics of Floating-Point Operations (<span class='codeph'>-fp-model=<var><value></var> </span>) Modify the Rounding Mode of Floating-point Operations (<span class='codeph'>-Xsrounding=<rounding_type></span>) Global Control of Exit FIFO Latency of Stall-free Clusters (<span class='codeph'>-Xssfc-exit-fifo-type=<var><value></var> </span>) Enable the Read-Only Cache for Read-Only Accessors (<span class='codeph'>-Xsread-only-cache-size=<var><N></var>)</span> Control Hardware Implementation of the Supported Data Types and Math Operations (<span class='codeph'>-Xsdsp-mode=<var><option></var> </span>)
Kernel Attributes x
Specify Schedule FMAX Target for Kernels Specify a Workgroup Size Specify Number of SIMD WorkItems Omit Hardware that Generates and Dispatches Kernel IDs Omit Hardware to Support the <span class='codeph'>no_global_work_offset</span> Attribute in <span class='codeph'>parallel_for</span> Kernels Reduce Kernel Area and Latency
Kernel Controls x
Pipes Extension
Pipes Extension x
Key Properties of a Pipe Accessing Pipes The <span class='codeph'>pipe</span> Class and its Use Example Code Using Blocking Inter-Kernel Pipes Example Code Using Non-Blocking Inter-Kernel Pipes I/O Pipes Characteristics of Pipes Restrictions of Pipes Guidelines for Designing Pipes Pipe and Atomic Fence
Loop Directives x
<span class='codeph'>disable_loop_pipelining</span> Attribute <span class='codeph'>initiation_interval</span> Attribute <span class='codeph'>ivdep</span> Attribute <span class='codeph'>loop_coalesce</span> Attribute <span class='codeph'>max_concurrency</span> Attribute <span class='codeph'>max_interleaving</span> Attribute <span class='codeph'>speculated_iterations</span> Attribute <span class='codeph'>unroll</span> Pragma Loop Fuse Functions and <span class='codeph'>nofusion</span> Attribute
System of Tasks Extension (<span class='codeph'>task_sequence</span>) x
Task Functions <span class='codeph'>task_sequence</span> Use Cases
Quick Reference x
Algorithmic C Data Types Floating Point Pragmas FPGA Accessor Properties FPGA Extensions FPGA Kernel Attributes FPGA Local Memory Function Latency Control Properties (Beta) FPGA LSU Controls FPGA Loop Directives FPGA Memory Attributes FPGA Optimization Flags Pipe API <span class='codeph'>task_sequence</span> Template Parameters and Function APIs
FPGA Optimization Guide for Intel® oneAPI Toolkits

The <span class='codeph'>pipe</span> Class and its Use

The pipe API exposed by the FPGA implementation is equivalent to the following class declaration: 

template <class name,
          class dataT,
          size_t min_capacity = 0>
class pipe {
public:
  // Blocking
  static dataT read();
  static void write(dataT data);
  // Non-blocking
  static dataT read(bool &success_code);
  static void write(dataT data, bool &success_code);
}

The following table describes the template parameters:

Template Parameters

Parameter

Description

name

The type that is the basis of a pipe identification. It is typically a user-defined class, in a user namespace. Forward declaration of the type is enough, and the type need not be defined.

dataT

The type of data packet contained within a pipe. This is the data type that is read during a successful pipe read() operation, or written during a successful pipe write() operation. The type must have a standard layout and be trivially copyable.

min_capacity

User-defined minimum number of words (in units of dataT) that the pipe must be able to store without any being read out. The compiler may create a pipe with a larger capacity due to performance considerations.

The pipe class exposes static methods for writing a data word to a pipe and reading a data word from a pipe. The reads and writes can be blocking or non-blocking depending on the parameters you pass to the read() and/or write() function.

NOTE:

A data word in this context is the data type that the pipe contains (dataT pipe template argument).

Example Code Using Blocking Inter-Kernel Pipes

When writing code with SYCL* pipes, use of the C++ type alias mechanism (using) is highly encouraged to avoid errors where slightly different pipe types inadvertently lead to unique pipes. The following code sample shows how to use pipes with blocking accessors to transfer data between two kernels: 

#include <sycl/sycl.hpp>
using namespace sycl;
constexpr int N = 3;
// Specialize a pipe type
using my_pipe = ext::intel::pipe<class some_pipe, int, 8>;
void producer(const std::array<int, N> &src) {
  queue q;
  // Launch the producer kernel
  buffer<int> src_buf = {std::begin(src), std::end(src)};
  q.submit([&](handler &cgh) {
    // Get read access to src array
    accessor rd_src_buf(src_buf, cgh, read_only);
    cgh.single_task<class producer>([=]() {
      for (int i = 0; i < N; i++) {
        // Blocking write an int to the pipe
        my_pipe::write(rd_src_buf[i]);
      }
    });
  });
}
void consumer(std::array<int, N> &dst) {
  queue q;
  // Launch the consumer kernel
  buffer<int> dst_buf = {std::begin(dst), std::end(dst)};
  q.submit([&](handler &cgh) {
    // Get write access to dst array
    accessor wr_dst_buf(dst_buf, cgh, write_only); 
    cgh.single_task<class consumer>([=]() {
      for (int i = 0; i < N; i++) {
        // Blocking read an int from the pipe
        wr_dst_buf[i] = my_pipe::read();
      }
    });
  });
}

The pipe data packet is of type int and the pipe has a depth of 8, as specified by the template parameters of my_pipe type. The pipe read() call blocks only when the pipe is empty, and the pipe write() call blocks only when the pipe is full.

NOTE:

The SYCL specification does not guarantee concurrent kernel execution. However, the Intel® oneAPI DPC++/C++ Compiler supports concurrent execution of kernels. Hence, you can modify your host application and kernel program to take advantage of this capability. The modifications increase the throughput of your application.

Example Code Using Non-Blocking Inter-Kernel Pipes

The code samples (Sample 1 and Sample 2) in this section illustrate how to use pipes with non-blocking writes and reads to transfer data between two concurrently running kernels: 

//Sample 1
// The Producer kernel reads data from a SYCL buffer and writes it to
// a pipe. This transfers the input data from the host to the Consumer kernel
// that is running concurrently.
event Producer(queue &q, buffer<int, 1> &input_buffer) {
  std::cout << "Enqueuing producer...\n";

  auto e = q.submit([&](handler &h) {
    accessor input_accessor(input_buffer, h, read_only);
    size_t num_elements = input_buffer.size();

    h.single_task<ProducerTutorial>([=]() {
      for (size_t i = 0; i < num_elements; ++i) {
        ProducerToConsumerPipe::write(input_accessor[i]);
      }
    });
  });

  return e;
}

For both pipes, the data packet is of type int. The pipes are different because the first template parameter is different. The non-blocking pipe write() and read() calls do not block. They respectively return a boolean value that indicates whether data is written successfully to the pipe (that is, the pipe is not full) or if the data is read successfully from the pipe (that is, the pipe is not empty).

Perform non-blocking pipe writes to facilitate applications where writes to a full FIFO buffer should not cause the kernel to stall until a slot in the FIFO buffer becomes free. Consider a scenario where your application has one data producer with two identical workers that consume the data. Assume the time each worker takes to process a message varies depending on the contents of the data. In this case, there might be a situation where one worker is busy while the other is free. A non-blocking write can facilitate work distribution such that both workers are busy. Like a non-blocking write, perform non-blocking reads to facilitate applications where data is not always available, and other operations need not wait for the data to become available.

NOTE:

You can mix blocking and non-blocking accessors for writing or reading data to or from pipes. For example, you can write data to a pipe using a blocking pipe write() call and read it from the other end using a non-blocking pipe read() call, and vice versa. 

//Sample 2
#include <sycl/sycl.hpp>
using namespace sycl;
constexpr size_t N = 16;
// Specialize the two pipe types, differentiated based on their first template
// parameter
using pipe1 = ext::intel::pipe<class some_pipe, int>;
using pipe2 = ext::intel::pipe<class other_pipe, int>;
// the producer kernels are not shown
void consumer(const std::array<int, N> &dst) {
  queue q;
  // Launch the consumer kernel
  buffer<int> dst_buf = {std::begin(dst), std::end(dst)};
  q.submit([&](handler &cgh) {
    // Get write access to src array
    accessor wr_dst_buf(dst_buf, cgh, write_only);
    cgh.single_task<class consumer>([=]() {
      int = 0;
      while (i < N) {
        bool valid0 = false, valid1 = false;
        auto data0 = pipe1::read(valid0);
        auto data1 = pipe2::read(valid1);
        if (valid0) {
          wr_dst_buf[i++] = process(data0);
        }
        if (valid1) {
          wr_dst_buf[i++] = process(data1);
        }
      }
    });
  });
}
NOTE:

For additional information, refer to FPGA tutorial samples "Pipe Array" and "Pipes" listed in the Intel® oneAPI Samples Browser on Linux* or Windows*, or access the code samples in GitHub.

Parent topic: Pipes Extension
Level Two Title