// Locking

void lock_p(global int * lock, int id)
{
    while (atomic_cmpxchg(lock, -1, id) != id) {};
}

void lock_v(global int * lock, int id)
{
    atomic_xchg(lock, -1);
}

#define LOCK_P(lock, id) while (atomic_cmpxchg(lock, -1, id) != id) {}
#define LOCK_V(lock, id) while (atomic_cmpxchg(lock, id, -1) != -1) {}

kernel void locktest(
    global int *lock,
    global int *counter)
{
    LOCK_P(lock, get_global_id(0));
    *counter = *counter + 1;
    LOCK_V(lock, get_global_id(0));
}