Is atomic_fetch_add_explicit atomic within threadgroup memory?

Question

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

score 2 · Answer 1 · answered Mar 01 '16 at 11:27

2

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

answered Mar 01 '16 at 11:27

Dickson

211
1
11

Is atomic_fetch_add_explicit atomic within threadgroup memory?

1 Answers1