source: CIVL/examples/cuda/cuda.cvh@ a8ca3d3

void __sync_init(int* in_barrier, int size) {
  barrier_size = size;
  for (int i=0; i<size; i++) in_barrier[i] = 0;
}
   
// model the synchronization of threads in the same block
void __syncthreads(int* in_barrier, int tid) {
  $atomic {
    in_barrier[tid] = 1; // I am in the barrier
    num_in_barrier++; // increment number in barrier
    if (num_in_barrier == barrier_size) { // I am last to enter
      for (int i=0; i<barrier_size; i++) 
        in_barrier[i] = 0; // release all
      num_in_barrier = 0; // now none are in barrier
    }
  }
  $when (in_barrier[tid] == 0); // wait till I am released
}