Structures with OpenCL 1.0

Hi Everyone,

For my implementation I should used OpenCL 1.0.

I am working on a kernel in which a pixel is compared to neighbouring pixels by weight(pixel values).

Based on certain threshold value like in the range of 200 to 50, for the pixels falling in the range, I would like to store the pixel coordinates values in a dataarray of structs of a two dimensional point system containing the x and y coordinate(like Point2D).

The array of struct has a buffer allocated along with predefined size.

My question,

How can I create an indexed array of a struct independent to the global refereneces like array[0],array[1]…independent of the index value from get_global()?


A simple way (although not the most efficient if most of the pixels are to be returned) is to use an atomic counter.
Your device has to support the cl_khr_global_int32_base_atomics extension.

#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable

__kernel myKernel(__global Point2D *foundCoords, __global int *count)

if (myCondition)
int idx = atom_inc(count);
foundCoords[idx] = pointToAdd;

When the kernel has ended, read the value of count to get the number of pixels found.

Hi utnapishtim,

Thank you very much for the suggestion! The solution you gave works fine.

Mostly there are many positives getting returned and matches to the original algorithm.


Atomic operations can be very slow on GPUs (before Kepler for NVidia for instance), so we try to avoid them if they occur frequently in a kernel.

Recent AMD GPUs have hardware atomic counters dedicated to this kind of operation. If the extension “cl_ext_atomic_counters_32” is supported, you can use counter32_t instead of int for the counter.

A way to avoid the use of atomics is that each work-item work on a row of pixels instead of a single pixel. Pre-allocate a buffer of Point2D at the size of the image, and a buffer of ints which will contain the number of detected coordinates for each row:

__kernel myKernel(__global Point2D *foundCoords, __global int *counts)
int count = 0;
int row = get_global_id(0);
__global Point2D *foundRowCoords = foundCoords + row * rowSize;

for (int i = 0; i < rowSize; i++)
    if (myCondition)
        foundRowCoords[count] = pointToAdd;

counts[row] = count;


The coordinates won’t be adjacent in foundCoords. You’ll have to analyze the results row by row, or you may compact the coordinates depending on your algorithm (which can be done in-place here).