tint: fix emitting duplicate structs for atomicCompareExchangeWeak

Bug: tint:1574
Change-Id: Id4ae2d2de9ac4678260f4ecfb3a0f779d170f9a4
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/92280
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>

src/tint/writer/glsl/generator_impl.cc

@@ -920,7 +920,7 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
         case sem::BuiltinType::kAtomicCompareExchangeWeak: {
             // Emit the builtin return type unique to this overload. This does not
             // exist in the AST, so it will not be generated in Generate().
-            if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
+            if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
                 return false;
             }
 
@@ -2822,6 +2822,14 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
     return true;
 }
 
+bool GeneratorImpl::EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* str) {
+    auto it = emitted_structs_.emplace(str);
+    if (!it.second) {
+        return true;
+    }
+    return EmitStructType(buffer, str);
+}
+
 bool GeneratorImpl::EmitStructMembers(TextBuffer* b, const sem::Struct* str, bool emit_offsets) {
     ScopedIndent si(b);
     for (auto* mem : str->Members()) {

src/tint/writer/glsl/generator_impl.h

@@ -411,6 +411,12 @@ class GeneratorImpl : public TextGenerator {
     /// @param ty the struct to generate
     /// @returns true if the struct is emitted
     bool EmitStructType(TextBuffer* buffer, const sem::Struct* ty);
+    /// Handles generating a structure declaration only the first time called. Subsequent calls are
+    /// a no-op and return true.
+    /// @param buffer the text buffer that the type declaration will be written to
+    /// @param ty the struct to generate
+    /// @returns true if the struct is emitted
+    bool EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* ty);
     /// Handles generating the members of a structure
     /// @param buffer the text buffer that the struct members will be written to
     /// @param ty the struct to generate
@@ -503,6 +509,7 @@ class GeneratorImpl : public TextGenerator {
     std::unordered_map<const sem::Vector*, std::string> dynamic_vector_write_;
     std::unordered_map<const sem::Vector*, std::string> int_dot_funcs_;
     std::unordered_map<const sem::Type*, std::string> float_modulo_funcs_;
+    std::unordered_set<const sem::Struct*> emitted_structs_;
     bool requires_oes_sample_variables_ = false;
     bool requires_default_precision_qualifier_ = false;
     Version version_;

src/tint/writer/hlsl/generator_impl.cc

@@ -1767,7 +1767,7 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
         case sem::BuiltinType::kAtomicCompareExchangeWeak: {
             // Emit the builtin return type unique to this overload. This does not
             // exist in the AST, so it will not be generated in Generate().
-            if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
+            if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
                 return false;
             }
 
@@ -3921,6 +3921,14 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
     return true;
 }
 
+bool GeneratorImpl::EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* str) {
+    auto it = emitted_structs_.emplace(str);
+    if (!it.second) {
+        return true;
+    }
+    return EmitStructType(buffer, str);
+}
+
 bool GeneratorImpl::EmitUnaryOp(std::ostream& out, const ast::UnaryOpExpression* expr) {
     switch (expr->op) {
         case ast::UnaryOp::kIndirection:

src/tint/writer/hlsl/generator_impl.h

@@ -411,6 +411,12 @@ class GeneratorImpl : public TextGenerator {
     /// @param ty the struct to generate
     /// @returns true if the struct is emitted
     bool EmitStructType(TextBuffer* buffer, const sem::Struct* ty);
+    /// Handles generating a structure declaration only the first time called. Subsequent calls are
+    /// a no-op and return true.
+    /// @param buffer the text buffer that the type declaration will be written to
+    /// @param ty the struct to generate
+    /// @returns true if the struct is emitted
+    bool EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* ty);
     /// Handles a unary op expression
     /// @param out the output of the expression stream
     /// @param expr the expression to emit
@@ -530,6 +536,7 @@ class GeneratorImpl : public TextGenerator {
     std::unordered_map<const sem::Matrix*, std::string> dynamic_matrix_vector_write_;
     std::unordered_map<const sem::Matrix*, std::string> dynamic_matrix_scalar_write_;
     std::unordered_map<const sem::Type*, std::string> value_or_one_if_zero_;
+    std::unordered_set<const sem::Struct*> emitted_structs_;
 };
 
 }  // namespace tint::writer::hlsl

src/tint/writer/msl/generator_impl.cc

@@ -826,28 +826,39 @@ bool GeneratorImpl::EmitAtomicCall(std::ostream& out,
             return call("atomic_exchange_explicit", true);
 
         case sem::BuiltinType::kAtomicCompareExchangeWeak: {
-            // Emit the builtin return type unique to this overload. This does not
-            // exist in the AST, so it will not be generated in Generate().
-            if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
-                return false;
-            }
-
             auto* ptr_ty = TypeOf(expr->args[0])->UnwrapRef()->As<sem::Pointer>();
             auto sc = ptr_ty->StorageClass();
+            auto* str = builtin->ReturnType()->As<sem::Struct>();
+
+            auto func = utils::GetOrCreate(
+                atomicCompareExchangeWeak_, ACEWKeyType{{sc, str}}, [&]() -> std::string {
+                    // Emit the builtin return type unique to this overload. This does not
+                    // exist in the AST, so it will not be generated in Generate().
+                    if (!EmitStructTypeOnce(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
+                        return "";
+                    }
 
-            auto func = utils::GetOrCreate(atomicCompareExchangeWeak_, sc, [&]() -> std::string {
                     auto name = UniqueIdentifier("atomicCompareExchangeWeak");
                     auto& buf = helpers_;
+                    auto* atomic_ty = builtin->Parameters()[0]->Type();
+                    auto* arg_ty = builtin->Parameters()[1]->Type();
 
-                line(&buf) << "template <typename A, typename T>";
                     {
                         auto f = line(&buf);
                         auto str_name = StructName(builtin->ReturnType()->As<sem::Struct>());
                         f << str_name << " " << name << "(";
-                    if (!EmitStorageClass(f, sc)) {
+                        if (!EmitTypeAndName(f, atomic_ty, "atomic")) {
                             return "";
                         }
-                    f << " A* atomic, T compare, T value) {";
+                        f << ", ";
+                        if (!EmitTypeAndName(f, arg_ty, "compare")) {
+                            return "";
+                        }
+                        f << ", ";
+                        if (!EmitTypeAndName(f, arg_ty, "value")) {
+                            return "";
+                        }
+                        f << ") {";
                     }
 
                     buf.IncrementIndent();
@@ -857,7 +868,13 @@ bool GeneratorImpl::EmitAtomicCall(std::ostream& out,
                         line(&buf);
                     });
 
-                line(&buf) << "T old_value = compare;";
+                    {
+                        auto f = line(&buf);
+                        if (!EmitTypeAndName(f, arg_ty, "old_value")) {
+                            return "";
+                        }
+                        f << " = compare;";
+                    }
                     line(&buf) << "bool exchanged = "
                                   "atomic_compare_exchange_weak_explicit(atomic, "
                                   "&old_value, value, memory_order_relaxed, "
@@ -866,6 +883,9 @@ bool GeneratorImpl::EmitAtomicCall(std::ostream& out,
                     return name;
                 });
 
+            if (func.empty()) {
+                return false;
+            }
             return call(func, false);
         }
 
@@ -2765,6 +2785,14 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
     return true;
 }
 
+bool GeneratorImpl::EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* str) {
+    auto it = emitted_structs_.emplace(str);
+    if (!it.second) {
+        return true;
+    }
+    return EmitStructType(buffer, str);
+}
+
 bool GeneratorImpl::EmitUnaryOp(std::ostream& out, const ast::UnaryOpExpression* expr) {
     // Handle `-e` when `e` is signed, so that we ensure that if `e` is the
     // largest negative value, it returns `e`.

src/tint/writer/msl/generator_impl.h

@@ -16,6 +16,7 @@
 #define SRC_TINT_WRITER_MSL_GENERATOR_IMPL_H_
 
 #include <string>
+#include <tuple>
 #include <unordered_map>
 #include <unordered_set>
 #include <vector>
@@ -332,6 +333,12 @@ class GeneratorImpl : public TextGenerator {
     /// @param str the struct to generate
     /// @returns true if the struct is emitted
     bool EmitStructType(TextBuffer* buffer, const sem::Struct* str);
+    /// Handles generating a structure declaration only the first time called. Subsequent calls are
+    /// a no-op and return true.
+    /// @param buffer the text buffer that the type declaration will be written to
+    /// @param ty the struct to generate
+    /// @returns true if the struct is emitted
+    bool EmitStructTypeOnce(TextBuffer* buffer, const sem::Struct* ty);
     /// Handles a unary op expression
     /// @param out the output of the expression stream
     /// @param expr the expression to emit
@@ -400,13 +407,13 @@ class GeneratorImpl : public TextGenerator {
     /// type.
     SizeAndAlign MslPackedTypeSizeAndAlign(const sem::Type* ty);
 
-    using StorageClassToString = std::unordered_map<ast::StorageClass, std::string>;
-
     std::function<bool()> emit_continuing_;
 
     /// Name of atomicCompareExchangeWeak() helper for the given pointer storage
-    /// class.
+    /// class and struct return type
-    StorageClassToString atomicCompareExchangeWeak_;
+    using ACEWKeyType =
+        utils::UnorderedKeyWrapper<std::tuple<ast::StorageClass, const sem::Struct*>>;
+    std::unordered_map<ACEWKeyType, std::string> atomicCompareExchangeWeak_;
 
     /// Unique name of the 'TINT_INVARIANT' preprocessor define. Non-empty only if
     /// an invariant attribute has been generated.
@@ -423,6 +430,7 @@ class GeneratorImpl : public TextGenerator {
     std::unordered_map<const sem::Builtin*, std::string> builtins_;
     std::unordered_map<const sem::Type*, std::string> unary_minus_funcs_;
     std::unordered_map<uint32_t, std::string> int_dot_funcs_;
+    std::unordered_set<const sem::Struct*> emitted_structs_;
 };
 
 }  // namespace tint::writer::msl

test/tint/bug/tint/1573.wgsl.expected.msl

@@ -6,9 +6,8 @@ struct atomic_compare_exchange_resultu32 {
   uint old_value;
   bool exchanged;
 };
-template <typename A, typename T>
+atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) {
-atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device A* atomic, T compare, T value) {
+  uint old_value = compare;
-  T old_value = compare;
   bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
   return {old_value, exchanged};
 }

test/tint/bug/tint/1574.wgsl

@@ -0,0 +1,39 @@
+@group(0) @binding(0)
+var<storage, read_write> a_u32 : atomic<u32>;
+
+@group(0) @binding(1)
+var<storage, read_write> a_i32 : atomic<i32>;
+
+var<workgroup> b_u32 : atomic<u32>;
+
+var<workgroup> b_i32 : atomic<i32>;
+
+
+@stage(compute) @workgroup_size(16)
+fn main() {
+  {
+    var value = 42u;
+    let r1 = atomicCompareExchangeWeak(&a_u32, 0u, value);
+    let r2 = atomicCompareExchangeWeak(&a_u32, 0u, value);
+    let r3 = atomicCompareExchangeWeak(&a_u32, 0u, value);
+  }
+  {
+    var value = 42;
+    let r1 = atomicCompareExchangeWeak(&a_i32, 0, value);
+    let r2 = atomicCompareExchangeWeak(&a_i32, 0, value);
+    let r3 = atomicCompareExchangeWeak(&a_i32, 0, value);
+  }
+  {
+    var value = 42u;
+    let r1 = atomicCompareExchangeWeak(&b_u32, 0u, value);
+    let r2 = atomicCompareExchangeWeak(&b_u32, 0u, value);
+    let r3 = atomicCompareExchangeWeak(&b_u32, 0u, value);
+  }
+  {
+    var value = 42;
+    let r1 = atomicCompareExchangeWeak(&b_i32, 0, value);
+    let r2 = atomicCompareExchangeWeak(&b_i32, 0, value);
+    let r3 = atomicCompareExchangeWeak(&b_i32, 0, value);
+  }
+
+}

test/tint/bug/tint/1574.wgsl.expected.glsl

@@ -0,0 +1,102 @@
+#version 310 es
+
+struct atomic_compare_exchange_resultu32 {
+  uint old_value;
+  bool exchanged;
+};
+
+struct atomic_compare_exchange_resulti32 {
+  int old_value;
+  bool exchanged;
+};
+
+
+struct a_u32_block {
+  uint inner;
+};
+
+layout(binding = 0, std430) buffer a_u32_block_1 {
+  uint inner;
+} a_u32;
+struct a_i32_block {
+  int inner;
+};
+
+layout(binding = 1, std430) buffer a_i32_block_1 {
+  int inner;
+} a_i32;
+shared uint b_u32;
+shared int b_i32;
+void tint_symbol(uint local_invocation_index) {
+  if ((local_invocation_index < 1u)) {
+    atomicExchange(b_u32, 0u);
+    atomicExchange(b_i32, 0);
+  }
+  barrier();
+  {
+    uint value = 42u;
+    atomic_compare_exchange_resultu32 atomic_compare_result;
+    atomic_compare_result.old_value = atomicCompSwap(a_u32.inner, 0u, value);
+    atomic_compare_result.exchanged = atomic_compare_result.old_value == 0u;
+    atomic_compare_exchange_resultu32 r1 = atomic_compare_result;
+    atomic_compare_exchange_resultu32 atomic_compare_result_1;
+    atomic_compare_result_1.old_value = atomicCompSwap(a_u32.inner, 0u, value);
+    atomic_compare_result_1.exchanged = atomic_compare_result_1.old_value == 0u;
+    atomic_compare_exchange_resultu32 r2 = atomic_compare_result_1;
+    atomic_compare_exchange_resultu32 atomic_compare_result_2;
+    atomic_compare_result_2.old_value = atomicCompSwap(a_u32.inner, 0u, value);
+    atomic_compare_result_2.exchanged = atomic_compare_result_2.old_value == 0u;
+    atomic_compare_exchange_resultu32 r3 = atomic_compare_result_2;
+  }
+  {
+    int value = 42;
+    atomic_compare_exchange_resulti32 atomic_compare_result_3;
+    atomic_compare_result_3.old_value = atomicCompSwap(a_i32.inner, 0, value);
+    atomic_compare_result_3.exchanged = atomic_compare_result_3.old_value == 0;
+    atomic_compare_exchange_resulti32 r1 = atomic_compare_result_3;
+    atomic_compare_exchange_resulti32 atomic_compare_result_4;
+    atomic_compare_result_4.old_value = atomicCompSwap(a_i32.inner, 0, value);
+    atomic_compare_result_4.exchanged = atomic_compare_result_4.old_value == 0;
+    atomic_compare_exchange_resulti32 r2 = atomic_compare_result_4;
+    atomic_compare_exchange_resulti32 atomic_compare_result_5;
+    atomic_compare_result_5.old_value = atomicCompSwap(a_i32.inner, 0, value);
+    atomic_compare_result_5.exchanged = atomic_compare_result_5.old_value == 0;
+    atomic_compare_exchange_resulti32 r3 = atomic_compare_result_5;
+  }
+  {
+    uint value = 42u;
+    atomic_compare_exchange_resultu32 atomic_compare_result_6;
+    atomic_compare_result_6.old_value = atomicCompSwap(b_u32, 0u, value);
+    atomic_compare_result_6.exchanged = atomic_compare_result_6.old_value == 0u;
+    atomic_compare_exchange_resultu32 r1 = atomic_compare_result_6;
+    atomic_compare_exchange_resultu32 atomic_compare_result_7;
+    atomic_compare_result_7.old_value = atomicCompSwap(b_u32, 0u, value);
+    atomic_compare_result_7.exchanged = atomic_compare_result_7.old_value == 0u;
+    atomic_compare_exchange_resultu32 r2 = atomic_compare_result_7;
+    atomic_compare_exchange_resultu32 atomic_compare_result_8;
+    atomic_compare_result_8.old_value = atomicCompSwap(b_u32, 0u, value);
+    atomic_compare_result_8.exchanged = atomic_compare_result_8.old_value == 0u;
+    atomic_compare_exchange_resultu32 r3 = atomic_compare_result_8;
+  }
+  {
+    int value = 42;
+    atomic_compare_exchange_resulti32 atomic_compare_result_9;
+    atomic_compare_result_9.old_value = atomicCompSwap(b_i32, 0, value);
+    atomic_compare_result_9.exchanged = atomic_compare_result_9.old_value == 0;
+    atomic_compare_exchange_resulti32 r1 = atomic_compare_result_9;
+    atomic_compare_exchange_resulti32 atomic_compare_result_10;
+    atomic_compare_result_10.old_value = atomicCompSwap(b_i32, 0, value);
+    atomic_compare_result_10.exchanged = atomic_compare_result_10.old_value == 0;
+    atomic_compare_exchange_resulti32 r2 = atomic_compare_result_10;
+    atomic_compare_exchange_resulti32 atomic_compare_result_11;
+    atomic_compare_result_11.old_value = atomicCompSwap(b_i32, 0, value);
+    atomic_compare_result_11.exchanged = atomic_compare_result_11.old_value == 0;
+    atomic_compare_exchange_resulti32 r3 = atomic_compare_result_11;
+  }
+}
+
+layout(local_size_x = 16, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol(gl_LocalInvocationIndex);
+  return;
+}

test/tint/bug/tint/1574.wgsl.expected.hlsl

@@ -0,0 +1,105 @@
+struct atomic_compare_exchange_resultu32 {
+  uint old_value;
+  bool exchanged;
+};
+struct atomic_compare_exchange_resulti32 {
+  int old_value;
+  bool exchanged;
+};
+RWByteAddressBuffer a_u32 : register(u0, space0);
+RWByteAddressBuffer a_i32 : register(u1, space0);
+groupshared uint b_u32;
+groupshared int b_i32;
+
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+struct atomic_compare_exchange_weak_ret_type {
+  uint old_value;
+  bool exchanged;
+};
+
+atomic_compare_exchange_weak_ret_type tint_atomicCompareExchangeWeak(RWByteAddressBuffer buffer, uint offset, uint compare, uint value) {
+  atomic_compare_exchange_weak_ret_type result=(atomic_compare_exchange_weak_ret_type)0;
+  buffer.InterlockedCompareExchange(offset, compare, value, result.old_value);
+  result.exchanged = result.old_value == compare;
+  return result;
+}
+
+
+struct atomic_compare_exchange_weak_ret_type_1 {
+  int old_value;
+  bool exchanged;
+};
+
+atomic_compare_exchange_weak_ret_type_1 tint_atomicCompareExchangeWeak_1(RWByteAddressBuffer buffer, uint offset, int compare, int value) {
+  atomic_compare_exchange_weak_ret_type_1 result=(atomic_compare_exchange_weak_ret_type_1)0;
+  buffer.InterlockedCompareExchange(offset, compare, value, result.old_value);
+  result.exchanged = result.old_value == compare;
+  return result;
+}
+
+
+void main_inner(uint local_invocation_index) {
+  if ((local_invocation_index < 1u)) {
+    uint atomic_result = 0u;
+    InterlockedExchange(b_u32, 0u, atomic_result);
+    int atomic_result_1 = 0;
+    InterlockedExchange(b_i32, 0, atomic_result_1);
+  }
+  GroupMemoryBarrierWithGroupSync();
+  {
+    uint value = 42u;
+    const atomic_compare_exchange_weak_ret_type r1 = tint_atomicCompareExchangeWeak(a_u32, 0u, 0u, value);
+    const atomic_compare_exchange_weak_ret_type r2 = tint_atomicCompareExchangeWeak(a_u32, 0u, 0u, value);
+    const atomic_compare_exchange_weak_ret_type r3 = tint_atomicCompareExchangeWeak(a_u32, 0u, 0u, value);
+  }
+  {
+    int value = 42;
+    const atomic_compare_exchange_weak_ret_type_1 r1 = tint_atomicCompareExchangeWeak_1(a_i32, 0u, 0, value);
+    const atomic_compare_exchange_weak_ret_type_1 r2 = tint_atomicCompareExchangeWeak_1(a_i32, 0u, 0, value);
+    const atomic_compare_exchange_weak_ret_type_1 r3 = tint_atomicCompareExchangeWeak_1(a_i32, 0u, 0, value);
+  }
+  {
+    uint value = 42u;
+    atomic_compare_exchange_resultu32 atomic_result_2 = (atomic_compare_exchange_resultu32)0;
+    uint atomic_compare_value = 0u;
+    InterlockedCompareExchange(b_u32, atomic_compare_value, value, atomic_result_2.old_value);
+    atomic_result_2.exchanged = atomic_result_2.old_value == atomic_compare_value;
+    const atomic_compare_exchange_resultu32 r1 = atomic_result_2;
+    atomic_compare_exchange_resultu32 atomic_result_3 = (atomic_compare_exchange_resultu32)0;
+    uint atomic_compare_value_1 = 0u;
+    InterlockedCompareExchange(b_u32, atomic_compare_value_1, value, atomic_result_3.old_value);
+    atomic_result_3.exchanged = atomic_result_3.old_value == atomic_compare_value_1;
+    const atomic_compare_exchange_resultu32 r2 = atomic_result_3;
+    atomic_compare_exchange_resultu32 atomic_result_4 = (atomic_compare_exchange_resultu32)0;
+    uint atomic_compare_value_2 = 0u;
+    InterlockedCompareExchange(b_u32, atomic_compare_value_2, value, atomic_result_4.old_value);
+    atomic_result_4.exchanged = atomic_result_4.old_value == atomic_compare_value_2;
+    const atomic_compare_exchange_resultu32 r3 = atomic_result_4;
+  }
+  {
+    int value = 42;
+    atomic_compare_exchange_resulti32 atomic_result_5 = (atomic_compare_exchange_resulti32)0;
+    int atomic_compare_value_3 = 0;
+    InterlockedCompareExchange(b_i32, atomic_compare_value_3, value, atomic_result_5.old_value);
+    atomic_result_5.exchanged = atomic_result_5.old_value == atomic_compare_value_3;
+    const atomic_compare_exchange_resulti32 r1 = atomic_result_5;
+    atomic_compare_exchange_resulti32 atomic_result_6 = (atomic_compare_exchange_resulti32)0;
+    int atomic_compare_value_4 = 0;
+    InterlockedCompareExchange(b_i32, atomic_compare_value_4, value, atomic_result_6.old_value);
+    atomic_result_6.exchanged = atomic_result_6.old_value == atomic_compare_value_4;
+    const atomic_compare_exchange_resulti32 r2 = atomic_result_6;
+    atomic_compare_exchange_resulti32 atomic_result_7 = (atomic_compare_exchange_resulti32)0;
+    int atomic_compare_value_5 = 0;
+    InterlockedCompareExchange(b_i32, atomic_compare_value_5, value, atomic_result_7.old_value);
+    atomic_result_7.exchanged = atomic_result_7.old_value == atomic_compare_value_5;
+    const atomic_compare_exchange_resulti32 r3 = atomic_result_7;
+  }
+}
+
+[numthreads(16, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  main_inner(tint_symbol.local_invocation_index);
+  return;
+}

test/tint/bug/tint/1574.wgsl.expected.msl

@@ -0,0 +1,75 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct atomic_compare_exchange_resultu32 {
+  uint old_value;
+  bool exchanged;
+};
+atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) {
+  uint old_value = compare;
+  bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
+  return {old_value, exchanged};
+}
+
+struct atomic_compare_exchange_resulti32 {
+  int old_value;
+  bool exchanged;
+};
+atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_2(device atomic_int* atomic, int compare, int value) {
+  int old_value = compare;
+  bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
+  return {old_value, exchanged};
+}
+
+atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_3(threadgroup atomic_uint* atomic, uint compare, uint value) {
+  uint old_value = compare;
+  bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
+  return {old_value, exchanged};
+}
+
+atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_4(threadgroup atomic_int* atomic, int compare, int value) {
+  int old_value = compare;
+  bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
+  return {old_value, exchanged};
+}
+
+void tint_symbol_inner(uint local_invocation_index, threadgroup atomic_uint* const tint_symbol_1, threadgroup atomic_int* const tint_symbol_2, device atomic_uint* const tint_symbol_3, device atomic_int* const tint_symbol_4) {
+  if ((local_invocation_index < 1u)) {
+    atomic_store_explicit(tint_symbol_1, 0u, memory_order_relaxed);
+    atomic_store_explicit(tint_symbol_2, 0, memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  {
+    uint value = 42u;
+    atomic_compare_exchange_resultu32 const r1 = atomicCompareExchangeWeak_1(tint_symbol_3, 0u, value);
+    atomic_compare_exchange_resultu32 const r2 = atomicCompareExchangeWeak_1(tint_symbol_3, 0u, value);
+    atomic_compare_exchange_resultu32 const r3 = atomicCompareExchangeWeak_1(tint_symbol_3, 0u, value);
+  }
+  {
+    int value = 42;
+    atomic_compare_exchange_resulti32 const r1 = atomicCompareExchangeWeak_2(tint_symbol_4, 0, value);
+    atomic_compare_exchange_resulti32 const r2 = atomicCompareExchangeWeak_2(tint_symbol_4, 0, value);
+    atomic_compare_exchange_resulti32 const r3 = atomicCompareExchangeWeak_2(tint_symbol_4, 0, value);
+  }
+  {
+    uint value = 42u;
+    atomic_compare_exchange_resultu32 const r1 = atomicCompareExchangeWeak_3(tint_symbol_1, 0u, value);
+    atomic_compare_exchange_resultu32 const r2 = atomicCompareExchangeWeak_3(tint_symbol_1, 0u, value);
+    atomic_compare_exchange_resultu32 const r3 = atomicCompareExchangeWeak_3(tint_symbol_1, 0u, value);
+  }
+  {
+    int value = 42;
+    atomic_compare_exchange_resulti32 const r1 = atomicCompareExchangeWeak_4(tint_symbol_2, 0, value);
+    atomic_compare_exchange_resulti32 const r2 = atomicCompareExchangeWeak_4(tint_symbol_2, 0, value);
+    atomic_compare_exchange_resulti32 const r3 = atomicCompareExchangeWeak_4(tint_symbol_2, 0, value);
+  }
+}
+
+kernel void tint_symbol(device atomic_uint* tint_symbol_7 [[buffer(0)]], device atomic_int* tint_symbol_8 [[buffer(1)]], uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_5;
+  threadgroup atomic_int tint_symbol_6;
+  tint_symbol_inner(local_invocation_index, &(tint_symbol_5), &(tint_symbol_6), tint_symbol_7, tint_symbol_8);
+  return;
+}
+

test/tint/bug/tint/1574.wgsl.expected.spvasm

@@ -0,0 +1,158 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 118
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %local_invocation_index_1
+               OpExecutionMode %main LocalSize 16 1 1
+               OpName %local_invocation_index_1 "local_invocation_index_1"
+               OpName %a_u32_block "a_u32_block"
+               OpMemberName %a_u32_block 0 "inner"
+               OpName %a_u32 "a_u32"
+               OpName %a_i32_block "a_i32_block"
+               OpMemberName %a_i32_block 0 "inner"
+               OpName %a_i32 "a_i32"
+               OpName %b_u32 "b_u32"
+               OpName %b_i32 "b_i32"
+               OpName %main_inner "main_inner"
+               OpName %local_invocation_index "local_invocation_index"
+               OpName %value "value"
+               OpName %__atomic_compare_exchange_resultu32 "__atomic_compare_exchange_resultu32"
+               OpMemberName %__atomic_compare_exchange_resultu32 0 "old_value"
+               OpMemberName %__atomic_compare_exchange_resultu32 1 "exchanged"
+               OpName %value_0 "value"
+               OpName %__atomic_compare_exchange_resulti32 "__atomic_compare_exchange_resulti32"
+               OpMemberName %__atomic_compare_exchange_resulti32 0 "old_value"
+               OpMemberName %__atomic_compare_exchange_resulti32 1 "exchanged"
+               OpName %value_1 "value"
+               OpName %value_2 "value"
+               OpName %main "main"
+               OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpDecorate %a_u32_block Block
+               OpMemberDecorate %a_u32_block 0 Offset 0
+               OpDecorate %a_u32 DescriptorSet 0
+               OpDecorate %a_u32 Binding 0
+               OpDecorate %a_i32_block Block
+               OpMemberDecorate %a_i32_block 0 Offset 0
+               OpDecorate %a_i32 DescriptorSet 0
+               OpDecorate %a_i32 Binding 1
+               OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0
+               OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4
+               OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0
+               OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
+%a_u32_block = OpTypeStruct %uint
+%_ptr_StorageBuffer_a_u32_block = OpTypePointer StorageBuffer %a_u32_block
+      %a_u32 = OpVariable %_ptr_StorageBuffer_a_u32_block StorageBuffer
+        %int = OpTypeInt 32 1
+%a_i32_block = OpTypeStruct %int
+%_ptr_StorageBuffer_a_i32_block = OpTypePointer StorageBuffer %a_i32_block
+      %a_i32 = OpVariable %_ptr_StorageBuffer_a_i32_block StorageBuffer
+%_ptr_Workgroup_uint = OpTypePointer Workgroup %uint
+      %b_u32 = OpVariable %_ptr_Workgroup_uint Workgroup
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+      %b_i32 = OpVariable %_ptr_Workgroup_int Workgroup
+       %void = OpTypeVoid
+         %15 = OpTypeFunction %void %uint
+     %uint_1 = OpConstant %uint 1
+       %bool = OpTypeBool
+     %uint_2 = OpConstant %uint 2
+     %uint_0 = OpConstant %uint 0
+         %29 = OpConstantNull %uint
+         %32 = OpConstantNull %int
+   %uint_264 = OpConstant %uint 264
+    %uint_42 = OpConstant %uint 42
+%_ptr_Function_uint = OpTypePointer Function %uint
+%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+     %int_42 = OpConstant %int 42
+%_ptr_Function_int = OpTypePointer Function %int
+%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+        %113 = OpTypeFunction %void
+ %main_inner = OpFunction %void None %15
+%local_invocation_index = OpFunctionParameter %uint
+         %19 = OpLabel
+      %value = OpVariable %_ptr_Function_uint Function %29
+    %value_0 = OpVariable %_ptr_Function_int Function %32
+    %value_1 = OpVariable %_ptr_Function_uint Function %29
+    %value_2 = OpVariable %_ptr_Function_int Function %32
+         %21 = OpULessThan %bool %local_invocation_index %uint_1
+               OpSelectionMerge %23 None
+               OpBranchConditional %21 %24 %23
+         %24 = OpLabel
+               OpAtomicStore %b_u32 %uint_2 %uint_0 %29
+               OpAtomicStore %b_i32 %uint_2 %uint_0 %32
+               OpBranch %23
+         %23 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+               OpStore %value %uint_42
+         %42 = OpAccessChain %_ptr_StorageBuffer_uint %a_u32 %uint_0
+         %43 = OpLoad %uint %value
+         %44 = OpAtomicCompareExchange %uint %42 %uint_1 %uint_0 %uint_0 %43 %29
+         %45 = OpIEqual %bool %44 %43
+         %38 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %44 %45
+         %48 = OpAccessChain %_ptr_StorageBuffer_uint %a_u32 %uint_0
+         %49 = OpLoad %uint %value
+         %50 = OpAtomicCompareExchange %uint %48 %uint_1 %uint_0 %uint_0 %49 %29
+         %51 = OpIEqual %bool %50 %49
+         %46 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %50 %51
+         %54 = OpAccessChain %_ptr_StorageBuffer_uint %a_u32 %uint_0
+         %55 = OpLoad %uint %value
+         %56 = OpAtomicCompareExchange %uint %54 %uint_1 %uint_0 %uint_0 %55 %29
+         %57 = OpIEqual %bool %56 %55
+         %52 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %56 %57
+               OpStore %value_0 %int_42
+         %65 = OpAccessChain %_ptr_StorageBuffer_int %a_i32 %uint_0
+         %66 = OpLoad %int %value_0
+         %67 = OpAtomicCompareExchange %int %65 %uint_1 %uint_0 %uint_0 %66 %32
+         %68 = OpIEqual %bool %67 %66
+         %61 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %67 %68
+         %71 = OpAccessChain %_ptr_StorageBuffer_int %a_i32 %uint_0
+         %72 = OpLoad %int %value_0
+         %73 = OpAtomicCompareExchange %int %71 %uint_1 %uint_0 %uint_0 %72 %32
+         %74 = OpIEqual %bool %73 %72
+         %69 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %73 %74
+         %77 = OpAccessChain %_ptr_StorageBuffer_int %a_i32 %uint_0
+         %78 = OpLoad %int %value_0
+         %79 = OpAtomicCompareExchange %int %77 %uint_1 %uint_0 %uint_0 %78 %32
+         %80 = OpIEqual %bool %79 %78
+         %75 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %79 %80
+               OpStore %value_1 %uint_42
+         %84 = OpLoad %uint %value_1
+         %85 = OpAtomicCompareExchange %uint %b_u32 %uint_2 %uint_0 %uint_0 %84 %29
+         %86 = OpIEqual %bool %85 %84
+         %82 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %85 %86
+         %89 = OpLoad %uint %value_1
+         %90 = OpAtomicCompareExchange %uint %b_u32 %uint_2 %uint_0 %uint_0 %89 %29
+         %91 = OpIEqual %bool %90 %89
+         %87 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %90 %91
+         %94 = OpLoad %uint %value_1
+         %95 = OpAtomicCompareExchange %uint %b_u32 %uint_2 %uint_0 %uint_0 %94 %29
+         %96 = OpIEqual %bool %95 %94
+         %92 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %95 %96
+               OpStore %value_2 %int_42
+        %100 = OpLoad %int %value_2
+        %101 = OpAtomicCompareExchange %int %b_i32 %uint_2 %uint_0 %uint_0 %100 %32
+        %102 = OpIEqual %bool %101 %100
+         %98 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %101 %102
+        %105 = OpLoad %int %value_2
+        %106 = OpAtomicCompareExchange %int %b_i32 %uint_2 %uint_0 %uint_0 %105 %32
+        %107 = OpIEqual %bool %106 %105
+        %103 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %106 %107
+        %110 = OpLoad %int %value_2
+        %111 = OpAtomicCompareExchange %int %b_i32 %uint_2 %uint_0 %uint_0 %110 %32
+        %112 = OpIEqual %bool %111 %110
+        %108 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %111 %112
+               OpReturn
+               OpFunctionEnd
+       %main = OpFunction %void None %113
+        %115 = OpLabel
+        %117 = OpLoad %uint %local_invocation_index_1
+        %116 = OpFunctionCall %void %main_inner %117
+               OpReturn
+               OpFunctionEnd

test/tint/bug/tint/1574.wgsl.expected.wgsl

@@ -0,0 +1,35 @@
+@group(0) @binding(0) var<storage, read_write> a_u32 : atomic<u32>;
+
+@group(0) @binding(1) var<storage, read_write> a_i32 : atomic<i32>;
+
+var<workgroup> b_u32 : atomic<u32>;
+
+var<workgroup> b_i32 : atomic<i32>;
+
+@stage(compute) @workgroup_size(16)
+fn main() {
+  {
+    var value = 42u;
+    let r1 = atomicCompareExchangeWeak(&(a_u32), 0u, value);
+    let r2 = atomicCompareExchangeWeak(&(a_u32), 0u, value);
+    let r3 = atomicCompareExchangeWeak(&(a_u32), 0u, value);
+  }
+  {
+    var value = 42;
+    let r1 = atomicCompareExchangeWeak(&(a_i32), 0, value);
+    let r2 = atomicCompareExchangeWeak(&(a_i32), 0, value);
+    let r3 = atomicCompareExchangeWeak(&(a_i32), 0, value);
+  }
+  {
+    var value = 42u;
+    let r1 = atomicCompareExchangeWeak(&(b_u32), 0u, value);
+    let r2 = atomicCompareExchangeWeak(&(b_u32), 0u, value);
+    let r3 = atomicCompareExchangeWeak(&(b_u32), 0u, value);
+  }
+  {
+    var value = 42;
+    let r1 = atomicCompareExchangeWeak(&(b_i32), 0, value);
+    let r2 = atomicCompareExchangeWeak(&(b_i32), 0, value);
+    let r3 = atomicCompareExchangeWeak(&(b_i32), 0, value);
+  }
+}

test/tint/builtins/gen/literal/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl

@@ -6,9 +6,8 @@ struct atomic_compare_exchange_resulti32 {
   int old_value;
   bool exchanged;
 };
-template <typename A, typename T>
+atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(device atomic_int* atomic, int compare, int value) {
-atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(device A* atomic, T compare, T value) {
+  int old_value = compare;
-  T old_value = compare;
   bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
   return {old_value, exchanged};
 }

test/tint/builtins/gen/literal/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl

@@ -6,9 +6,8 @@ struct atomic_compare_exchange_resultu32 {
   uint old_value;
   bool exchanged;
 };
-template <typename A, typename T>
+atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) {
-atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device A* atomic, T compare, T value) {
+  uint old_value = compare;
-  T old_value = compare;
   bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
   return {old_value, exchanged};
 }

test/tint/builtins/gen/literal/atomicCompareExchangeWeak/83580d.wgsl.expected.msl

@@ -6,9 +6,8 @@ struct atomic_compare_exchange_resultu32 {
   uint old_value;
   bool exchanged;
 };
-template <typename A, typename T>
+atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(threadgroup atomic_uint* atomic, uint compare, uint value) {
-atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) {
+  uint old_value = compare;
-  T old_value = compare;
   bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
   return {old_value, exchanged};
 }

test/tint/builtins/gen/literal/atomicCompareExchangeWeak/e88938.wgsl.expected.msl

@@ -6,9 +6,8 @@ struct atomic_compare_exchange_resulti32 {
   int old_value;
   bool exchanged;
 };
-template <typename A, typename T>
+atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(threadgroup atomic_int* atomic, int compare, int value) {
-atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) {
+  int old_value = compare;
-  T old_value = compare;
   bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
   return {old_value, exchanged};
 }

test/tint/builtins/gen/var/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl

@@ -6,9 +6,8 @@ struct atomic_compare_exchange_resulti32 {
   int old_value;
   bool exchanged;
 };
-template <typename A, typename T>
+atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(device atomic_int* atomic, int compare, int value) {
-atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(device A* atomic, T compare, T value) {
+  int old_value = compare;
-  T old_value = compare;
   bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
   return {old_value, exchanged};
 }

test/tint/builtins/gen/var/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl

@@ -6,9 +6,8 @@ struct atomic_compare_exchange_resultu32 {
   uint old_value;
   bool exchanged;
 };
-template <typename A, typename T>
+atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device atomic_uint* atomic, uint compare, uint value) {
-atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device A* atomic, T compare, T value) {
+  uint old_value = compare;
-  T old_value = compare;
   bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
   return {old_value, exchanged};
 }

test/tint/builtins/gen/var/atomicCompareExchangeWeak/83580d.wgsl.expected.msl

@@ -6,9 +6,8 @@ struct atomic_compare_exchange_resultu32 {
   uint old_value;
   bool exchanged;
 };
-template <typename A, typename T>
+atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(threadgroup atomic_uint* atomic, uint compare, uint value) {
-atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) {
+  uint old_value = compare;
-  T old_value = compare;
   bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
   return {old_value, exchanged};
 }

test/tint/builtins/gen/var/atomicCompareExchangeWeak/e88938.wgsl.expected.msl

@@ -6,9 +6,8 @@ struct atomic_compare_exchange_resulti32 {
   int old_value;
   bool exchanged;
 };
-template <typename A, typename T>
+atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(threadgroup atomic_int* atomic, int compare, int value) {
-atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) {
+  int old_value = compare;
-  T old_value = compare;
   bool exchanged = atomic_compare_exchange_weak_explicit(atomic, &old_value, value, memory_order_relaxed, memory_order_relaxed);
   return {old_value, exchanged};
 }