Unexpected Atomic behaviour

Question

quaternionboy OP

Created Apr ’21

Replies 1

Boosts 0

Participants 2

With a threadsPerGrid of (10,1,1) and:

Code Block [[kernel]] void compute_shader (device atomic_int& incremental [[buffer(0)]],
                ushort lid [[thread_position_in_threadgroup]] ){
   
  threadgroup atomic_int local_atomic {0};
   
  atomic_fetch_add_explicit(&local_atomic, 1, memory_order_relaxed);
   
  threadgroup_barrier(mem_flags::mem_threadgroup);
   
  if(lid == 0) {
    int local_non_atomic = atomic_load_explicit(&local_atomic, memory_order_relaxed);
    atomic_fetch_add_explicit(&incremental, local_non_atomic, memory_order_relaxed);
  }
}

I expect:

Code Block 10
20
30
...

But I get:

Code Block 1125974026
1125974036
-2000908258
-832823256 ...

github.com/quaternionboy/Atomic

Answered by Graphics and Games Engineer in 669557022

There are known issues initializing local threadgroup vars when declared.

Please try initializing local_atomic in the following way:

Code Block threadgroup atomic_int local_atomic;
if (lid == 0) atomic_store_explicit(&local_atomic, 0, memory_order_relaxed);
threadgroup_barrier(mem_flags::mem_threadgroup);

Boost

Answer 1

Graphics and Games Engineer OP

Apple

Apr ’21

Accepted Answer

There are known issues initializing local threadgroup vars when declared.

Please try initializing local_atomic in the following way:

Code Block threadgroup atomic_int local_atomic;
if (lid == 0) atomic_store_explicit(&local_atomic, 0, memory_order_relaxed);
threadgroup_barrier(mem_flags::mem_threadgroup);

0