Visible to Intel only — GUID: GUID-8ABC418E-2F96-4A0F-849E-92A37971D97A
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
See Also
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-8ABC418E-2F96-4A0F-849E-92A37971D97A
Using Specialization in Branching
You can improve the performance of both CPU and Intel® Graphics devices by converting the uniform conditions that are equal across all work-items into compile time branches, a techniques known as specialization.
The approach, which is sometimes referred as Uber-Shader in the pixel shader context, is to have a single kernel that implements all needed behaviors, and to let the host logic disable the paths that are not currently required. However, setting constants to branch on calculations wastes the device facilities, as the data is still being calculated before it is thrown away. Consider a preprocess approach instead, using #ifndef blocks.
Original kernel that uses constants to branch:
__kernel void foo(__constant int* src,
__global int* dst, unsigned char bFullFrame, unsigned char bAlpha)
{
…
if(bFullFrame)//uniform condition (equal for all work-items
{
…
if(bAlpha) //uniform condition
{
…
}
else
{
…
}
else
{
…
}
}
The same kernel with compile time branches:
__kernel void foo(__constant int* src,
__global int* dst)
{
…
#ifdef bFullFrame
{
…
#ifdef bAlpha
{
…
}
#else
{
…
}
#endif
#else
{
…
}
#endif
}
Also consider similar optimization for other constants.
Finally, avoid or minimize use of branching in short computations with using min, max, clamp or select built-ins instead of “if and else”.
Also, optimizing specifically for the OpenCL™ Intel Graphics device, ensure all conditionals are evaluated outside of code branches (for the CPU device it does not make any difference).
For example, the following code demonstrates conditional evaluation in the conditional blocks:
if(x && y || (z && functionCall(x, y, z))
{
// do something
}
else
{
// do something else
}
The following code demonstrates the conditional evaluation moved outside of the conditional blocks:
//improves compilation time for Intel® Graphics device
bool comparison = x && y || (z && functionCall(x, y, z));
if(comparison)
{
// do something
}
else
{
// do something else
}