wgpu/naga/tests/out/msl/bounds-check-zero-atomic.msl

// language: metal1.0
#include <metal_stdlib>
#include <simd/simd.h>

using metal::uint;
struct DefaultConstructible {
    template<typename T>
    operator T() && {
        return T {};
    }
};

struct _mslBufferSizes {
    uint size0;
};

struct type_1 {
    metal::atomic_uint inner[10];
};
typedef metal::atomic_uint type_2[1];
struct Globals {
    metal::atomic_uint a;
    type_1 b;
    type_2 c;
};

uint fetch_add_atomic(
    device Globals& globals,
    constant _mslBufferSizes& _buffer_sizes
) {
    uint _e3 = metal::atomic_fetch_add_explicit(&globals.a, 1u, metal::memory_order_relaxed);
    return _e3;
}

uint fetch_add_atomic_static_sized_array(
    int i,
    device Globals& globals,
    constant _mslBufferSizes& _buffer_sizes
) {
    uint _e5 = uint(i) < 10 ? metal::atomic_fetch_add_explicit(&globals.b.inner[i], 1u, metal::memory_order_relaxed) : DefaultConstructible();
    return _e5;
}

uint fetch_add_atomic_dynamic_sized_array(
    int i_1,
    device Globals& globals,
    constant _mslBufferSizes& _buffer_sizes
) {
    uint _e5 = uint(i_1) < 1 + (_buffer_sizes.size0 - 44 - 4) / 4 ? metal::atomic_fetch_add_explicit(&globals.c[i_1], 1u, metal::memory_order_relaxed) : DefaultConstructible();
    return _e5;
}

uint exchange_atomic(
    device Globals& globals,
    constant _mslBufferSizes& _buffer_sizes
) {
    uint _e3 = metal::atomic_exchange_explicit(&globals.a, 1u, metal::memory_order_relaxed);
    return _e3;
}

uint exchange_atomic_static_sized_array(
    int i_2,
    device Globals& globals,
    constant _mslBufferSizes& _buffer_sizes
) {
    uint _e5 = uint(i_2) < 10 ? metal::atomic_exchange_explicit(&globals.b.inner[i_2], 1u, metal::memory_order_relaxed) : DefaultConstructible();
    return _e5;
}

uint exchange_atomic_dynamic_sized_array(
    int i_3,
    device Globals& globals,
    constant _mslBufferSizes& _buffer_sizes
) {
    uint _e5 = uint(i_3) < 1 + (_buffer_sizes.size0 - 44 - 4) / 4 ? metal::atomic_exchange_explicit(&globals.c[i_3], 1u, metal::memory_order_relaxed) : DefaultConstructible();
    return _e5;
}

uint fetch_add_atomic_dynamic_sized_array_static_index(
    device Globals& globals,
    constant _mslBufferSizes& _buffer_sizes
) {
    uint _e4 = uint(1000) < 1 + (_buffer_sizes.size0 - 44 - 4) / 4 ? metal::atomic_fetch_add_explicit(&globals.c[1000], 1u, metal::memory_order_relaxed) : DefaultConstructible();
    return _e4;
}

uint exchange_atomic_dynamic_sized_array_static_index(
    device Globals& globals,
    constant _mslBufferSizes& _buffer_sizes
) {
    uint _e4 = uint(1000) < 1 + (_buffer_sizes.size0 - 44 - 4) / 4 ? metal::atomic_exchange_explicit(&globals.c[1000], 1u, metal::memory_order_relaxed) : DefaultConstructible();
    return _e4;
}