Visible to Intel only — GUID: GUID-EC4B2E93-8D7A-4456-B4B9-D34B3D596A4F
Legal Information
Getting Help and Support
Introduction
Check-list for OpenCL™ Optimizations
Tips and Tricks for Kernel Development
Application-Level Optimizations
Debugging OpenCL™ Kernels on Linux* OS
Performance Debugging with Intel® SDK for OpenCL™ Applications
Coding for the Intel® Architecture Processors
Why Optimizing Kernels Is Important?
Avoid Spurious Operations in Kernels
Avoid Handling Edge Conditions in Kernels
Use the Preprocessor for Constants
Prefer (32-bit) Signed Integer Data Types
Prefer Row-Wise Data Accesses
Use Built-In Functions
Avoid Extracting Vector Components
Task-Parallel Programming Model Hints
Common Mistakes in OpenCL™ Applications
Introduction for OpenCL™ Coding on Intel® Architecture Processors
Vectorization Basics for Intel® Architecture Processors
Vectorization: SIMD Processing Within a Work Group
Benefitting from Implicit Vectorization
Vectorizer Knobs
Targeting a Different CPU Architecture
Using Vector Data Types
Writing Kernels to Directly Target the Intel® Architecture Processors
Work-Group Size Considerations
Threading: Achieving Work-Group Level Parallelism
Efficient Data Layout
Using the Blocking Technique
Intel® Turbo Boost Technology Support
Global Memory Size
Visible to Intel only — GUID: GUID-EC4B2E93-8D7A-4456-B4B9-D34B3D596A4F
Prefer Row-Wise Data Accesses
OpenCL™ enables you to submit kernels on one-, two-, or three-dimensional index space. Consider using one-dimensional ranges for reasons of cache locality and saving index computations.
If two- or three-dimensional range naturally fits your data dimensions, try to keep work-items scanning along rows, not columns. For example, the following code is not optimized (it might trigger gathers instructions):
__kernel void smooth(__constant float* input,
uint image_width, uint image_height,
__global float* output)
{
int myX = get_global_id(1);
int myY = get_global_id(0);
int myPixel = myY * image_width + myX;
float data = input[myPixel];
…
}
In this code example, the image height is the first dimension and the image width is the second dimension. The resulting column-wise data access is inefficient, since Intel® OpenCL™ implementation initially iterates over the first dimension.
Below is more optimal version, because of more memory-friendly (sequential) access.
__kernel void smooth(__constant float* input,
uint image_width, uint image_height,
__global float* output)
{
int myX = get_global_id(0);
int myY = get_global_id(1);
int myPixel = myY * image_width + myX;
float data = input[myPixel];
…
}
In the example above, the first dimension is the image width and the second is the image height.
The same rule applies if each work-item calculates several elements. To optimize performance, make sure work-items read from consecutive memory addresses.
Finally, if you run two-dimensional NDRange, prefer the data access to be consecutive along dimension zero.