Visible to Intel only — GUID: uir1520273525684
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: uir1520273525684
Ixiasoft
10.4. On-Chip Storage of Reused Data
For Intel® Stratix® 10 designs, the presence of cached LSUs prevent certain optimizations. Intel® recommends that you avoid inferring caching LSUs.
The unoptimized code below creates a cached burst-coalesced LSU that consumes more resources and disables other optimizations:
kernel void cached (global int * restrict in,
global int * restrict out) {
int i = get_global_id(0);
int idx = out[i];
int cached_value = in[idx]; // Burst-coalesced cached LSU
out[i] = cached_value;
}To prevent the caching, mark the memory pointer as volatile, as shown in the optimized code below:
kernel void not_cached (global volatile int * restrict in,
global int * restrict out) {
int i = get_global_id(0);
int idx = out[i];
int not_cached_value = in[idx];
out[i] = not_cached_value;
}For more information about optimizing load-store units, refer to the Load-Store Units section.
Intel® also recommends that you use on-chip storage to achieve optimal results. Note that compared to non- Intel® Stratix® 10 devices, Intel® Stratix® 10 has a larger M20K to ALM ratio, allowing you to create larger local memory systems.
The unoptimized code below has a function that receives a pointer from the array in global memory. In this case, the offline compiler modifies the array and then stores it back to memory. Then, in a subsequent iteration of the outer loop, the array is reused.
void cached_array(int N, int M, int BUF_SIZE,
global float2 * global_buf)
{
for (uint i = 0; i< N; i++) {
float2 data_array[BUF_SIZE];
for (uint j= 0; j < M; j++) {
data_array[i] = global_buf [j]; //Load value
... //do work to modify data_array
global_buf[j] = data_array[i];
}
}
}To prevent unnecessary global memory accesses, define a private array in the kernel to declare the array on chip. The function accesses the private array instead of the one declared on chip. As a result, the array receives on-chip local memory storage. Accessing this on-chip local memory does not require accesses to global memory.
Optimized code:
void local_array(int N, int M, int BUF_SIZE,
global float2 * global_buf)
{
float2 local_buf[BUF_SIZE];
populate_buf(local_buf, global_buf);
for (uint i = 0; i< N; i++) {
float2 data_array[BUF_SIZE];
for (uint j= 0; j < M; j++) {
data_array[i] = local_buf[j];//Load value
... //do work to modify data_array
local_buf[j] = data_array[i];
}
}
}
Related Information