6.3.1 Use the global ID instead of the loop counter

In OpenCL, you use kernels to perform the equivalent of loop iterations. This means that there is no loop counter to use in computations. The global ID of the work-item provides the equivalent of the loop counter. Use the global ID to perform any computations based on the loop counter.


You can include loops in OpenCL kernels, but they can only iterate over the data for that work-item, not the entire NDRange.

Simplified loop example

Example showing a simple loop in C that assigns the value of the loop counter to each array element.

Loop example in C:

The following loop fills an array with numbers.

void SetElements(void)
    int loop_count;
    int my_array[4096];
    for (loop_count = 0; loop_count < 4096; loop_count++)
        my_array[loop_count] = loop_count;
    printf("Final count %d\n", loop_count);

This loop is parallelizable because the loop elements are all independent. There is no main loop counter loop_count in the OpenCL kernel, so it is replaced by the global ID.

The equivalent code in an OpenCL kernel:
__kernel void example(__global int * restrict my_array)
    int id;
    id = get_global_id(0);
    my_array[id] = id;
Non-ConfidentialPDF file icon PDF version101574_0302_00_en
Copyright © 2019 Arm Limited or its affiliates. All rights reserved.