Visible to Intel only — GUID: GUID-EB1F9802-5038-49D2-9888-497E6ECBD6D8
Legal Information
Getting Help and Support
Introduction
Coding for the Intel® Processor Graphics
Platform-Level Considerations
Application-Level Optimizations
Optimizing OpenCL™ Usage with Intel® Processor Graphics
Check-list for OpenCL™ Optimizations
Performance Debugging
Using Multiple OpenCL™ Devices
Coding for the Intel® CPU OpenCL™ Device
OpenCL™ Kernel Development for Intel® CPU OpenCL™ device
Mapping Memory Objects
Using Buffers and Images Appropriately
Using Floating Point for Calculations
Using Compiler Options for Optimizations
Using Built-In Functions
Loading and Storing Data in Greatest Chunks
Applying Shared Local Memory
Using Specialization in Branching
Considering native_ and half_ Versions of Math Built-Ins
Using the Restrict Qualifier for Kernel Arguments
Avoiding Handling Edge Conditions in Kernels
Using Shared Context for Multiple OpenCL™ Devices
Sharing Resources Efficiently
Synchronization Caveats
Writing to a Shared Resource
Partitioning the Work
Keeping Kernel Sources the Same
Basic Frequency Considerations
Eliminating Device Starvation
Limitations of Shared Context with Respect to Extensions
Why Optimizing Kernel Code Is Important?
Avoid Spurious Operations in Kernel Code
Perform Initialization in a Separate Task
Use Preprocessor for Constants
Use Signed Integer Data Types
Use Row-Wise Data Accesses
Tips for Auto-Vectorization
Local Memory Usage
Avoid Extracting Vector Components
Task-Parallel Programming Model Hints
Visible to Intel only — GUID: GUID-EB1F9802-5038-49D2-9888-497E6ECBD6D8
Avoid Spurious Operations in Kernel Code
Since every line in kernel code is executed many times, make sure you have no spurious instructions in your kernel code.
Spurious instructions are not always obvious. Consider the following kernel:
__kernel void foo(const __global int* data, const uint dataSize)
{
size_t tid = get_global_id(0);
size_t gridSize = get_global_size(0);
size_t workPerItem = dataSize / gridSize;
size_t myStart = tid * workPerItem;
for (size_t i = myStart; i < myStart + workPerItem; ++i)
{
//actual work
}
}
In this kernel, the for loop is used to reduce the number of work-items and the overhead of keeping them. In this example every work-item recalculates the limit to the index i, but this number is identical for all work-items. Since the sizes of the dataset and the NDRange dimensions are known before the kernel launch, consider calculating the amount of work per item on the host once, and then pass the result as constant parameter.
In addition, using size_t for indices makes vectorization of indexing arithmetic less efficient. To improve performance, when your index fits the 32-bit integer range, use int data type, as shown in the following example:
__kernel void foo(const __global int* data, const uint workPerItem)
{
int tid = get_global_id(0);
int gridSize = get_global_size(0);
//int workPerItem = dataSize / gridSize;
int myStart = tid * workPerItem;
for (int i = myStart; i < mystart + workPerItem; ++i)