mirror of
https://github.com/gfx-rs/wgpu.git
synced 2025-12-08 21:26:17 +00:00
195 lines
13 KiB
Plaintext
195 lines
13 KiB
Plaintext
// language: metal1.0
|
|
#include <metal_stdlib>
|
|
#include <simd/simd.h>
|
|
|
|
using metal::uint;
|
|
|
|
struct type_2 {
|
|
metal::atomic_int inner[2];
|
|
};
|
|
struct Struct {
|
|
metal::atomic_uint atomic_scalar;
|
|
type_2 atomic_arr;
|
|
};
|
|
struct _atomic_compare_exchange_resultUint4_ {
|
|
uint old_value;
|
|
bool exchanged;
|
|
char _pad2[3];
|
|
};
|
|
struct _atomic_compare_exchange_resultSint4_ {
|
|
int old_value;
|
|
bool exchanged;
|
|
char _pad2[3];
|
|
};
|
|
|
|
template <typename A>
|
|
_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
|
|
device A *atomic_ptr,
|
|
uint cmp,
|
|
uint v
|
|
) {
|
|
bool swapped = metal::atomic_compare_exchange_weak_explicit(
|
|
atomic_ptr, &cmp, v,
|
|
metal::memory_order_relaxed, metal::memory_order_relaxed
|
|
);
|
|
return _atomic_compare_exchange_resultUint4_{cmp, swapped};
|
|
}
|
|
template <typename A>
|
|
_atomic_compare_exchange_resultUint4_ naga_atomic_compare_exchange_weak_explicit(
|
|
threadgroup A *atomic_ptr,
|
|
uint cmp,
|
|
uint v
|
|
) {
|
|
bool swapped = metal::atomic_compare_exchange_weak_explicit(
|
|
atomic_ptr, &cmp, v,
|
|
metal::memory_order_relaxed, metal::memory_order_relaxed
|
|
);
|
|
return _atomic_compare_exchange_resultUint4_{cmp, swapped};
|
|
}
|
|
|
|
template <typename A>
|
|
_atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit(
|
|
device A *atomic_ptr,
|
|
int cmp,
|
|
int v
|
|
) {
|
|
bool swapped = metal::atomic_compare_exchange_weak_explicit(
|
|
atomic_ptr, &cmp, v,
|
|
metal::memory_order_relaxed, metal::memory_order_relaxed
|
|
);
|
|
return _atomic_compare_exchange_resultSint4_{cmp, swapped};
|
|
}
|
|
template <typename A>
|
|
_atomic_compare_exchange_resultSint4_ naga_atomic_compare_exchange_weak_explicit(
|
|
threadgroup A *atomic_ptr,
|
|
int cmp,
|
|
int v
|
|
) {
|
|
bool swapped = metal::atomic_compare_exchange_weak_explicit(
|
|
atomic_ptr, &cmp, v,
|
|
metal::memory_order_relaxed, metal::memory_order_relaxed
|
|
);
|
|
return _atomic_compare_exchange_resultSint4_{cmp, swapped};
|
|
}
|
|
|
|
struct cs_mainInput {
|
|
};
|
|
kernel void cs_main(
|
|
metal::uint3 id [[thread_position_in_threadgroup]]
|
|
, device metal::atomic_uint& storage_atomic_scalar [[user(fake0)]]
|
|
, device type_2& storage_atomic_arr [[user(fake0)]]
|
|
, device Struct& storage_struct [[user(fake0)]]
|
|
, threadgroup metal::atomic_uint& workgroup_atomic_scalar
|
|
, threadgroup type_2& workgroup_atomic_arr
|
|
, threadgroup Struct& workgroup_struct
|
|
) {
|
|
if (metal::all(id == metal::uint3(0u))) {
|
|
metal::atomic_store_explicit(&workgroup_atomic_scalar, 0, metal::memory_order_relaxed);
|
|
for (int __i0 = 0; __i0 < 2; __i0++) {
|
|
metal::atomic_store_explicit(&workgroup_atomic_arr.inner[__i0], 0, metal::memory_order_relaxed);
|
|
}
|
|
metal::atomic_store_explicit(&workgroup_struct.atomic_scalar, 0, metal::memory_order_relaxed);
|
|
for (int __i0 = 0; __i0 < 2; __i0++) {
|
|
metal::atomic_store_explicit(&workgroup_struct.atomic_arr.inner[__i0], 0, metal::memory_order_relaxed);
|
|
}
|
|
}
|
|
metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
|
|
metal::atomic_store_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
metal::atomic_store_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
metal::atomic_store_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
metal::atomic_store_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
metal::atomic_store_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
metal::atomic_store_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
metal::atomic_store_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
metal::atomic_store_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
|
|
uint l0_ = metal::atomic_load_explicit(&storage_atomic_scalar, metal::memory_order_relaxed);
|
|
int l1_ = metal::atomic_load_explicit(&storage_atomic_arr.inner[1], metal::memory_order_relaxed);
|
|
uint l2_ = metal::atomic_load_explicit(&storage_struct.atomic_scalar, metal::memory_order_relaxed);
|
|
int l3_ = metal::atomic_load_explicit(&storage_struct.atomic_arr.inner[1], metal::memory_order_relaxed);
|
|
uint l4_ = metal::atomic_load_explicit(&workgroup_atomic_scalar, metal::memory_order_relaxed);
|
|
int l5_ = metal::atomic_load_explicit(&workgroup_atomic_arr.inner[1], metal::memory_order_relaxed);
|
|
uint l6_ = metal::atomic_load_explicit(&workgroup_struct.atomic_scalar, metal::memory_order_relaxed);
|
|
int l7_ = metal::atomic_load_explicit(&workgroup_struct.atomic_arr.inner[1], metal::memory_order_relaxed);
|
|
metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
|
|
uint _e51 = metal::atomic_fetch_add_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e55 = metal::atomic_fetch_add_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e59 = metal::atomic_fetch_add_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e64 = metal::atomic_fetch_add_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e67 = metal::atomic_fetch_add_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e71 = metal::atomic_fetch_add_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e75 = metal::atomic_fetch_add_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e80 = metal::atomic_fetch_add_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
|
|
uint _e83 = metal::atomic_fetch_sub_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e87 = metal::atomic_fetch_sub_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e91 = metal::atomic_fetch_sub_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e96 = metal::atomic_fetch_sub_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e99 = metal::atomic_fetch_sub_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e103 = metal::atomic_fetch_sub_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e107 = metal::atomic_fetch_sub_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e112 = metal::atomic_fetch_sub_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
|
|
uint _e115 = metal::atomic_fetch_max_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e119 = metal::atomic_fetch_max_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e123 = metal::atomic_fetch_max_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e128 = metal::atomic_fetch_max_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e131 = metal::atomic_fetch_max_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e135 = metal::atomic_fetch_max_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e139 = metal::atomic_fetch_max_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e144 = metal::atomic_fetch_max_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
|
|
uint _e147 = metal::atomic_fetch_min_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e151 = metal::atomic_fetch_min_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e155 = metal::atomic_fetch_min_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e160 = metal::atomic_fetch_min_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e163 = metal::atomic_fetch_min_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e167 = metal::atomic_fetch_min_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e171 = metal::atomic_fetch_min_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e176 = metal::atomic_fetch_min_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
|
|
uint _e179 = metal::atomic_fetch_and_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e183 = metal::atomic_fetch_and_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e187 = metal::atomic_fetch_and_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e192 = metal::atomic_fetch_and_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e195 = metal::atomic_fetch_and_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e199 = metal::atomic_fetch_and_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e203 = metal::atomic_fetch_and_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e208 = metal::atomic_fetch_and_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
|
|
uint _e211 = metal::atomic_fetch_or_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e215 = metal::atomic_fetch_or_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e219 = metal::atomic_fetch_or_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e224 = metal::atomic_fetch_or_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e227 = metal::atomic_fetch_or_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e231 = metal::atomic_fetch_or_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e235 = metal::atomic_fetch_or_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e240 = metal::atomic_fetch_or_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
metal::threadgroup_barrier(metal::mem_flags::mem_threadgroup);
|
|
uint _e243 = metal::atomic_fetch_xor_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e247 = metal::atomic_fetch_xor_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e251 = metal::atomic_fetch_xor_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e256 = metal::atomic_fetch_xor_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e259 = metal::atomic_fetch_xor_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e263 = metal::atomic_fetch_xor_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e267 = metal::atomic_fetch_xor_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e272 = metal::atomic_fetch_xor_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e275 = metal::atomic_exchange_explicit(&storage_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e279 = metal::atomic_exchange_explicit(&storage_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e283 = metal::atomic_exchange_explicit(&storage_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e288 = metal::atomic_exchange_explicit(&storage_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e291 = metal::atomic_exchange_explicit(&workgroup_atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e295 = metal::atomic_exchange_explicit(&workgroup_atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
uint _e299 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_scalar, 1u, metal::memory_order_relaxed);
|
|
int _e304 = metal::atomic_exchange_explicit(&workgroup_struct.atomic_arr.inner[1], 1, metal::memory_order_relaxed);
|
|
_atomic_compare_exchange_resultUint4_ _e308 = naga_atomic_compare_exchange_weak_explicit(&storage_atomic_scalar, 1u, 2u);
|
|
_atomic_compare_exchange_resultSint4_ _e313 = naga_atomic_compare_exchange_weak_explicit(&storage_atomic_arr.inner[1], 1, 2);
|
|
_atomic_compare_exchange_resultUint4_ _e318 = naga_atomic_compare_exchange_weak_explicit(&storage_struct.atomic_scalar, 1u, 2u);
|
|
_atomic_compare_exchange_resultSint4_ _e324 = naga_atomic_compare_exchange_weak_explicit(&storage_struct.atomic_arr.inner[1], 1, 2);
|
|
_atomic_compare_exchange_resultUint4_ _e328 = naga_atomic_compare_exchange_weak_explicit(&workgroup_atomic_scalar, 1u, 2u);
|
|
_atomic_compare_exchange_resultSint4_ _e333 = naga_atomic_compare_exchange_weak_explicit(&workgroup_atomic_arr.inner[1], 1, 2);
|
|
_atomic_compare_exchange_resultUint4_ _e338 = naga_atomic_compare_exchange_weak_explicit(&workgroup_struct.atomic_scalar, 1u, 2u);
|
|
_atomic_compare_exchange_resultSint4_ _e344 = naga_atomic_compare_exchange_weak_explicit(&workgroup_struct.atomic_arr.inner[1], 1, 2);
|
|
return;
|
|
}
|