1

The following code snippet in a Metal compute kernel suggests that atomic_fetch_add_explicit does not have an atomic read-modify-write within threadgroup memory.

The value of i is not unique within the threadgroup as I expect it to be.

Am I using it wrong??

threadgroup atomic_int index;  
atomic_store_explicit( &index, 0, memory_order_relaxed );  
threadgroup_barrier( mem_flags::mem_none );  
int i = atomic_fetch_add_explicit( &index, 1, memory_order_relaxed );  

1 Answer 1

2

This is indeed correct and functions atomically as expected. The error was in my code verifying the uniqueness of i.

Your Answer

By clicking “Post Your Answer”, you agree to our terms of service and acknowledge you have read our privacy policy.

Not the answer you're looking for? Browse other questions tagged or ask your own question.