What is best way to implement an atomic add with clamp?

OpenCL
1

I need to increment a uint by a calculated amount “inc” between 0 and UINT_MAX. These calculations occur in parallel, and several threads may acess the same value.

I can do this with atom_add(pmap, inc), but there is a distinct possiblity that the value may wrap.

Is there an efficient way to prevent the roll over by clamping the value?

Currently I increment the value, then retrieve the value and test that it is > inc.
This would be simple if atom_add() return the new value, but atom_add returns the “old” value.

As it is, this two step process is not guaranteed to be correct.

Is there a simple way to do what I need?

Thanks

2

Use atom_cmpxchg() in a loop until it succeeds.


__kernel void foo(volatile __global int* bar)
{
    int old, new;

    do
    {
        old = *bar;
        new = anything_you_want(old);
    }
    while(atomic_cmpxchg(bar, old, new) != old);
}