OpenCL: synchronisation between kernels

Question

So I am trying to do something like this:

In my host code I have a loop

for(int i = 0; i < params.maxIters; i++){
    launch1DKernel(kernel…)
}

and in my kernel I thave

__kernel k(__global int * i, …){
    somearray[i * size + offset] = some_val;
    if(get_global_id(0) == 0){
        (*i) = (*i) + 1;
    }
}

This seems not a good a pproach beause I want to gurantee that the all the work-items of the same iteration shares the i. And the above code does not gurantee that the i won't be increnmented by the next iteration's work item 0. Am I right about this?

Answer 1

In OpenCL, global memory consistency is only guaranteed at the start of a kernel invocations. So, if one work-item writes a value to global memory, you cannot guarantee that this new value will be visible to work-items that are in other work-groups. This new value will be visible to all work-items by the next kernel invocation however, which in your example would be the next iteration.

So, yes, the *i value will be incremented by the next iteration's work-item 0, but no, you cannot share that value with other work-items within the same iteration (unless they are in the same work-group, in which case you need to use a barrier).

There's a more detailed description of the OpenCL memory model in the specification, which is worth a read to get a good handle on this stuff.

Answer 2

You need to use the atomic_* functions to ensure that global memory will be what you expect it to be.

Be warned though these are very slow.