Intel® FPGA SDK for OpenCL™ Standard Edition: Best Practices Guide

Download PDF
ID 683176
Date 9/24/2018
Public
Document Table of Contents
Document Table of Contents x
1. Introduction to Standard Edition Best Practices Guide 2. Reviewing Your Kernel's report.html File 3. OpenCL Kernel Design Best Practices 4. Profiling Your Kernel to Identify Performance Bottlenecks 5. Strategies for Improving Single Work-Item Kernel Performance 6. Strategies for Improving NDRange Kernel Data Processing Efficiency 7. Strategies for Improving Memory Access Efficiency 8. Strategies for Optimizing FPGA Area Usage A. Additional Information
1. Introduction to Standard Edition Best Practices Guide x
1.1. FPGA Overview 1.2. Pipelines 1.3. Single Work-Item Kernel versus NDRange Kernel 1.4. Multi-Threaded Host Application
2. Reviewing Your Kernel's report.html File x
2.1. High Level Design Report Layout 2.2. Reviewing the Report Summary 2.3. Reviewing Loop Information 2.4. Reviewing Area Information 2.5. Verifying Information on Memory Replication and Stalls 2.6. Optimizing an OpenCL Design Example Based on Information in the HTML Report 2.7. HTML Report: Area Report Messages 2.8. HTML Report: Kernel Design Concepts
2.3. Reviewing Loop Information x
2.3.1. Loop Analysis Report of an OpenCL Design Example 2.3.2. Changing the Memory Access Pattern Example 2.3.3. Reducing the Area Consumed by Nested Loops Using loop_coalesce
2.4. Reviewing Area Information x
2.4.1. Area Analysis by Source 2.4.2. Area Analysis of System
2.5. Verifying Information on Memory Replication and Stalls x
2.5.1. Features of the System Viewer 2.5.2. Features of the Kernel Memory Viewer
2.7. HTML Report: Area Report Messages x
2.7.1. Area Report Message for Board Interface 2.7.2. Area Report Message for Function Overhead 2.7.3. Area Report Message for State 2.7.4. Area Report Message for Feedback 2.7.5. Area Report Message for Constant Memory 2.7.6. Area Report Messages for Private Variable Storage
2.8. HTML Report: Kernel Design Concepts x
2.8.1. Kernels 2.8.2. Global Memory Interconnect 2.8.3. Local Memory 2.8.4. Nested Loops 2.8.5. Loops in a Single Work-Item Kernel 2.8.6. Channels 2.8.7. Load-Store Units
3. OpenCL Kernel Design Best Practices x
3.1. Transferring Data Via Channels or OpenCL Pipes 3.2. Unrolling Loops 3.3. Optimizing Floating-Point Operations 3.4. Allocating Aligned Memory 3.5. Aligning a Struct with or without Padding 3.6. Maintaining Similar Structures for Vector Type Elements 3.7. Avoiding Pointer Aliasing 3.8. Avoid Expensive Functions 3.9. Avoiding Work-Item ID-Dependent Backward Branching
3.1. Transferring Data Via Channels or OpenCL Pipes x
3.1.1. Characteristics of Channels and Pipes Default Behavior Concurrent Execution of Multiple OpenCL Kernels Implicit Kernel Synchronization Data Persistence Across Invocations Imposed Work-Item Order 3.1.2. Execution Order for Channels and Pipes 3.1.3. Optimizing Buffer Inference for Channels or Pipes 3.1.4. Best Practices for Channels and Pipes
3.3. Optimizing Floating-Point Operations x
3.3.1. Floating-Point versus Fixed-Point Representations
4. Profiling Your Kernel to Identify Performance Bottlenecks x
4.1. Best Practices 4.2. GUI 4.3. Interpreting the Profiling Information 4.4. Limitations
4.2. GUI x
4.2.1. Source Code Tab 4.2.2. Kernel Execution Tab 4.2.3. Autorun Captures Tab
4.2.1. Source Code Tab x
4.2.1.1. Tool Tip Options
4.3. Interpreting the Profiling Information x
4.3.1. Stall, Occupancy, Bandwidth 4.3.2. Activity 4.3.3. Cache Hit 4.3.4. Profiler Analyses of Example OpenCL Design Scenarios 4.3.5. Autorun Profiler Data
4.3.1. Stall, Occupancy, Bandwidth x
4.3.1.1. Stalling Channels
4.3.4. Profiler Analyses of Example OpenCL Design Scenarios x
4.3.4.1. High Stall Percentage 4.3.4.2. Low Occupancy Percentage 4.3.4.3. Low Bandwidth Efficiency 4.3.4.4. High Stall and High Occupancy Percentages 4.3.4.5. No Stalls, Low Occupancy Percentage, and Low Bandwidth Efficiency 4.3.4.6. No Stalls, High Occupancy Percentage, and Low Bandwidth Efficiency 4.3.4.7. Stalling Channels 4.3.4.8. High Stall and Low Occupancy Percentages
5. Strategies for Improving Single Work-Item Kernel Performance x
5.1. Addressing Single Work-Item Kernel Dependencies Based on Optimization Report Feedback 5.2. Removing Loop-Carried Dependencies Caused by Accesses to Memory Arrays 5.3. Good Design Practices for Single Work-Item Kernel
5.1. Addressing Single Work-Item Kernel Dependencies Based on Optimization Report Feedback x
5.1.1. Removing Loop-Carried Dependency 5.1.2. Relaxing Loop-Carried Dependency 5.1.3. Simplifying Loop-Carried Dependency 5.1.4. Transferring Loop-Carried Dependency to Local Memory 5.1.5. Removing Loop-Carried Dependency by Inferring Shift Registers
6. Strategies for Improving NDRange Kernel Data Processing Efficiency x
6.1. Specifying a Maximum Work-Group Size or a Required Work-Group Size 6.2. Kernel Vectorization 6.3. Multiple Compute Units 6.4. Combination of Compute Unit Replication and Kernel SIMD Vectorization 6.5. Reviewing Kernel Properties and Loop Unroll Status in the HTML Report
6.2. Kernel Vectorization x
6.2.1. Static Memory Coalescing
6.3. Multiple Compute Units x
6.3.1. Compute Unit Replication versus Kernel SIMD Vectorization
7. Strategies for Improving Memory Access Efficiency x
7.1. General Guidelines on Optimizing Memory Accesses 7.2. Optimize Global Memory Accesses 7.3. Performing Kernel Computations Using Constant, Local or Private Memory 7.4. Improving Kernel Performance by Banking the Local Memory 7.5. Optimizing Accesses to Local Memory by Controlling the Memory Replication Factor 7.6. Minimizing the Memory Dependencies for Loop Pipelining
7.2. Optimize Global Memory Accesses x
7.2.1. Contiguous Memory Accesses 7.2.2. Manual Partitioning of Global Memory
7.2.2. Manual Partitioning of Global Memory x
7.2.2.1. Heterogeneous Memory Buffers
7.3. Performing Kernel Computations Using Constant, Local or Private Memory x
7.3.1. Constant Cache Memory 7.3.2. Preloading Data to Local Memory 7.3.3. Storing Variables and Arrays in Private Memory
7.4. Improving Kernel Performance by Banking the Local Memory x
7.4.1. Optimizing the Geometric Configuration of Local Memory Banks Based on Array Index
8. Strategies for Optimizing FPGA Area Usage x
8.1. Compilation Considerations 8.2. Board Variant Selection Considerations 8.3. Memory Access Considerations 8.4. Arithmetic Operation Considerations 8.5. Data Type Selection Considerations
A. Additional Information x
A.1. Document Revision History for the Standard Edition Best Practices Guide
1. Introduction to Standard Edition Best Practices Guide
2. Reviewing Your Kernel's report.html File
3. OpenCL Kernel Design Best Practices
4. Profiling Your Kernel to Identify Performance Bottlenecks
5. Strategies for Improving Single Work-Item Kernel Performance
6. Strategies for Improving NDRange Kernel Data Processing Efficiency
7. Strategies for Improving Memory Access Efficiency
8. Strategies for Optimizing FPGA Area Usage
A. Additional Information

