writer/msl: Implement atomics

Common logic between the HLSL, WGSL and MSL writers has been moved into
the TextGenerator base class.

Fixed: tint:892
Change-Id: I0f469516947fe64817ce6251e436da74e5e176e8
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56068
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Reviewed-by: James Price <jrprice@google.com>

test/BUILD.gn

@@ -305,6 +305,7 @@ tint_unittests_source_set("tint_unittests_core_src") {
     "../src/utils/unique_vector_test.cc",
     "../src/writer/append_vector_test.cc",
     "../src/writer/float_to_string_test.cc",
+    "../src/writer/text_generator_test.cc",
   ]
 
   deps = [

test/intrinsics/gen/atomicAdd/794055.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicAdd_794055(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : i32 = atomicAdd(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicAdd_794055(threadgroup atomic_int* const tint_symbol_1) {
+  int res = atomic_fetch_add_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicAdd_794055(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicAdd_794055(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicAdd/8a199a.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicAdd_8a199a() {
-  var res : u32 = atomicAdd(&(sb_rw.arg_0), 1u);
+void atomicAdd_8a199a(device SB_RW& sb_rw) {
+  uint res = atomic_fetch_add_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicAdd_8a199a();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicAdd_8a199a(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicAdd_8a199a();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicAdd_8a199a(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32

test/intrinsics/gen/atomicAdd/d32fe4.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicAdd_d32fe4() {
-  var res : i32 = atomicAdd(&(sb_rw.arg_0), 1);
+void atomicAdd_d32fe4(device SB_RW& sb_rw) {
+  int res = atomic_fetch_add_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicAdd_d32fe4();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicAdd_d32fe4(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicAdd_d32fe4();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicAdd_d32fe4(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32

test/intrinsics/gen/atomicAdd/d5db1d.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicAdd_d5db1d(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : u32 = atomicAdd(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicAdd_d5db1d(threadgroup atomic_uint* const tint_symbol_1) {
+  uint res = atomic_fetch_add_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicAdd_d5db1d(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicAdd_d5db1d(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicAnd/152966.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicAnd_152966() {
-  var res : i32 = atomicAnd(&(sb_rw.arg_0), 1);
+void atomicAnd_152966(device SB_RW& sb_rw) {
+  int res = atomic_fetch_and_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicAnd_152966();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicAnd_152966(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicAnd_152966();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicAnd_152966(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32

test/intrinsics/gen/atomicAnd/34edd3.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicAnd_34edd3(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : u32 = atomicAnd(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicAnd_34edd3(threadgroup atomic_uint* const tint_symbol_1) {
+  uint res = atomic_fetch_and_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicAnd_34edd3(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicAnd_34edd3(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicAnd/45a819.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicAnd_45a819(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : i32 = atomicAnd(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicAnd_45a819(threadgroup atomic_int* const tint_symbol_1) {
+  int res = atomic_fetch_and_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicAnd_45a819(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicAnd_45a819(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicAnd/85a8d9.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicAnd_85a8d9() {
-  var res : u32 = atomicAnd(&(sb_rw.arg_0), 1u);
+void atomicAnd_85a8d9(device SB_RW& sb_rw) {
+  uint res = atomic_fetch_and_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicAnd_85a8d9();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicAnd_85a8d9(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicAnd_85a8d9();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicAnd_85a8d9(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32

test/intrinsics/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl

@@ -1,25 +1,23 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicCompareExchangeWeak_12871c() {
-  var res : vec2<i32> = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1, 1);
+void atomicCompareExchangeWeak_12871c(device SB_RW& sb_rw) {
+  int prev_value = 1;
+  bool matched = atomic_compare_exchange_weak_explicit(&(sb_rw.arg_0), &prev_value, 1, memory_order_relaxed, memory_order_relaxed);
+  int2 res = int2(prev_value, matched);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicCompareExchangeWeak_12871c();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicCompareExchangeWeak_12871c(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicCompareExchangeWeak_12871c();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicCompareExchangeWeak_12871c(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32

test/intrinsics/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl

@@ -1,25 +1,23 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicCompareExchangeWeak_6673da() {
-  var res : vec2<u32> = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1u, 1u);
+void atomicCompareExchangeWeak_6673da(device SB_RW& sb_rw) {
+  uint prev_value = 1u;
+  bool matched = atomic_compare_exchange_weak_explicit(&(sb_rw.arg_0), &prev_value, 1u, memory_order_relaxed, memory_order_relaxed);
+  uint2 res = uint2(prev_value, matched);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicCompareExchangeWeak_6673da();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicCompareExchangeWeak_6673da(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicCompareExchangeWeak_6673da();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicCompareExchangeWeak_6673da(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32

test/intrinsics/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl

@@ -1,14 +1,19 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicCompareExchangeWeak_89ea3b(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : vec2<i32> = atomicCompareExchangeWeak(&(*(tint_symbol)), 1, 1);
+using namespace metal;
+void atomicCompareExchangeWeak_89ea3b(threadgroup atomic_int* const tint_symbol_1) {
+  int prev_value = 1;
+  bool matched = atomic_compare_exchange_weak_explicit(&(*(tint_symbol_1)), &prev_value, 1, memory_order_relaxed, memory_order_relaxed);
+  int2 res = int2(prev_value, matched);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicCompareExchangeWeak_89ea3b(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicCompareExchangeWeak_89ea3b(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl

@@ -1,14 +1,19 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicCompareExchangeWeak_b2ab2c(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : vec2<u32> = atomicCompareExchangeWeak(&(*(tint_symbol)), 1u, 1u);
+using namespace metal;
+void atomicCompareExchangeWeak_b2ab2c(threadgroup atomic_uint* const tint_symbol_1) {
+  uint prev_value = 1u;
+  bool matched = atomic_compare_exchange_weak_explicit(&(*(tint_symbol_1)), &prev_value, 1u, memory_order_relaxed, memory_order_relaxed);
+  uint2 res = uint2(prev_value, matched);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicCompareExchangeWeak_b2ab2c(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicCompareExchangeWeak_b2ab2c(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicExchange/0a5dca.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicExchange_0a5dca(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : u32 = atomicExchange(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicExchange_0a5dca(threadgroup atomic_uint* const tint_symbol_1) {
+  uint res = atomic_exchange_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicExchange_0a5dca(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicExchange_0a5dca(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicExchange/d59712.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicExchange_d59712() {
-  var res : u32 = atomicExchange(&(sb_rw.arg_0), 1u);
+void atomicExchange_d59712(device SB_RW& sb_rw) {
+  uint res = atomic_exchange_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicExchange_d59712();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicExchange_d59712(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicExchange_d59712();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicExchange_d59712(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32

test/intrinsics/gen/atomicExchange/e114ba.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicExchange_e114ba(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : i32 = atomicExchange(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicExchange_e114ba(threadgroup atomic_int* const tint_symbol_1) {
+  int res = atomic_exchange_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicExchange_e114ba(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicExchange_e114ba(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicExchange/f2e22f.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicExchange_f2e22f() {
-  var res : i32 = atomicExchange(&(sb_rw.arg_0), 1);
+void atomicExchange_f2e22f(device SB_RW& sb_rw) {
+  int res = atomic_exchange_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicExchange_f2e22f();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicExchange_f2e22f(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicExchange_f2e22f();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicExchange_f2e22f(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32

test/intrinsics/gen/atomicLoad/0806ad.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicLoad_0806ad() {
-  var res : i32 = atomicLoad(&(sb_rw.arg_0));
+void atomicLoad_0806ad(device SB_RW& sb_rw) {
+  int res = atomic_load_explicit(&(sb_rw.arg_0), memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicLoad_0806ad();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicLoad_0806ad(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicLoad_0806ad();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicLoad_0806ad(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32

test/intrinsics/gen/atomicLoad/361bf1.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicLoad_361bf1(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : u32 = atomicLoad(&(*(tint_symbol)));
+using namespace metal;
+void atomicLoad_361bf1(threadgroup atomic_uint* const tint_symbol_1) {
+  uint res = atomic_load_explicit(&(*(tint_symbol_1)), memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicLoad_361bf1(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicLoad_361bf1(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicLoad/afcc03.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicLoad_afcc03(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : i32 = atomicLoad(&(*(tint_symbol)));
+using namespace metal;
+void atomicLoad_afcc03(threadgroup atomic_int* const tint_symbol_1) {
+  int res = atomic_load_explicit(&(*(tint_symbol_1)), memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicLoad_afcc03(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicLoad_afcc03(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicLoad/fe6cc3.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicLoad_fe6cc3() {
-  var res : u32 = atomicLoad(&(sb_rw.arg_0));
+void atomicLoad_fe6cc3(device SB_RW& sb_rw) {
+  uint res = atomic_load_explicit(&(sb_rw.arg_0), memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicLoad_fe6cc3();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicLoad_fe6cc3(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicLoad_fe6cc3();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicLoad_fe6cc3(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32

test/intrinsics/gen/atomicMax/51b9be.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicMax_51b9be() {
-  var res : u32 = atomicMax(&(sb_rw.arg_0), 1u);
+void atomicMax_51b9be(device SB_RW& sb_rw) {
+  uint res = atomic_fetch_max_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicMax_51b9be();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicMax_51b9be(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicMax_51b9be();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicMax_51b9be(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32

test/intrinsics/gen/atomicMax/92aa72.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicMax_92aa72() {
-  var res : i32 = atomicMax(&(sb_rw.arg_0), 1);
+void atomicMax_92aa72(device SB_RW& sb_rw) {
+  int res = atomic_fetch_max_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicMax_92aa72();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicMax_92aa72(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicMax_92aa72();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicMax_92aa72(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32

test/intrinsics/gen/atomicMax/a89cc3.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicMax_a89cc3(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : i32 = atomicMax(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicMax_a89cc3(threadgroup atomic_int* const tint_symbol_1) {
+  int res = atomic_fetch_max_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicMax_a89cc3(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicMax_a89cc3(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicMax/beccfc.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicMax_beccfc(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : u32 = atomicMax(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicMax_beccfc(threadgroup atomic_uint* const tint_symbol_1) {
+  uint res = atomic_fetch_max_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicMax_beccfc(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicMax_beccfc(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicMin/278235.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicMin_278235(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : i32 = atomicMin(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicMin_278235(threadgroup atomic_int* const tint_symbol_1) {
+  int res = atomic_fetch_min_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicMin_278235(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicMin_278235(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicMin/69d383.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicMin_69d383(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : u32 = atomicMin(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicMin_69d383(threadgroup atomic_uint* const tint_symbol_1) {
+  uint res = atomic_fetch_min_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicMin_69d383(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicMin_69d383(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicMin/8e38dc.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicMin_8e38dc() {
-  var res : i32 = atomicMin(&(sb_rw.arg_0), 1);
+void atomicMin_8e38dc(device SB_RW& sb_rw) {
+  int res = atomic_fetch_min_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicMin_8e38dc();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicMin_8e38dc(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicMin_8e38dc();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicMin_8e38dc(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32

test/intrinsics/gen/atomicMin/c67a74.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicMin_c67a74() {
-  var res : u32 = atomicMin(&(sb_rw.arg_0), 1u);
+void atomicMin_c67a74(device SB_RW& sb_rw) {
+  uint res = atomic_fetch_min_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicMin_c67a74();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicMin_c67a74(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicMin_c67a74();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicMin_c67a74(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32

test/intrinsics/gen/atomicOr/5e3d61.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicOr_5e3d61(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : u32 = atomicOr(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicOr_5e3d61(threadgroup atomic_uint* const tint_symbol_1) {
+  uint res = atomic_fetch_or_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicOr_5e3d61(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicOr_5e3d61(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicOr/5e95d4.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicOr_5e95d4() {
-  var res : u32 = atomicOr(&(sb_rw.arg_0), 1u);
+void atomicOr_5e95d4(device SB_RW& sb_rw) {
+  uint res = atomic_fetch_or_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicOr_5e95d4();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicOr_5e95d4(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicOr_5e95d4();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicOr_5e95d4(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32

test/intrinsics/gen/atomicOr/8d96a0.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicOr_8d96a0() {
-  var res : i32 = atomicOr(&(sb_rw.arg_0), 1);
+void atomicOr_8d96a0(device SB_RW& sb_rw) {
+  int res = atomic_fetch_or_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicOr_8d96a0();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicOr_8d96a0(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicOr_8d96a0();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicOr_8d96a0(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32

test/intrinsics/gen/atomicOr/d09248.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicOr_d09248(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : i32 = atomicOr(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicOr_d09248(threadgroup atomic_int* const tint_symbol_1) {
+  int res = atomic_fetch_or_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicOr_d09248(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicOr_d09248(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicStore/726882.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicStore_726882(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  atomicStore(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicStore_726882(threadgroup atomic_uint* const tint_symbol_1) {
+  atomic_store_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicStore_726882(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicStore_726882(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicStore/8bea94.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicStore_8bea94(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  atomicStore(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicStore_8bea94(threadgroup atomic_int* const tint_symbol_1) {
+  atomic_store_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicStore_8bea94(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicStore_8bea94(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicStore/cdc29e.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicStore_cdc29e() {
-  atomicStore(&(sb_rw.arg_0), 1u);
+void atomicStore_cdc29e(device SB_RW& sb_rw) {
+  atomic_store_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicStore_cdc29e();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicStore_cdc29e(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicStore_cdc29e();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicStore_cdc29e(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32

test/intrinsics/gen/atomicStore/d1e9a6.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicStore_d1e9a6() {
-  atomicStore(&(sb_rw.arg_0), 1);
+void atomicStore_d1e9a6(device SB_RW& sb_rw) {
+  atomic_store_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicStore_d1e9a6();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicStore_d1e9a6(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicStore_d1e9a6();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicStore_d1e9a6(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32

test/intrinsics/gen/atomicXor/54510e.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicXor_54510e() {
-  var res : u32 = atomicXor(&(sb_rw.arg_0), 1u);
+void atomicXor_54510e(device SB_RW& sb_rw) {
+  uint res = atomic_fetch_xor_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicXor_54510e();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicXor_54510e(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicXor_54510e();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicXor_54510e(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32

test/intrinsics/gen/atomicXor/75dc95.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicXor_75dc95(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : i32 = atomicXor(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicXor_75dc95(threadgroup atomic_int* const tint_symbol_1) {
+  int res = atomic_fetch_xor_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicXor_75dc95(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicXor_75dc95(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope

test/intrinsics/gen/atomicXor/c1b78c.wgsl.expected.msl

@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicXor_c1b78c() {
-  var res : i32 = atomicXor(&(sb_rw.arg_0), 1);
+void atomicXor_c1b78c(device SB_RW& sb_rw) {
+  int res = atomic_fetch_xor_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicXor_c1b78c();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicXor_c1b78c(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicXor_c1b78c();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicXor_c1b78c(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32

test/intrinsics/gen/atomicXor/c8e6be.wgsl.expected.msl

@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicXor_c8e6be(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : u32 = atomicXor(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicXor_c8e6be(threadgroup atomic_uint* const tint_symbol_1) {
+  uint res = atomic_fetch_xor_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicXor_c8e6be(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicXor_c8e6be(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope