Visible to Intel only — GUID: GUID-50C00560-3371-43F9-B428-9F4EAC8819F8
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
See Also
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-50C00560-3371-43F9-B428-9F4EAC8819F8
Loading and Storing Data in Greatest Chunks
“Saturating” the available memory bandwidth is very important. Bytes data types actually load integer data types (DWORDS), but also trigger instructions to pack and unpack data. Using (u)int4 or float4 for buffers saves a lot of compute, even if you unpack data manually afterward. In other words, you should avoid using uchar4 or char4. See the example below:
__kernel void amp (__constant uchar4* src, __global uchar4* dst)
…
uint4 tempSrc = convert_uint4(src[offset]);//Load one RGBA8 pixel
…
//some processing
…
dst[offset] = convert_uchar4(tempDst);
}
Consider data accesses by using int4 data type:
__kernel void amp (__constant uint4* src, __global uint4* dst)
…
uint4 tempSrc = src[offset]; // Load 4 RGBA8 pixels
…
//some processing in uint4
uint r0 = (tempSrc.x & 0xff);//Red component of 1st pixel
uint r1 = (tempSrc.y & 0xff);//Red component of 2nd pixel
…
tempSrc.x >>= 8;
tempSrc.y >>= 8;
…
tempSrc.x >>= 8;
tempSrc.y >>= 8;
…
uint a0 = (tempSrc.x & 0xff);// Alpha component of 1st pixel
uint a1 = (tempSrc.y & 0xff);// Alpha component of 2nd pixel
//any calculations on the individual components
…
uint4 final = 0; // repack them:
final.x = (r0) | ((g0) << 8) | ((b0) << 16) | ((a0) << 16);//first pixel
final.y = (r1) | ((g1) << 8) | ((b1) << 16) | ((a1) << 16);//second pixel
…
dst[offset] = final;
}
NOTE:
The global size is 1/4th of the original size in the second example above.
If your kernel operates on floating-point data, consider using float4 data type, which gets four times as much data in one load. It also helps to ensure that the kernel has enough work to do, amortizing the work-item scheduling overheads.
For the CPU device this optimization is equivalent to explicit (manual) vectorization, see the “Using Vector Data Types” section for more information.
Accessing data in greater chunks can improve the Intel® Graphics device data throughput, but it might slightly reduce the CPU device performance as also explained in the “Using Vector Data Types” section.