Visible to Intel only — GUID: GUID-94835769-D664-4A2D-B667-3DF3297AA4AC
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
See Also
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-94835769-D664-4A2D-B667-3DF3297AA4AC
Avoid Spurious Operations in Kernels
As 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 evident. 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. However, in this example, every work-item recalculates the amount of indices to iterate on while this number is identical for all work-items.
The size of the dataset and the NDRange dimensions is known before kernel launch. Calculate the amount of work per item on the host once and then pass the result as a constant parameter.
Using size_t for indices makes vectorization of indexing arithmetic less efficient. To improve performance, use the int data type, when your index fits the 32-bit integer range. Consider 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)="" …=""
See Also
OpenCL™ 1.2 Specification at https://www.khronos.org/registry/OpenCL/specs/opencl-1.2.pdf
Overview Presentations of the OpenCL™ Standard at https://www.khronos.org/registry/OpenCL