Visible to Intel only — GUID: GUID-CC7B9348-C8F8-4E99-B846-5FE79B0B7828
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
See Also
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-CC7B9348-C8F8-4E99-B846-5FE79B0B7828
Use the Preprocessor for Constants
Consider the following example of using the preprocessor for constants:
__kernel void exponentor(__global int* data, const uint exponent)
{
int tid = get_global_id(0);
int base = data[tid];
for (int i = 1; i < exponent; ++i)
{
data[tid] *= base;
}
}
The number of iterations for the inner for loop is determined at run time, after the kernel is issued for execution. However, you can use the OpenCL™ dynamic compilation feature to ensure the exponent is known at kernel compile time, which is done during the host run time. In this case, the kernel appears as follows:
__kernel void exponentor(__global int* data)
{
int tid = get_global_id(0);
int base = data[tid];
for (int i = 1; i < EXPONENT; ++i)
{
data[tid] *= base;
}
}
Capitalization indicates that exponent became a preprocessor macro.
The original version of the host code passes exponent_val through kernel arguments as follows:
clSetKernelArg(kernel, 1, exponent_val);
The updated version uses a compilation step:
sprintf(buildOptions, “-DEXPONENT=%u”, exponent_val); clBuildProgram(program, <...>, buildOptions, <...>);
Thus, the value of exponent is passed during preprocessing of the kernel code. Besides saving stack space used by the kernel, this also enables the compiler to perform optimizations, such as loop unrolling or elimination. This technique is often useful for transferring parameters like image dimensions to video-processing kernels, where the value is only known at host run time, but does not change once it is defined.
NOTE:
NOTE: This approach requires recompiling the program every time the value of exponent_val changes. Prefer the variable-passing approach if you expect to change this value often.