Visible to Intel only — GUID: ewa1417026580901
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: ewa1417026580901
Ixiasoft
4.5. Aligning a Struct with or without Padding
A properly aligned struct helps the Intel® FPGA SDK for OpenCL™ Offline Compiler generate the most efficient hardware.
A proper struct alignment means that the alignment can be evenly divided by the struct size.
Important: Ensure a 4-byte alignment for the data structures. struct alignments smaller than four bytes result in larger and slower hardware. Hardware efficiency increases with the increasing alignment. In the following example, the Pixel_s structure is only one-byte aligned but the Pixel structure is four-byte aligned due to the presence of a four-byte not_used integer:
typedef struct {
char r,g,b,alpha;
} Pixel_s;
typedef union {
Pixel_s p;
int not_used;
} Pixel;typedef struct {
char r,g,b,alpha;
} __attribute__((aligned(4))) Pixel;The offline compiler conforms with the ISO C standard that requires the alignment of a struct to satisfy all of the following criteria:
- The alignment must be an integer multiple of the lowest common multiple between the alignments of all struct members.
- The alignment must be a power of two.
You may set the struct alignment by including the aligned(N) attribute in your kernel code. Without an aligned attribute, the offline compiler determines the alignment of each struct in an array of struct based on the size of the struct. Consider the following example:
__kernel void test (struct mystruct* A,
struct mystruct* B)
{
A[get_global_id(0)] = B[get_global_id(0)];
}If the size of mystruct is 101 bytes, each load or store access is 1-byte aligned. If the size of mystruct is 128 bytes, each load or store access is 128-byte aligned, which generates the most efficient hardware.
When the struct fields are not aligned within the struct, the offline compiler inserts padding to align them. Inserting padding between struct fields affects hardware efficiency in the following manner:
- Increases the size of the struct
- Might affect the alignment
To prevent the offline compiler from inserting padding, include the packed attribute in your kernel code. The aforementioned ISO C standard applies when determining the alignment of a packed or unpacked struct. Consider the following example:
struct mystruct1
{
char a;
int b;
};The size of mystruct1 is 8 bytes. Therefore, the struct is 8-byte aligned, resulting in efficient accesses in the kernel. Now consider another example:
struct mystruct2
{
char a;
int b;
int c;
};The size of mystruct2 is 12 bytes and the struct is 4-byte aligned. Because the struct fields are padded and the struct is unaligned, accesses in the kernel are inefficient.
Following is an example of a struct that includes the packed attribute:
struct __attribute__((packed)) mystruct3
{
char a;
int b;
int c;
};The size of mystruct3 is 16 bytes. Because mystruct3 is aligned and there is no padding between struct fields, accesses in this kernel are more efficient than accesses in mystruct3.
To include both the aligned(N) and packed attributes in a struct, consider the following example:
struct __attribute__((packed)) __attribute__((aligned(16))) mystruct5
{
char a;
int b;
int c;
};The size of mystruct5 is 9 bytes. Because of the aligned(16) attribute, the struct is stored at 16-byte aligned addresses in an array. Because mystruct5 is 16-byte aligned and has no padding, accesses in this kernel is efficient.
For more information about struct alignment and the aligned(N) and packed attributes, refer to the following documents:
- Section 6.11.1 of the OpenCL Specification version 1.2
- Disabling Insertion of Data Structure Padding section of the Intel® FPGA SDK for OpenCL™ Programming Guide
- Specifying the Alignment of a Struct section of the Intel® FPGA SDK for OpenCL™ Programming Guide