3.1.1. Characteristics of Channels and Pipes

To implement channels or pipes in your OpenCL™ kernel program, keep in mind their respective characteristics that are specific to the .

Default Behavior

The default behavior of channels is blocking. The default behavior of pipes is nonblocking.

Concurrent Execution of Multiple OpenCL Kernels

You can execute multiple OpenCL kernels concurrently. To enable concurrent execution, modify the host code to instantiate multiple command queues. Each concurrently executing kernel is associated with a separate command queue.

Important:

Pipe-specific considerations:

The OpenCL pipe modifications outlined in Ensuring Compatibility with Other OpenCL SDKs in the Standard Edition Programming Guide allow you to run your kernel on the SDK. However, they do not maximize the kernel throughput. The OpenCL Specification version 2.0 requires that pipe writes occur before pipe reads so that the kernel is not reading from an empty pipe. As a result, the kernels cannot execute concurrently. Because the supports concurrent execution, you can modify your host application and kernel program to take advantage of this capability. The modifications increase the throughput of your application; however, you can no longer port your kernel to another SDK. Despite this limitation, the modifications are minimal, and it does not require much effort to maintain both types of code.

To enable concurrent execution of kernels containing pipes, replace the depth attribute in your kernel code with the blocking attribute (that is, __attribute__((blocking))). The blocking attribute introduces a blocking behavior in the read_pipe and write_pipe function calls. The call site blocks kernel execution until the other end of the pipe is ready.

