Visible to Intel only — GUID: ewa1397680164688
Ixiasoft
1. Introduction to Intel® FPGA SDK for OpenCL™ Pro Edition Best Practices Guide
2. Reviewing Your Kernel's report.html File
3. OpenCL Kernel Design Concepts
4. OpenCL Kernel Design Best Practices
5. Profiling Your Kernel to Identify Performance Bottlenecks
6. Strategies for Improving Single Work-Item Kernel Performance
7. Strategies for Improving NDRange Kernel Data Processing Efficiency
8. Strategies for Improving Memory Access Efficiency
9. Strategies for Optimizing FPGA Area Usage
10. Strategies for Optimizing Intel® Stratix® 10 OpenCL Designs
11. Strategies for Improving Performance in Your Host Application
12. Intel® FPGA SDK for OpenCL™ Pro Edition Best Practices Guide Archives
A. Document Revision History for the Intel® FPGA SDK for OpenCL™ Pro Edition Best Practices Guide
2.1. High-Level Design Report Layout
2.2. Reviewing the Summary Report
2.3. Viewing Throughput Bottlenecks in the Design
2.4. Using Views
2.5. Analyzing Throughput
2.6. Reviewing Area Information
2.7. Optimizing an OpenCL Design Example Based on Information in the HTML Report
2.8. Accessing HLD FPGA Reports in JSON Format
4.1. Transferring Data Via Intel® FPGA SDK for OpenCL™ Channels or OpenCL Pipes
4.2. Unrolling Loops
4.3. Optimizing Floating-Point Operations
4.4. Allocating Aligned Memory
4.5. Aligning a Struct with or without Padding
4.6. Maintaining Similar Structures for Vector Type Elements
4.7. Avoiding Pointer Aliasing
4.8. Avoid Expensive Functions
4.9. Avoiding Work-Item ID-Dependent Backward Branching
5.1. Best Practices for Profiling Your Kernel
5.2. Instrumenting the Kernel Pipeline with Performance Counters (-profile)
5.3. Obtaining Profiling Data During Runtime
5.4. Reducing Area Resource Use While Profiling
5.5. Temporal Performance Collection
5.6. Performance Data Types
5.7. Interpreting the Profiling Information
5.8. Profiler Analyses of Example OpenCL Design Scenarios
5.9. Intel® FPGA Dynamic Profiler for OpenCL™ Limitations
8.1. General Guidelines on Optimizing Memory Accesses
8.2. Optimize Global Memory Accesses
8.3. Performing Kernel Computations Using Constant, Local or Private Memory
8.4. Improving Kernel Performance by Banking the Local Memory
8.5. Optimizing Accesses to Local Memory by Controlling the Memory Replication Factor
8.6. Minimizing the Memory Dependencies for Loop Pipelining
8.7. Static Memory Coalescing
Visible to Intel only — GUID: ewa1397680164688
Ixiasoft
9.3. Memory Access Considerations
Intel® recommends kernel programming strategies that can improve memory access efficiency and reduce area usage of your OpenCL™ kernel.
- Minimize the number of access points to external memory.
If possible, structure your kernel such that it reads its input from one location, processes the data internally, and then writes the output to another location.
- Instead of relying on local or global memory accesses, structure your kernel as a single work-item with shift register inference whenever possible.
- Instead of creating a kernel that writes data to external memory and a kernel that reads data from external memory, implement the Intel® FPGA SDK for OpenCL™ channels extension between the kernels for direct data transfer.
- If your OpenCL application includes many separate constant data accesses, declare the corresponding pointers using __constant instead of __global const. Declaration using __global const creates a private cache for each load or store operation. However, declaration using __constant creates a single constant cache on the chip only.
CAUTION:If your kernel targets a Cyclone® V device (for example, Cyclone V SoC), declaring __constant pointer kernel arguments might degrade FPGA performance.
- If your kernel passes a small number of constant arguments, pass them as values instead of pointers to global memory.
For example, instead of passing __constant int * coef and then dereferencing coef with index 0 to 10, pass coef as a value (int16 coef). If coef was the only __constant pointer argument, passing it as a value eliminates the constant cache and the corresponding load and store operations completely.
- Conditionally shifting large shift registers inside pipelined loops leads to the creation of inefficient hardware. For example, the following kernel consumes more resources when the if (K > 5) condition is present:
#define SHIFT_REG_LEN 1024 __kernel void bad_shift_reg (__global int * restrict src, __global int * restrict dst, int K) { float shift_reg[SHIFT_REG_LEN]; int sum = 0; for (unsigned i = 0; i < K; i++) { sum += shift_reg[0]; shift_reg[SHIFT_REG_LEN-1] = src[i]; // This condition will cause sever area bloat. if (K > 5) { #pragma unroll for (int m = 0; m < SHIFT_REG_LEN-1 ; m++) { shift_reg[m] = shift_reg[m + 1]; } } dst[i] = sum; } }
Attention: Conditionally accessing a shift register does not degrade hardware efficiency. If it is necessary to implement conditional shifting of a large shift register in your kernel, consider modifying your code so that it uses local memory.