Visible to Intel only — GUID: GUID-6A6003F1-BAD5-4648-AB10-A8959D4A8B29
Legal Information
Getting Help and Support
Introduction
Coding for the Intel® Processor Graphics
Platform-Level Considerations
Application-Level Optimizations
Optimizing OpenCL™ Usage with Intel® Processor Graphics
Check-list for OpenCL™ Optimizations
Performance Debugging
Using Multiple OpenCL™ Devices
Coding for the Intel® CPU OpenCL™ Device
OpenCL™ Kernel Development for Intel® CPU OpenCL™ device
Mapping Memory Objects
Using Buffers and Images Appropriately
Using Floating Point for Calculations
Using Compiler Options for Optimizations
Using Built-In Functions
Loading and Storing Data in Greatest Chunks
Applying Shared Local Memory
Using Specialization in Branching
Considering native_ and half_ Versions of Math Built-Ins
Using the Restrict Qualifier for Kernel Arguments
Avoiding Handling Edge Conditions in Kernels
Using Shared Context for Multiple OpenCL™ Devices
Sharing Resources Efficiently
Synchronization Caveats
Writing to a Shared Resource
Partitioning the Work
Keeping Kernel Sources the Same
Basic Frequency Considerations
Eliminating Device Starvation
Limitations of Shared Context with Respect to Extensions
Why Optimizing Kernel Code Is Important?
Avoid Spurious Operations in Kernel Code
Perform Initialization in a Separate Task
Use Preprocessor for Constants
Use Signed Integer Data Types
Use Row-Wise Data Accesses
Tips for Auto-Vectorization
Local Memory Usage
Avoid Extracting Vector Components
Task-Parallel Programming Model Hints
Visible to Intel only — GUID: GUID-6A6003F1-BAD5-4648-AB10-A8959D4A8B29
Use Preprocessor for Constants
Consider the following kernel:
__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 runtime, after the kernel is issued for execution. However, you can use 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;
}
}
The capitalization indicates that EXPONENT is 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 the 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.
NOTE:
This approach requires recompiling the program every time the value of exponent_val changes. If you expect to change this value often, this approach is not advised. However, 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.