Visible to Intel only — GUID: GUID-2B97945D-6BEA-4D4A-84E4-A7A396821001
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
See Also
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-2B97945D-6BEA-4D4A-84E4-A7A396821001
Prefer (32-bit) Signed Integer Data Types
Many image-processing kernels operate on uchar input. To avoid overflows, those eight-bit input values are typically converted and processed as 16- or 32-bit integer values. Use signed data types (shorts and ints) in both cases to enable the compiler to utilize a larger set of SIMD instructions.
Using size_t, which is another unsigned type for indices, makes the vectorization of indexing arithmetic less efficient. To improve performance, use the int data type for work-item parameters and loop counter, when your index fits the 32-bit integer range. Consider the following example:
__kernel void foo(__constant int* data, const uint workPerItem)
{
int tid = get_global_id(0);
int gridSize = get_global_size(0);
for (int i = myStart; i <mystart += workperitem; ++i) …
Also when the compiler generates the scatter or gather instructions on non-consecutive memory accesses, it needs to safely cast to the int32 since gather and scatter instructions use the int32 indices. Explicit casting of the indices to the int32 in a kernel simplifies the compiler job.