// 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));
  }