Compare and Swap

See: compare and swap

int atomicCAS(int* address, int compare, int val);

Its semantics is like this

__device__ int atomicCAS(int* address, int compare, int val) {
  __lock(address) {
    int old = *address;
    if (old == compare) { *address = val; }
    return old;
  }
}

Example: Implement atomic add with atomicCAS

__device__ int atomicAdd(int* address, int val) {
 int old = *address;
 int assumed;
 do {
   // If the value at *address didn't change, imcrement it
   assumed = old;
   old = atomicCAS(address, assumed, val + assumed);
 }
 return old;
}