static void
activate(struct sml_control *control)
{
unsigned int old;
/* lock; all updates so far must be acquired */
old = fetch_and(acquire, &control->state, ~INACTIVE_FLAG);
if (IS_ACTIVE(old)) {
mutex_lock(&control->inactive_wait_lock);
pthread_cleanup_push(cleanup_mutex_unlock,
&control->inactive_wait_lock);
while ((old = fetch_and(acquire, &control->state,
~INACTIVE_FLAG),
IS_ACTIVE(old)))
cond_wait(&control->inactive_wait_cond,
&control->inactive_wait_lock);
pthread_cleanup_pop(1);
}
assert(IS_ACTIVE(load_relaxed(&control->state)));
}
__host__ __device__
int_type operator&=(int_type val) volatile
{
return fetch_and(val) + val;
}
T operator &= (T value)
{
return fetch_and(value) & value;
}