is the initial value of a threadgroup atomic_uint zero?

Question

Is the initial value of a threadgroup atomic_uint zero? I didn't see anything in the MSL spec. Or do I need to do something like the following to initialize it to zero?

threadgroup atomic_uint flags;
if(localIndex == 0) {
    atomic_store_explicit(&flags, 0, memory_order_relaxed);
}

threadgroup_barrier(mem_flags::mem_threadgroup);

Answer 1

Local threadgroup memory is not initialized by default, so you will have to do something like you've done:

compute_kernel(texture2d<float, access::read> inTexture [[texture(0)]],
                  volatile device atomic_uint* outBins [[buffer(0)]],
                  uint2 gid [[thread_position_in_grid]],
                  uint2 threadIndex [[thread_position_in_threadgroup]]) 
{
    threadgroup atomic_uint flags;
    if (all(threadIndex == 0)) { // Only do the initialization on one thread
        atomic_store_explicit(&flags + i, 0, memory_order_relaxed);
    }
    threadgroup_barrier(mem_flags::mem_threadgroup); // Make all threads wait until initialization is done
    
    // flags is now initialized and usable
}