Visible to Intel only — GUID: ewa1457384630094
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: ewa1457384630094
Ixiasoft
8.5. Optimizing Accesses to Local Memory by Controlling the Memory Replication Factor
The memory replication factor is the number of M20K memory blocks that your design uses to implement the local memory system. To control the memory replication factor, use the max_replicates kernel attribute in your OpenCL™ kernel.
Intel® 's M20K memory blocks have two physical ports. The number of logical ports that are available in each M20K block depends on the degree of pumping. Pumping is a measure of the clock frequency of the M20K blocks relative to the rest of the design.
Consider the following code example where the singlepump attribute is applied to a local memory system, lmem, which has three read accesses and one write access. The singlepump attribute indicates that the M20K blocks runs at the same frequency as the rest of the design.
__kernel void three_copies(int raddr, int waddr) {
int __attribute__((memory,
numbanks(1),
singlepump,
max_replicates(3)))
lmem[16];
lmem[waddr] = lmem[raddr] + lmem[raddr + 1] + lmem[raddr + 2];
// do something with lmem
}
Figure 85. Accesses to Single-Pumped M20K Memory Blocks
The compiler creates an arbitration-free network, as shown in Figure 85. Each single-pumped M20K block has two logical ports available. Each write port in the local memory system must be connected to all M20K blocks that your design uses to implement the memory system. Each read port in the local memory system must be connected to one M20K block. Because of these connection constraints, there must be three M20K blocks to implement the specified number of ports in lmem.
Note: If you change max_replicates(3) to max_replicates(1), you observes one M20K block with arbitration between the three reads.
If you include the doublepump kernel attribute in your local variable declaration, you specify that the M20K memory blocks runs at double the frequency as the rest of the design.
__kernel void three_copies(int raddr, int waddr) {
int __attribute__((memory,
numbanks(1),
doublepump))
lmem[16];
lmem[waddr] = lmem[raddr] + lmem[raddr + 1] + lmem[raddr + 2];
// do something with lmem
}
Figure 86. Accesses to Double-Pumped M20K Memory Blocks
Each double-pumped M20K block has four logical ports available. As such, there only needs to be one M20K block to implement three read ports and one write port in lmem.
Attention:
- Double pumping the memory increases resource overhead. Use the doublepump kernel attribute only if it results in actual M20K savings, improves performance, or both.
- Stores must be connected to every replicate. Hence, if there are more than three stores, the memory is not replicated. Local memory replication works well with single store.
- Because the entire memory system is replicated, you might observe potentially large M20K memory blocks.
Related Information