Results 1 to 2 of 2

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

  1. #1
    Junior Member
    Join Date
    Jan 2012
    Posts
    1

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

    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. #2
    Senior Member
    Join Date
    May 2010
    Location
    Toronto, Canada
    Posts
    845

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

    Use atom_cmpxchg() in a loop until it succeeds.

    Code :
    __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);
    }
    Disclaimer: Employee of Qualcomm Canada. Any opinions expressed here are personal and do not necessarily reflect the views of my employer. LinkedIn profile.

Similar Threads

  1. Atomic Function
    By fici_sg in forum Suggestions for next release
    Replies: 0
    Last Post: 01-03-2013, 08:02 AM

Posting Permissions

  • You may not post new threads
  • You may not post replies
  • You may not post attachments
  • You may not edit your posts
  •