If you add both the blocking attribute and the depth attribute to your kernel, the read_pipe calls will only block when the pipe is empty, and the write_pipe calls will only block when the pipe is full. Blocking behavior causes an implicit synchronization between the kernels, which forces the kernels to run in lock step with each other.

Implicit Kernel Synchronization

Synchronize the kernels implicitly via blocking channel calls or blocking pipe calls. Consider the following examples:

Table 5.  Blocking Channel and Pipe Calls for Kernel Synchronization
Kernel with Blocking Channel Call Kernel with Blocking Pipe Call 
channel int c0; 

__kernel
void producer (__global int * in_buf) 
{ 
  for (int i = 0; i < 10; i++) 
  {   
    write_channel_intel (c0, in_buf[i]);
  } 
}

__kernel
void consumer (__global int * ret_buf) 
{                 
  for (int i = 0; i < 10; i++) 
  { 
    ret_buf[i] = read_channel_intel(c0);
  }
} 
__kernel
void producer (__global int * in_buf,
  write_only pipe int __attribute__
  ((blocking)) c0) 
{
  for (int i = 0; i < 10; i++) 
{
    write_pipe (c0, &in_buf[i]);
  }
}

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

You can synchronize the kernels such that a producer kernel writes data and a consumer kernel reads the data during each loop iteration. If the write_channel_intel or write_pipe call in producer does not write any data, consumer blocks and waits at the read_channel_intel or read_pipe call until producer sends valid data, and vice versa.

Data Persistence Across Invocations

After the write_channel_intel call writes data to a channel or the write_pipe call writes data to a pipe, the data is persistent across work-groups and NDRange invocations. Data that a work-item writes to a channel or a pipe remains in that channel or pipe until another work-item reads from it. In addition, the order of data in a channel or a pipe is equivalent to the sequence of write operations to that channel or pipe, and the order is independent of the work-item that performs the write operation.

For example, if multiple work-items try to access a channel or a pipe simultaneously, only one work-item can access it. The write_channel_intel call or write_pipe call writes the particular work-item data, called DATAX, to the channel or pipe, respectively. Similarly, the first work-item to access the channel or pipe reads DATAX from it. This sequential order of read and write operations makes channels and pipes an effective way to share data between kernels.

Imposed Work-Item Order

The SDK imposes a work-item order to maintain the consistency of the read and write operations for a channel or a pipe.

Related Information
Level Two Title