• 2019 Update 4
  • 03/20/2019
  • Public Content
Contents

Avoid Spurious Operations in Kernel Code

Since every line in kernel code is executed many times, make sure you have no spurious instructions in your kernel code.
Spurious instructions are not always obvious. Consider the following kernel:
__kernel void foo(const __global int* data, const uint dataSize) { size_t tid = get_global_id(0); size_t gridSize = get_global_size(0); size_t workPerItem = dataSize / gridSize; size_t myStart = tid * workPerItem; for (size_t i = myStart; i < myStart + workPerItem; ++i) { //actual work } }
In this kernel, the
for loop
is used to reduce the number of work-items and the overhead of keeping them. In this example every work-item recalculates the limit to the index
i
, but this number is identical for all work-items. Since the sizes of the dataset and the NDRange dimensions are known before the kernel launch, consider calculating the amount of work per item on the host once, and then pass the result as constant parameter.
In addition, using
size_t
for indices makes vectorization of indexing arithmetic less efficient. To improve performance, when your index fits the
32
-bit integer range, use
int
data type, as shown in the following example:
__kernel void foo(const __global int* data, const uint workPerItem) { int tid = get_global_id(0); int gridSize = get_global_size(0); //int workPerItem = dataSize / gridSize; int myStart = tid * workPerItem; for (int i = myStart; i < mystart + workPerItem; ++i)

Product and Performance Information

1

Intel's compilers may or may not optimize to the same degree for non-Intel microprocessors for optimizations that are not unique to Intel microprocessors. These optimizations include SSE2, SSE3, and SSSE3 instruction sets and other optimizations. Intel does not guarantee the availability, functionality, or effectiveness of any optimization on microprocessors not manufactured by Intel. Microprocessor-dependent optimizations in this product are intended for use with Intel microprocessors. Certain optimizations not specific to Intel microarchitecture are reserved for Intel microprocessors. Please refer to the applicable product User and Reference Guides for more information regarding the specific instruction sets covered by this notice.

Notice revision #20110804