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

5.1.3. Simplifying Loop-Carried Dependency

In cases where you cannot remove or relax the loop-carried dependency in your kernel, you might be able to simplify the dependency to improve single work-item kernel performance.

Consider the following kernel example:

 1 #define N 128
 2 #define NUM_CH 3
 3 
 4 channel uchar CH_DATA_IN[NUM_CH];
 5 channel uchar CH_DATA_OUT;
 6 
 7 __kernel void unoptimized()
 8 {
 9   unsigned storage = 0;
10   unsigned num_bytes = 0;
11 
12   for (unsigned i = 0; i < N; i++) {
13 
14     #pragma unroll
15     for (unsigned j = 0; j < NUM_CH; j++) {
16       if (num_bytes < NUM_CH) {
17         bool valid = false;
18         uchar data_in = read_channel_nb_intel(CH_DATA_IN[j], &valid);
19         if (valid) {
20           storage <<= 8;
21           storage |= data_in;
22           num_bytes++;
23         }
24       }
25     }
26 
27     if (num_bytes >= 1) {
28       num_bytes -= 1;
29       uchar data_out = storage >> (num_bytes*8);
30       write_channel_intel(CH_DATA_OUT, data_out);
31     }
32   }
33 }

This kernel reads one byte of data from three input channels in a nonblocking fashion. It then writes the data one byte at a time to an output channel. It uses the variable storage to store up to 4 bytes of data, and uses the variable num_bytes to keep track of how many bytes are stored in storage. If storage has space available, then the kernel reads a byte of data from one of the channels and stores it in the least significant byte of storage.

The optimization report below indicates that there is a loop-carried dependency on the variable num_bytes: 

==================================================================================
Kernel: unoptimized
==================================================================================
The kernel is compiled for single work-item execution.

Loop Report:

 + Loop "Block1" (file unoptimized3.cl line 12)
 | Pipelined with successive iterations launched every 7 cycles due to:
 |
 |     Data dependency on variable num_bytes  (file unoptimized3.cl line 10)
 |     Largest Critical Path Contributors:
 |      16%: Integer Compare Operation  (file unoptimized3.cl line 16)
 |      16%: Integer Compare Operation  (file unoptimized3.cl line 16)
 |      16%: Integer Compare Operation  (file unoptimized3.cl line 16)
 |       7%: Integer Compare Operation  (file unoptimized3.cl line 27)
 |       6%: Add Operation  (file unoptimized3.cl line 10, line 22, line 28)
 |       6%: Add Operation  (file unoptimized3.cl line 10, line 22, line 28)
 |       6%: Add Operation  (file unoptimized3.cl line 10, line 22, line 28)
 |       3%: Non-Blocking Channel Read Operation  (file unoptimized3.cl line 18)
 |       3%: Non-Blocking Channel Read Operation  (file unoptimized3.cl line 18)
 |       3%: Non-Blocking Channel Read Operation  (file unoptimized3.cl line 18)
 |
 |
 |-+ Fully unrolled loop (file unoptimized3.cl line 15)
     Loop was fully unrolled due to "#pragma unroll" annotation.

The computation path of num_bytes is as follows:

  1. Comparison on line 16 (if (num_bytes < NUM_CH)).
  2. Computation of variable valid by the nonblocking channel read operation on line 18 (uchar data_in = read_channel_nb_intel(CH_DATA_IN[j], &valid)) for the comparison on line 19.
  3. Addition on line 22 (num_bytes++).
  4. Comparison on line 27 (if (num_bytes >= 1)).
  5. Subtraction on line 28 (num_bytes -= 1).

Because of the unroll pragma on line 14, the unrolls the loop, causing the comparisons and additions in the loop body to replicate three times. The optimization report shows that the comparisons are the most expensive operations on the computation path of num_bytes, followed by the additions on line 22.

To simplify the loop-carried dependency on num_bytes, consider restructuring the application to perform the following tasks:

  1. Ensure that the kernel reads from the channels only if there is enough space available in storage, in the event that all channel read operations return data (that is, there is at least 3 bytes of empty space in storage).
    Setting this condition simplifies the computation path of the variable num_bytes by reducing the number of comparisons.
  2. Increase the size of storage from 4 bytes to 8 bytes to satisfy the 3-byte space threshold more easily.
Below is the restructured kernel optimized: 
 1 #define N 128
 2 #define NUM_CH 3
 3 
 4 channel uchar CH_DATA_IN[NUM_CH];
 5 channel uchar CH_DATA_OUT;
 6 
 7 __kernel void optimized()
 8 {
 9   // Change storage to 64 bits
10   ulong storage = 0;
11   unsigned num_bytes = 0;
12 
13   for (unsigned i = 0; i < N; i++) {
14 
15     // Ensure that we have enough space if we read from ALL channels
16     if (num_bytes <= (8-NUM_CH)) {
17       #pragma unroll
18       for (unsigned j = 0; j < NUM_CH; j++) {
19         bool valid = false;
20         uchar data_in = read_channel_nb_intel(CH_DATA_IN[j], &valid);
21         if (valid) {
22           storage <<= 8;
23           storage |= data_in;
24           num_bytes++;
25         }
26       }
27     }
28 
29     if (num_bytes >= 1) {
30       num_bytes -= 1;
31       uchar data_out = storage >> (num_bytes*8);
32       write_channel_intel(CH_DATA_OUT, data_out);
33     }
34   }
35 }

An optimization report similar to the one below indicates the successful simplification of the loop-carried dependency on the variable num_bytes: 

==================================================================================
Kernel: optimized
==================================================================================
The kernel is compiled for single work-item execution.

Loop Report:

 + Loop "Block1" (file optimized3.cl line 13)
 | Pipelined well. Successive iterations are launched every cycle.
 |
 |
 |-+ Fully unrolled loop (file optimized3.cl line 18)
     Loop was fully unrolled due to "#pragma unroll" annotation.
Level Two Title