As 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 evident. Consider the following
__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)
In this kernel, the
loop is used to reduce the number
of work-items and the overhead of keeping them. However, in this example,
every work-item recalculates the amount of indices to iterate on while
this number is identical for all work-items.
The size of the dataset and the NDRange dimensions is known before kernel
launch. Calculate the amount of work per item on the host once and then
pass the result as a constant parameter.
for indices makes vectorization of indexing
arithmetic less efficient. To improve performance, use the
data type, when your index fits the 32-bit integer range. Consider the
__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)="" …=""
OpenCL™ 1.2 Specification at
Overview Presentations of the OpenCL™ Standard at