atomicExchange() and atomicCompareExchangeWeak() failed on Metal

[ciaochaos]

@group(0)
@binding(0)
var<storage, read_write> v_indices: array<atomic<u32>>; // this is used as both input and output for convenience

// The Collatz Conjecture states that for any integer n:
// If n is even, n = n/2
// If n is odd, n = 3n+1
// And repeat this process for each new n, you will always eventually reach 1.
// Though the conjecture has not been proven, no counterexample has ever been found.
// This function returns how many times this recurrence needs to be applied to reach 1.
fn collatz_iterations(n_base: u32) -> u32{
    var n: u32 = n_base;
    var i: u32 = 0u;
    loop {
        if (n <= 1u) {
            break;
        }
        if (n % 2u == 0u) {
            n = n / 2u;
        }
        else {
            // Overflow? (i.e. 3*n + 1 > 0xffffffffu?)
            if (n >= 1431655765u) {   // 0x55555555u
                return 4294967295u;   // 0xffffffffu
            }

            n = 3u * n + 1u;
        }
        i = i + 1u;
    }
    return i;
}

@compute
@workgroup_size(1)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
    atomicStore(&v_indices[global_id.x], collatz_iterations(v_indices[global_id.x]));
    atomicExchange(&v_indices[global_id.x], collatz_iterations(v_indices[global_id.x]));
}

Caused by:
    In Device::create_compute_pipeline
    Internal error: Metal: program_source:67:49: error: cannot take the address of an rvalue of type 'metal::uint' (aka 'unsigned int')
    uint _e14 = metal::atomic_exchange_explicit(&uint(_e8) < 1 + (_buffer_sizes.size0 - 0 - 4) / 4 ? v_indices[_e8] : DefaultConstructible(), _e13, metal::memory_order_relaxed);

macOS 12.3, M1