Visible to Intel only — GUID: GUID-F92FD3B8-375C-45EE-8E20-2CE268436FCF
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
Infinite Execution of OpenCL™ Kernels
See Also
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-F92FD3B8-375C-45EE-8E20-2CE268436FCF
Common Mistakes in OpenCL™ Applications
This topic describes several cases of OpenCL™ applications undefined behavior that you might encounter if you do not follow the OpenCL™ specification.
Infinite Execution of OpenCL™ Kernels
- Independent forward progress between work items within a workgroup
Compilation of kernels assuming independent forward progress of the work-items may produce non-terminating code.
Consider the following example:
__kernel void kern(__local uint *flag) { size_t id = get_local_id(0); if (id==0) { flag[0] = 0; } barrier( CLK_LOCAL_MEM_FENCE ); while ( flag[0]<1 ) { if (id==0) { atomic_inc( &flag[0] ); } } }According to the OpenCL 2.0 specification, section 3.2.2:
"In the absence of work-group functions (e.g. a barrier), work-items within a workgroup may be serialized. In the presence of work-group functions, work-items within a workgroup may be serialized before any given work-group function, between dynamically encountered pairs of work-group functions and between a work-group function and the end of the kernel."
For example, if the order of work-items execution is 3, 2, 1, 0 and the work items are serialized, the first executed work item with id = 3 will never exit the loop.
The section 3.2.2 also states:
"The work-items within a single work-group execute concurrently but not necessarily in parallel (i.e. they are not guaranteed to make independent forward progress)."
Redesign your code to comply with specification and not to rely on independent forward progress of work items within a work group.
- barrier() divergence
Compilation of kernels containing divergent barriers may produce non-terminating code.
A branch is considered divergent when some work items within a work group take it and other do not take it.
A simplified code looks as follows:
__kernel void kern() { size_t lid = get_local_id(0); if (lid %2) { // any condition that would force only some work items to take the branch ... barrier(); } }According to the OpenCL™ 1.2 specification, section 3.4.3:
"Note that the work-group barrier must be encountered by all workitems of a work-group executing the kernel or by none at all".
It is harder to follow this rule if you place a barrier inside a loop. Redesign your code to avoid hitting the barrier only by a subset of work items within a work group.
See Also
OpenCL 1.2 Specification at https://www.khronos.org/registry/OpenCL/specs/opencl-1.2.pdf
OpenCL 2.0 Specification at https://www.khronos.org/registry/OpenCL/specs/opencl-2.0.pdf