Implement atomicCompareExchangeWeak returning struct instead of vec2

Also fixed implementation of this atomic in GLSL. It was emitting code
that would not compile because, as for HLSL, we must pass in the
variable directly to atomic funcs, not via an in/out arg to a function.

Bug: tint:1185
Change-Id: Id0e9f99d6368717511ef3a94473634c512e10cb8
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/91881
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>

src/tint/intrinsics.def

@@ -118,6 +118,8 @@ type __modf_result
 type __frexp_result
 [[display("__frexp_result_vec{N}")]] type __frexp_result_vec<N: num>
 
+type __atomic_compare_exchange_result<T>
+
 ////////////////////////////////////////////////////////////////////////////////
 // Type matchers                                                              //
 //                                                                            //
@@ -603,7 +605,7 @@ fn textureLoad(texture: texture_external, coords: vec2<i32>) -> vec4<f32>
 [[stage("fragment", "compute")]] fn atomicOr<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T
 [[stage("fragment", "compute")]] fn atomicXor<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T
 [[stage("fragment", "compute")]] fn atomicExchange<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T) -> T
-[[stage("fragment", "compute")]] fn atomicCompareExchangeWeak<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T, T) -> vec2<T>
+[[stage("fragment", "compute")]] fn atomicCompareExchangeWeak<T: iu32, S: workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T, T) -> __atomic_compare_exchange_result<T>
 
 ////////////////////////////////////////////////////////////////////////////////
 // Type constructors                                                          //

src/tint/resolver/intrinsic_table.cc

@@ -722,6 +722,14 @@ bool match_frexp_result_vec(const sem::Type* ty, Number& N) {
     return true;
 }
 
+bool match_atomic_compare_exchange_result(const sem::Type* ty, const sem::Type*& T) {
+    if (ty->Is<Any>()) {
+        T = ty;
+        return true;
+    }
+    return false;
+}
+
 struct NameAndType {
     std::string name;
     sem::Type* type;
@@ -779,6 +787,13 @@ const sem::Struct* build_frexp_result_vec(MatchState& state, Number& n) {
                         {{"sig", vec_f32}, {"exp", vec_i32}});
 }
 
+const sem::Struct* build_atomic_compare_exchange_result(MatchState& state, const sem::Type* ty) {
+    return build_struct(
+        state, "__atomic_compare_exchange_result" + ty->FriendlyName(state.builder.Symbols()),
+        {{"old_value", const_cast<sem::Type*>(ty)},
+         {"exchanged", state.builder.create<sem::Bool>()}});
+}
+
 /// ParameterInfo describes a parameter
 struct ParameterInfo {
     /// The parameter usage (parameter name in definition file)

src/tint/transform/decompose_memory_access.cc

@@ -644,14 +644,34 @@ struct DecomposeMemoryAccess::State {
                     << el_ty->TypeInfo().name;
             }
 
-            auto* ret_ty = CreateASTTypeFor(ctx, intrinsic->ReturnType());
+            const ast::Type* ret_ty = nullptr;
-            auto* func =
+
-                b.create<ast::Function>(b.Sym(), params, ret_ty, nullptr,
+            // For intrinsics that return a struct, there is no AST node for it, so create one now.
-                                        ast::AttributeList{
+            if (intrinsic->Type() == sem::BuiltinType::kAtomicCompareExchangeWeak) {
-                                            atomic,
+                auto* str = intrinsic->ReturnType()->As<sem::Struct>();
-                                            b.Disable(ast::DisabledValidation::kFunctionHasNoBody),
+                TINT_ASSERT(Transform, str && str->Declaration() == nullptr);
-                                        },
+
-                                        ast::AttributeList{});
+                ast::StructMemberList ast_members;
+                ast_members.reserve(str->Members().size());
+                for (auto& m : str->Members()) {
+                    ast_members.push_back(
+                        b.Member(ctx.Clone(m->Name()), CreateASTTypeFor(ctx, m->Type())));
+                }
+
+                auto name = b.Symbols().New("atomic_compare_exchange_weak_ret_type");
+                auto* new_str = b.Structure(name, std::move(ast_members));
+                ret_ty = b.ty.Of(new_str);
+            } else {
+                ret_ty = CreateASTTypeFor(ctx, intrinsic->ReturnType());
+            }
+
+            auto* func = b.create<ast::Function>(
+                b.Symbols().New(std::string{"tint_"} + intrinsic->str()), params, ret_ty, nullptr,
+                ast::AttributeList{
+                    atomic,
+                    b.Disable(ast::DisabledValidation::kFunctionHasNoBody),
+                },
+                ast::AttributeList{});
 
             b.AST().AddFunction(func);
             return func->symbol;
@@ -753,6 +773,10 @@ const DecomposeMemoryAccess::Intrinsic* DecomposeMemoryAccess::Intrinsic::Clone(
                                                                          storage_class, type);
 }
 
+bool DecomposeMemoryAccess::Intrinsic::IsAtomic() const {
+    return op != Op::kLoad && op != Op::kStore;
+}
+
 DecomposeMemoryAccess::DecomposeMemoryAccess() = default;
 DecomposeMemoryAccess::~DecomposeMemoryAccess() = default;
 

src/tint/transform/decompose_memory_access.h

@@ -89,6 +89,9 @@ class DecomposeMemoryAccess final : public Castable<DecomposeMemoryAccess, Trans
         /// @return the newly cloned object
         const Intrinsic* Clone(CloneContext* ctx) const override;
 
+        /// @return true if op is atomic
+        bool IsAtomic() const;
+
         /// The op of the intrinsic
         const Op op;
 

src/tint/transform/decompose_memory_access_test.cc

@@ -2467,95 +2467,105 @@ struct SB {
 @group(0) @binding(0) var<storage, read_write> sb : SB;
 
 @internal(intrinsic_atomic_store_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32)
+fn tint_atomicStore(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32)
 
 @internal(intrinsic_atomic_load_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> i32
+fn tint_atomicLoad(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> i32
 
 @internal(intrinsic_atomic_add_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_2(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+fn tint_atomicAdd(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
 
 @internal(intrinsic_atomic_sub_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+fn tint_atomicSub(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
 
 @internal(intrinsic_atomic_max_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_4(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+fn tint_atomicMax(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
 
 @internal(intrinsic_atomic_min_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_5(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+fn tint_atomicMin(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
 
 @internal(intrinsic_atomic_and_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+fn tint_atomicAnd(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
 
 @internal(intrinsic_atomic_or_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_7(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+fn tint_atomicOr(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
 
 @internal(intrinsic_atomic_xor_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_8(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+fn tint_atomicXor(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
 
 @internal(intrinsic_atomic_exchange_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_9(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+fn tint_atomicExchange(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+
+struct atomic_compare_exchange_weak_ret_type {
+  old_value : i32,
+  exchanged : bool,
+}
 
 @internal(intrinsic_atomic_compare_exchange_weak_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_10(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32, param_2 : i32) -> vec2<i32>
+fn tint_atomicCompareExchangeWeak(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32, param_2 : i32) -> atomic_compare_exchange_weak_ret_type
 
 @internal(intrinsic_atomic_store_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_11(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32)
+fn tint_atomicStore_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32)
 
 @internal(intrinsic_atomic_load_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> u32
+fn tint_atomicLoad_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> u32
 
 @internal(intrinsic_atomic_add_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+fn tint_atomicAdd_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
 
 @internal(intrinsic_atomic_sub_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_14(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+fn tint_atomicSub_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
 
 @internal(intrinsic_atomic_max_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_15(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+fn tint_atomicMax_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
 
 @internal(intrinsic_atomic_min_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_16(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+fn tint_atomicMin_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
 
 @internal(intrinsic_atomic_and_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+fn tint_atomicAnd_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
 
 @internal(intrinsic_atomic_or_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+fn tint_atomicOr_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
 
 @internal(intrinsic_atomic_xor_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+fn tint_atomicXor_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
 
 @internal(intrinsic_atomic_exchange_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+fn tint_atomicExchange_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+
+struct atomic_compare_exchange_weak_ret_type_1 {
+  old_value : u32,
+  exchanged : bool,
+}
 
 @internal(intrinsic_atomic_compare_exchange_weak_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32, param_2 : u32) -> vec2<u32>
+fn tint_atomicCompareExchangeWeak_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32, param_2 : u32) -> atomic_compare_exchange_weak_ret_type_1
 
 @stage(compute) @workgroup_size(1)
 fn main() {
-  tint_symbol(sb, 16u, 123);
+  tint_atomicStore(sb, 16u, 123);
-  tint_symbol_1(sb, 16u);
+  tint_atomicLoad(sb, 16u);
-  tint_symbol_2(sb, 16u, 123);
+  tint_atomicAdd(sb, 16u, 123);
-  tint_symbol_3(sb, 16u, 123);
+  tint_atomicSub(sb, 16u, 123);
-  tint_symbol_4(sb, 16u, 123);
+  tint_atomicMax(sb, 16u, 123);
-  tint_symbol_5(sb, 16u, 123);
+  tint_atomicMin(sb, 16u, 123);
-  tint_symbol_6(sb, 16u, 123);
+  tint_atomicAnd(sb, 16u, 123);
-  tint_symbol_7(sb, 16u, 123);
+  tint_atomicOr(sb, 16u, 123);
-  tint_symbol_8(sb, 16u, 123);
+  tint_atomicXor(sb, 16u, 123);
-  tint_symbol_9(sb, 16u, 123);
+  tint_atomicExchange(sb, 16u, 123);
-  tint_symbol_10(sb, 16u, 123, 345);
+  tint_atomicCompareExchangeWeak(sb, 16u, 123, 345);
-  tint_symbol_11(sb, 20u, 123u);
+  tint_atomicStore_1(sb, 20u, 123u);
-  tint_symbol_12(sb, 20u);
+  tint_atomicLoad_1(sb, 20u);
-  tint_symbol_13(sb, 20u, 123u);
+  tint_atomicAdd_1(sb, 20u, 123u);
-  tint_symbol_14(sb, 20u, 123u);
+  tint_atomicSub_1(sb, 20u, 123u);
-  tint_symbol_15(sb, 20u, 123u);
+  tint_atomicMax_1(sb, 20u, 123u);
-  tint_symbol_16(sb, 20u, 123u);
+  tint_atomicMin_1(sb, 20u, 123u);
-  tint_symbol_17(sb, 20u, 123u);
+  tint_atomicAnd_1(sb, 20u, 123u);
-  tint_symbol_18(sb, 20u, 123u);
+  tint_atomicOr_1(sb, 20u, 123u);
-  tint_symbol_19(sb, 20u, 123u);
+  tint_atomicXor_1(sb, 20u, 123u);
-  tint_symbol_20(sb, 20u, 123u);
+  tint_atomicExchange_1(sb, 20u, 123u);
-  tint_symbol_21(sb, 20u, 123u, 345u);
+  tint_atomicCompareExchangeWeak_1(sb, 20u, 123u, 345u);
 }
 )";
 
@@ -2604,95 +2614,105 @@ struct SB {
 
     auto* expect = R"(
 @internal(intrinsic_atomic_store_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32)
+fn tint_atomicStore(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32)
 
 @internal(intrinsic_atomic_load_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> i32
+fn tint_atomicLoad(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> i32
 
 @internal(intrinsic_atomic_add_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_2(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+fn tint_atomicAdd(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
 
 @internal(intrinsic_atomic_sub_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+fn tint_atomicSub(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
 
 @internal(intrinsic_atomic_max_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_4(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+fn tint_atomicMax(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
 
 @internal(intrinsic_atomic_min_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_5(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+fn tint_atomicMin(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
 
 @internal(intrinsic_atomic_and_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+fn tint_atomicAnd(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
 
 @internal(intrinsic_atomic_or_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_7(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+fn tint_atomicOr(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
 
 @internal(intrinsic_atomic_xor_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_8(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+fn tint_atomicXor(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
 
 @internal(intrinsic_atomic_exchange_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_9(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+fn tint_atomicExchange(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32) -> i32
+
+struct atomic_compare_exchange_weak_ret_type {
+  old_value : i32,
+  exchanged : bool,
+}
 
 @internal(intrinsic_atomic_compare_exchange_weak_storage_i32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_10(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32, param_2 : i32) -> vec2<i32>
+fn tint_atomicCompareExchangeWeak(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : i32, param_2 : i32) -> atomic_compare_exchange_weak_ret_type
 
 @internal(intrinsic_atomic_store_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_11(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32)
+fn tint_atomicStore_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32)
 
 @internal(intrinsic_atomic_load_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_12(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> u32
+fn tint_atomicLoad_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32) -> u32
 
 @internal(intrinsic_atomic_add_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_13(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+fn tint_atomicAdd_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
 
 @internal(intrinsic_atomic_sub_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_14(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+fn tint_atomicSub_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
 
 @internal(intrinsic_atomic_max_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_15(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+fn tint_atomicMax_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
 
 @internal(intrinsic_atomic_min_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_16(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+fn tint_atomicMin_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
 
 @internal(intrinsic_atomic_and_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_17(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+fn tint_atomicAnd_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
 
 @internal(intrinsic_atomic_or_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_18(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+fn tint_atomicOr_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
 
 @internal(intrinsic_atomic_xor_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_19(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+fn tint_atomicXor_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
 
 @internal(intrinsic_atomic_exchange_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_20(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+fn tint_atomicExchange_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32) -> u32
+
+struct atomic_compare_exchange_weak_ret_type_1 {
+  old_value : u32,
+  exchanged : bool,
+}
 
 @internal(intrinsic_atomic_compare_exchange_weak_storage_u32) @internal(disable_validation__function_has_no_body)
-fn tint_symbol_21(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32, param_2 : u32) -> vec2<u32>
+fn tint_atomicCompareExchangeWeak_1(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, offset : u32, param_1 : u32, param_2 : u32) -> atomic_compare_exchange_weak_ret_type_1
 
 @stage(compute) @workgroup_size(1)
 fn main() {
-  tint_symbol(sb, 16u, 123);
+  tint_atomicStore(sb, 16u, 123);
-  tint_symbol_1(sb, 16u);
+  tint_atomicLoad(sb, 16u);
-  tint_symbol_2(sb, 16u, 123);
+  tint_atomicAdd(sb, 16u, 123);
-  tint_symbol_3(sb, 16u, 123);
+  tint_atomicSub(sb, 16u, 123);
-  tint_symbol_4(sb, 16u, 123);
+  tint_atomicMax(sb, 16u, 123);
-  tint_symbol_5(sb, 16u, 123);
+  tint_atomicMin(sb, 16u, 123);
-  tint_symbol_6(sb, 16u, 123);
+  tint_atomicAnd(sb, 16u, 123);
-  tint_symbol_7(sb, 16u, 123);
+  tint_atomicOr(sb, 16u, 123);
-  tint_symbol_8(sb, 16u, 123);
+  tint_atomicXor(sb, 16u, 123);
-  tint_symbol_9(sb, 16u, 123);
+  tint_atomicExchange(sb, 16u, 123);
-  tint_symbol_10(sb, 16u, 123, 345);
+  tint_atomicCompareExchangeWeak(sb, 16u, 123, 345);
-  tint_symbol_11(sb, 20u, 123u);
+  tint_atomicStore_1(sb, 20u, 123u);
-  tint_symbol_12(sb, 20u);
+  tint_atomicLoad_1(sb, 20u);
-  tint_symbol_13(sb, 20u, 123u);
+  tint_atomicAdd_1(sb, 20u, 123u);
-  tint_symbol_14(sb, 20u, 123u);
+  tint_atomicSub_1(sb, 20u, 123u);
-  tint_symbol_15(sb, 20u, 123u);
+  tint_atomicMax_1(sb, 20u, 123u);
-  tint_symbol_16(sb, 20u, 123u);
+  tint_atomicMin_1(sb, 20u, 123u);
-  tint_symbol_17(sb, 20u, 123u);
+  tint_atomicAnd_1(sb, 20u, 123u);
-  tint_symbol_18(sb, 20u, 123u);
+  tint_atomicOr_1(sb, 20u, 123u);
-  tint_symbol_19(sb, 20u, 123u);
+  tint_atomicXor_1(sb, 20u, 123u);
-  tint_symbol_20(sb, 20u, 123u);
+  tint_atomicExchange_1(sb, 20u, 123u);
-  tint_symbol_21(sb, 20u, 123u, 345u);
+  tint_atomicCompareExchangeWeak_1(sb, 20u, 123u, 345u);
 }
 
 @group(0) @binding(0) var<storage, read_write> sb : SB;

src/tint/transform/manager.cc

@@ -49,7 +49,7 @@ Output Manager::Run(const Program* program, const DataMap& data) const {
     Output out;
     for (const auto& transform : transforms_) {
         if (!transform->ShouldRun(in, data)) {
-            TINT_IF_PRINT_PROGRAM(std::cout << "Skipping " << transform->TypeInfo().name);
+            TINT_IF_PRINT_PROGRAM(std::cout << "Skipping " << transform->TypeInfo().name << std::endl);
             continue;
         }
         TINT_IF_PRINT_PROGRAM(print_program("Input to", transform.get()));

src/tint/writer/glsl/generator_impl.cc

@@ -911,39 +911,56 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
             return true;
         }
         case sem::BuiltinType::kAtomicCompareExchangeWeak: {
-            return CallBuiltinHelper(
+            // Emit the builtin return type unique to this overload. This does not
-                out, expr, builtin, [&](TextBuffer* b, const std::vector<std::string>& params) {
+            // exist in the AST, so it will not be generated in Generate().
-                    {
+            if (!EmitStructType(&helpers_, builtin->ReturnType()->As<sem::Struct>())) {
-                        auto pre = line(b);
+                return false;
-                        if (!EmitTypeAndName(pre, builtin->ReturnType(), ast::StorageClass::kNone,
+            }
-                                             ast::Access::kUndefined, "result")) {
+
-                            return false;
+            auto* dest = expr->args[0];
-                        }
+            auto* compare_value = expr->args[1];
-                        pre << ";";
+            auto* value = expr->args[2];
+
+            std::string result = UniqueIdentifier("atomic_compare_result");
+
+            {
+                auto pre = line();
+                if (!EmitTypeAndName(pre, builtin->ReturnType(), ast::StorageClass::kNone,
+                                     ast::Access::kUndefined, result)) {
+                    return false;
+                }
+                pre << ";";
+            }
+            {
+                auto pre = line();
+                pre << result << ".old_value = atomicCompSwap";
+                {
+                    ScopedParen sp(pre);
+                    if (!EmitExpression(pre, dest)) {
+                        return false;
                     }
-                    {
+                    pre << ", ";
-                        auto pre = line(b);
+                    if (!EmitExpression(pre, compare_value)) {
-                        pre << "result.x = atomicCompSwap";
+                        return false;
-                        {
-                            ScopedParen sp(pre);
-                            pre << params[0];
-                            pre << ", " << params[1];
-                            pre << ", " << params[2];
-                        }
-                        pre << ";";
                     }
-                    {
+                    pre << ", ";
-                        auto pre = line(b);
+                    if (!EmitExpression(pre, value)) {
-                        pre << "result.y = result.x == " << params[2] << " ? ";
+                        return false;
-                        if (TypeOf(expr->args[2])->Is<sem::U32>()) {
-                            pre << "1u : 0u;";
-                        } else {
-                            pre << "1 : 0;";
-                        }
                     }
-                    line(b) << "return result;";
+                }
-                    return true;
+                pre << ";";
-                });
+            }
+            {
+                auto pre = line();
+                pre << result << ".exchanged = " << result << ".old_value == ";
+                if (!EmitExpression(pre, compare_value)) {
+                    return false;
+                }
+                pre << ";";
+            }
+
+            out << result;
+            return true;
         }
 
         case sem::BuiltinType::kAtomicAdd:

src/tint/writer/glsl/generator_impl.h

@@ -174,14 +174,6 @@ class GeneratorImpl : public TextGenerator {
     /// @param builtin the semantic information for the barrier builtin
     /// @returns true if the call expression is emitted
     bool EmitBarrierCall(std::ostream& out, const sem::Builtin* builtin);
-    /// Handles generating an atomic intrinsic call for a storage buffer variable
-    /// @param out the output of the expression stream
-    /// @param expr the call expression
-    /// @param intrinsic the atomic intrinsic
-    /// @returns true if the call expression is emitted
-    bool EmitStorageAtomicCall(std::ostream& out,
-                               const ast::CallExpression* expr,
-                               const transform::DecomposeMemoryAccess::Intrinsic* intrinsic);
     /// Handles generating an atomic builtin call for a workgroup variable
     /// @param out the output of the expression stream
     /// @param expr the call expression

src/tint/writer/hlsl/generator_impl.cc

@@ -968,7 +968,10 @@ bool GeneratorImpl::EmitFunctionCall(std::ostream& out,
             case ast::StorageClass::kUniform:
                 return EmitUniformBufferAccess(out, expr, intrinsic);
             case ast::StorageClass::kStorage:
-                return EmitStorageBufferAccess(out, expr, intrinsic);
+                if (!intrinsic->IsAtomic()) {
+                    return EmitStorageBufferAccess(out, expr, intrinsic);
+                }
+                break;
             default:
                 TINT_UNREACHABLE(Writer, diagnostics_)
                     << "unsupported DecomposeMemoryAccess::Intrinsic storage class:"
@@ -1445,19 +1448,10 @@ bool GeneratorImpl::EmitStorageBufferAccess(
                 << static_cast<int>(intrinsic->type);
             return false;
         }
-
+        default:
-        case Op::kAtomicLoad:
+            // Break out to error case below/
-        case Op::kAtomicStore:
+            // Note that atomic intrinsics are generated as functions.
-        case Op::kAtomicAdd:
+            break;
-        case Op::kAtomicSub:
-        case Op::kAtomicMax:
-        case Op::kAtomicMin:
-        case Op::kAtomicAnd:
-        case Op::kAtomicOr:
-        case Op::kAtomicXor:
-        case Op::kAtomicExchange:
-        case Op::kAtomicCompareExchangeWeak:
-            return EmitStorageAtomicCall(out, expr, intrinsic);
     }
 
     TINT_UNREACHABLE(Writer, diagnostics_)
@@ -1465,32 +1459,127 @@ bool GeneratorImpl::EmitStorageBufferAccess(
     return false;
 }
 
-bool GeneratorImpl::EmitStorageAtomicCall(
+bool GeneratorImpl::EmitStorageAtomicIntrinsic(
-    std::ostream& out,
+    const ast::Function* func,
-    const ast::CallExpression* expr,
     const transform::DecomposeMemoryAccess::Intrinsic* intrinsic) {
     using Op = transform::DecomposeMemoryAccess::Intrinsic::Op;
 
-    auto* result_ty = TypeOf(expr);
+    const sem::Function* sem_func = builder_.Sem().Get(func);
+    auto* result_ty = sem_func->ReturnType();
+    const auto& params = sem_func->Parameters();
+    const auto name = builder_.Symbols().NameFor(func->symbol);
+    auto& buf = *current_buffer_;
 
-    auto& buf = helpers_;
+    auto rmw = [&](const char* hlsl) -> bool {
+        {
+            auto fn = line(&buf);
+            if (!EmitTypeAndName(fn, result_ty, ast::StorageClass::kNone, ast::Access::kUndefined,
+                                 name)) {
+                return false;
+            }
+            fn << "(RWByteAddressBuffer buffer, uint offset, ";
+            if (!EmitTypeAndName(fn, result_ty, ast::StorageClass::kNone, ast::Access::kUndefined,
+                                 "value")) {
+                return false;
+            }
+            fn << ") {";
+        }
 
-    // generate_helper() generates a helper function that translates the
+        buf.IncrementIndent();
-    // DecomposeMemoryAccess::Intrinsic call into the corresponding HLSL
+        TINT_DEFER({
-    // atomic intrinsic function.
+            buf.DecrementIndent();
-    auto generate_helper = [&]() -> std::string {
+            line(&buf) << "}";
-        auto rmw = [&](const char* wgsl, const char* hlsl) -> std::string {
+            line(&buf);
-            auto name = UniqueIdentifier(wgsl);
+        });
+
+        {
+            auto l = line(&buf);
+            if (!EmitTypeAndName(l, result_ty, ast::StorageClass::kNone, ast::Access::kUndefined,
+                                 "original_value")) {
+                return false;
+            }
+            l << " = 0;";
+        }
+        {
+            auto l = line(&buf);
+            l << "buffer." << hlsl << "(offset, ";
+            if (intrinsic->op == Op::kAtomicSub) {
+                l << "-";
+            }
+            l << "value, original_value);";
+        }
+        line(&buf) << "return original_value;";
+        return true;
+    };
+
+    switch (intrinsic->op) {
+        case Op::kAtomicAdd:
+            return rmw("InterlockedAdd");
+
+        case Op::kAtomicSub:
+            // Use add with the operand negated.
+            return rmw("InterlockedAdd");
+
+        case Op::kAtomicMax:
+            return rmw("InterlockedMax");
+
+        case Op::kAtomicMin:
+            return rmw("InterlockedMin");
+
+        case Op::kAtomicAnd:
+            return rmw("InterlockedAnd");
+
+        case Op::kAtomicOr:
+            return rmw("InterlockedOr");
+
+        case Op::kAtomicXor:
+            return rmw("InterlockedXor");
+
+        case Op::kAtomicExchange:
+            return rmw("InterlockedExchange");
+
+        case Op::kAtomicLoad: {
+            // HLSL does not have an InterlockedLoad, so we emulate it with
+            // InterlockedOr using 0 as the OR value
             {
                 auto fn = line(&buf);
                 if (!EmitTypeAndName(fn, result_ty, ast::StorageClass::kNone,
                                      ast::Access::kUndefined, name)) {
-                    return "";
+                    return false;
                 }
-                fn << "(RWByteAddressBuffer buffer, uint offset, ";
+                fn << "(RWByteAddressBuffer buffer, uint offset) {";
-                if (!EmitTypeAndName(fn, result_ty, ast::StorageClass::kNone,
+            }
+
+            buf.IncrementIndent();
+            TINT_DEFER({
+                buf.DecrementIndent();
+                line(&buf) << "}";
+                line(&buf);
+            });
+
+            {
+                auto l = line(&buf);
+                if (!EmitTypeAndName(l, result_ty, ast::StorageClass::kNone,
                                      ast::Access::kUndefined, "value")) {
-                    return "";
+                    return false;
+                }
+                l << " = 0;";
+            }
+
+            line(&buf) << "buffer.InterlockedOr(offset, 0, value);";
+            line(&buf) << "return value;";
+            return true;
+        }
+        case Op::kAtomicStore: {
+            // HLSL does not have an InterlockedStore, so we emulate it with
+            // InterlockedExchange and discard the returned value
+            auto* value_ty = params[2]->Type()->UnwrapRef();
+            {
+                auto fn = line(&buf);
+                fn << "void " << name << "(RWByteAddressBuffer buffer, uint offset, ";
+                if (!EmitTypeAndName(fn, value_ty, ast::StorageClass::kNone,
+                                     ast::Access::kUndefined, "value")) {
+                    return false;
                 }
                 fn << ") {";
             }
@@ -1504,191 +1593,73 @@ bool GeneratorImpl::EmitStorageAtomicCall(
 
             {
                 auto l = line(&buf);
-                if (!EmitTypeAndName(l, result_ty, ast::StorageClass::kNone,
+                if (!EmitTypeAndName(l, value_ty, ast::StorageClass::kNone, ast::Access::kUndefined,
-                                     ast::Access::kUndefined, "original_value")) {
+                                     "ignored")) {
-                    return "";
+                    return false;
                 }
-                l << " = 0;";
+                l << ";";
             }
+            line(&buf) << "buffer.InterlockedExchange(offset, value, ignored);";
+            return true;
+        }
+        case Op::kAtomicCompareExchangeWeak: {
+            // NOTE: We don't need to emit the return type struct here as DecomposeMemoryAccess
+            // already added it to the AST, and it should have already been emitted by now.
+            auto* value_ty = params[2]->Type()->UnwrapRef();
             {
+                auto fn = line(&buf);
+                if (!EmitTypeAndName(fn, result_ty, ast::StorageClass::kNone,
+                                     ast::Access::kUndefined, name)) {
+                    return false;
+                }
+                fn << "(RWByteAddressBuffer buffer, uint offset, ";
+                if (!EmitTypeAndName(fn, value_ty, ast::StorageClass::kNone,
+                                     ast::Access::kUndefined, "compare")) {
+                    return false;
+                }
+                fn << ", ";
+                if (!EmitTypeAndName(fn, value_ty, ast::StorageClass::kNone,
+                                     ast::Access::kUndefined, "value")) {
+                    return false;
+                }
+                fn << ") {";
+            }
+
+            buf.IncrementIndent();
+            TINT_DEFER({
+                buf.DecrementIndent();
+                line(&buf) << "}";
+                line(&buf);
+            });
+
+            {  // T result = {0};
                 auto l = line(&buf);
-                l << "buffer." << hlsl << "(offset, ";
+                if (!EmitTypeAndName(l, result_ty, ast::StorageClass::kNone,
-                if (intrinsic->op == Op::kAtomicSub) {
+                                     ast::Access::kUndefined, "result")) {
-                    l << "-";
+                    return false;
                 }
-                l << "value, original_value);";
+                l << "=";
+                if (!EmitZeroValue(l, result_ty)) {
+                    return false;
+                }
+                l << ";";
             }
-            line(&buf) << "return original_value;";
-            return name;
-        };
 
-        switch (intrinsic->op) {
+            line(&buf) << "buffer.InterlockedCompareExchange(offset, compare, value, "
-            case Op::kAtomicAdd:
+                          "result.old_value);";
-                return rmw("atomicAdd", "InterlockedAdd");
+            line(&buf) << "result.exchanged = result.old_value == compare;";
+            line(&buf) << "return result;";
 
-            case Op::kAtomicSub:
+            return true;
-                // Use add with the operand negated.
-                return rmw("atomicSub", "InterlockedAdd");
-
-            case Op::kAtomicMax:
-                return rmw("atomicMax", "InterlockedMax");
-
-            case Op::kAtomicMin:
-                return rmw("atomicMin", "InterlockedMin");
-
-            case Op::kAtomicAnd:
-                return rmw("atomicAnd", "InterlockedAnd");
-
-            case Op::kAtomicOr:
-                return rmw("atomicOr", "InterlockedOr");
-
-            case Op::kAtomicXor:
-                return rmw("atomicXor", "InterlockedXor");
-
-            case Op::kAtomicExchange:
-                return rmw("atomicExchange", "InterlockedExchange");
-
-            case Op::kAtomicLoad: {
-                // HLSL does not have an InterlockedLoad, so we emulate it with
-                // InterlockedOr using 0 as the OR value
-                auto name = UniqueIdentifier("atomicLoad");
-                {
-                    auto fn = line(&buf);
-                    if (!EmitTypeAndName(fn, result_ty, ast::StorageClass::kNone,
-                                         ast::Access::kUndefined, name)) {
-                        return "";
-                    }
-                    fn << "(RWByteAddressBuffer buffer, uint offset) {";
-                }
-
-                buf.IncrementIndent();
-                TINT_DEFER({
-                    buf.DecrementIndent();
-                    line(&buf) << "}";
-                    line(&buf);
-                });
-
-                {
-                    auto l = line(&buf);
-                    if (!EmitTypeAndName(l, result_ty, ast::StorageClass::kNone,
-                                         ast::Access::kUndefined, "value")) {
-                        return "";
-                    }
-                    l << " = 0;";
-                }
-
-                line(&buf) << "buffer.InterlockedOr(offset, 0, value);";
-                line(&buf) << "return value;";
-                return name;
-            }
-            case Op::kAtomicStore: {
-                // HLSL does not have an InterlockedStore, so we emulate it with
-                // InterlockedExchange and discard the returned value
-                auto* value_ty = TypeOf(expr->args[2])->UnwrapRef();
-                auto name = UniqueIdentifier("atomicStore");
-                {
-                    auto fn = line(&buf);
-                    fn << "void " << name << "(RWByteAddressBuffer buffer, uint offset, ";
-                    if (!EmitTypeAndName(fn, value_ty, ast::StorageClass::kNone,
-                                         ast::Access::kUndefined, "value")) {
-                        return "";
-                    }
-                    fn << ") {";
-                }
-
-                buf.IncrementIndent();
-                TINT_DEFER({
-                    buf.DecrementIndent();
-                    line(&buf) << "}";
-                    line(&buf);
-                });
-
-                {
-                    auto l = line(&buf);
-                    if (!EmitTypeAndName(l, value_ty, ast::StorageClass::kNone,
-                                         ast::Access::kUndefined, "ignored")) {
-                        return "";
-                    }
-                    l << ";";
-                }
-                line(&buf) << "buffer.InterlockedExchange(offset, value, ignored);";
-                return name;
-            }
-            case Op::kAtomicCompareExchangeWeak: {
-                auto* value_ty = TypeOf(expr->args[2])->UnwrapRef();
-
-                auto name = UniqueIdentifier("atomicCompareExchangeWeak");
-                {
-                    auto fn = line(&buf);
-                    if (!EmitTypeAndName(fn, result_ty, ast::StorageClass::kNone,
-                                         ast::Access::kUndefined, name)) {
-                        return "";
-                    }
-                    fn << "(RWByteAddressBuffer buffer, uint offset, ";
-                    if (!EmitTypeAndName(fn, value_ty, ast::StorageClass::kNone,
-                                         ast::Access::kUndefined, "compare")) {
-                        return "";
-                    }
-                    fn << ", ";
-                    if (!EmitTypeAndName(fn, value_ty, ast::StorageClass::kNone,
-                                         ast::Access::kUndefined, "value")) {
-                        return "";
-                    }
-                    fn << ") {";
-                }
-
-                buf.IncrementIndent();
-                TINT_DEFER({
-                    buf.DecrementIndent();
-                    line(&buf) << "}";
-                    line(&buf);
-                });
-
-                {  // T result = {0, 0};
-                    auto l = line(&buf);
-                    if (!EmitTypeAndName(l, result_ty, ast::StorageClass::kNone,
-                                         ast::Access::kUndefined, "result")) {
-                        return "";
-                    }
-                    l << " = {0, 0};";
-                }
-                line(&buf) << "buffer.InterlockedCompareExchange(offset, compare, "
-                              "value, result.x);";
-                line(&buf) << "result.y = result.x == compare;";
-                line(&buf) << "return result;";
-                return name;
-            }
-            default:
-                break;
         }
-        TINT_UNREACHABLE(Writer, diagnostics_)
+        default:
-            << "unsupported atomic DecomposeMemoryAccess::Intrinsic::Op: "
+            break;
-            << static_cast<int>(intrinsic->op);
-        return "";
-    };
-
-    auto func = utils::GetOrCreate(dma_intrinsics_, DMAIntrinsic{intrinsic->op, intrinsic->type},
-                                   generate_helper);
-    if (func.empty()) {
-        return false;
     }
 
-    out << func;
+    TINT_UNREACHABLE(Writer, diagnostics_)
-    {
+        << "unsupported atomic DecomposeMemoryAccess::Intrinsic::Op: "
-        ScopedParen sp(out);
+        << static_cast<int>(intrinsic->op);
-        bool first = true;
+    return false;
-        for (auto* arg : expr->args) {
-            if (!first) {
-                out << ", ";
-            }
-            first = false;
-            if (!EmitExpression(out, arg)) {
-                return false;
-            }
-        }
-    }
-
-    return true;
 }
 
 bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
@@ -1788,6 +1759,12 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
             return 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* dest = expr->args[0];
             auto* compare_value = expr->args[1];
             auto* value = expr->args[2];
@@ -1807,7 +1784,7 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
                 pre << ";";
             }
 
-            {  // InterlockedCompareExchange(dst, compare, value, result.x);
+            {  // InterlockedCompareExchange(dst, compare, value, result.old_value);
                 auto pre = line();
                 pre << "InterlockedCompareExchange";
                 {
@@ -1819,14 +1796,13 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
                     if (!EmitExpression(pre, value)) {
                         return false;
                     }
-                    pre << ", " << result << ".x";
+                    pre << ", " << result << ".old_value";
                 }
                 pre << ";";
             }
 
-            {  // result.y = result.x == compare;
+            // result.exchanged = result.old_value == compare;
-                line() << result << ".y = " << result << ".x == " << compare << ";";
+            line() << result << ".exchanged = " << result << ".old_value == " << compare << ";";
-            }
 
             out << result;
             return true;
@@ -2740,6 +2716,17 @@ bool GeneratorImpl::EmitIf(const ast::IfStatement* stmt) {
 bool GeneratorImpl::EmitFunction(const ast::Function* func) {
     auto* sem = builder_.Sem().Get(func);
 
+    // Emit storage atomic helpers
+    if (auto* intrinsic =
+            ast::GetAttribute<transform::DecomposeMemoryAccess::Intrinsic>(func->attributes)) {
+        if (intrinsic->storage_class == ast::StorageClass::kStorage && intrinsic->IsAtomic()) {
+            if (!EmitStorageAtomicIntrinsic(func, intrinsic)) {
+                return false;
+            }
+        }
+        return true;
+    }
+
     if (ast::HasAttribute<ast::InternalAttribute>(func->attributes)) {
         // An internal function. Do not emit.
         return true;
@@ -3755,13 +3742,9 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
         ScopedIndent si(b);
         for (auto* mem : str->Members()) {
             auto mem_name = builder_.Symbols().NameFor(mem->Name());
-
             auto* ty = mem->Type();
-
             auto out = line(b);
-
             std::string pre, post;
-
             if (auto* decl = mem->Declaration()) {
                 for (auto* attr : decl->attributes) {
                     if (auto* location = attr->As<ast::LocationAttribute>()) {
@@ -3826,7 +3809,6 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
     }
 
     line(b) << "};";
-
     return true;
 }
 

src/tint/writer/hlsl/generator_impl.h

@@ -187,6 +187,12 @@ class GeneratorImpl : public TextGenerator {
     bool EmitStorageAtomicCall(std::ostream& out,
                                const ast::CallExpression* expr,
                                const transform::DecomposeMemoryAccess::Intrinsic* intrinsic);
+    /// Handles generating the helper function for the atomic intrinsic function
+    /// @param func the function
+    /// @param intrinsic the atomic intrinsic
+    /// @returns true if the function is emitted
+    bool EmitStorageAtomicIntrinsic(const ast::Function* func,
+                                    const transform::DecomposeMemoryAccess::Intrinsic* intrinsic);
     /// Handles generating an atomic intrinsic call for a workgroup variable
     /// @param out the output of the expression stream
     /// @param expr the call expression
@@ -511,7 +517,6 @@ class GeneratorImpl : public TextGenerator {
 
     TextBuffer helpers_;  // Helper functions emitted at the top of the output
     std::function<bool()> emit_continuing_;
-    std::unordered_map<DMAIntrinsic, std::string, DMAIntrinsic::Hasher> dma_intrinsics_;
     std::unordered_map<const sem::Matrix*, std::string> matrix_scalar_ctors_;
     std::unordered_map<const sem::Builtin*, std::string> builtins_;
     std::unordered_map<const sem::Struct*, std::string> structure_builders_;

src/tint/writer/msl/generator_impl.cc

@@ -806,6 +806,12 @@ 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();
 
@@ -816,7 +822,8 @@ bool GeneratorImpl::EmitAtomicCall(std::ostream& out,
                 line(&buf) << "template <typename A, typename T>";
                 {
                     auto f = line(&buf);
-                    f << "vec<T, 2> " << name << "(";
+                    auto str_name = StructName(builtin->ReturnType()->As<sem::Struct>());
+                    f << str_name << " " << name << "(";
                     if (!EmitStorageClass(f, sc)) {
                         return "";
                     }
@@ -830,12 +837,12 @@ bool GeneratorImpl::EmitAtomicCall(std::ostream& out,
                     line(&buf);
                 });
 
-                line(&buf) << "T prev_value = compare;";
+                line(&buf) << "T old_value = compare;";
-                line(&buf) << "bool matched = "
+                line(&buf) << "bool exchanged = "
                               "atomic_compare_exchange_weak_explicit(atomic, "
-                              "&prev_value, value, memory_order_relaxed, "
+                              "&old_value, value, memory_order_relaxed, "
                               "memory_order_relaxed);";
-                line(&buf) << "return {prev_value, matched};";
+                line(&buf) << "return {old_value, exchanged};";
                 return name;
             });
 

src/tint/writer/spirv/builder.cc

@@ -3201,42 +3201,12 @@ bool Builder::GenerateAtomicBuiltin(const sem::Call* call,
                 return false;
             }
 
-            // zero := T(0)
+            // result := __atomic_compare_exchange_result<T>(original_value, values_equal)
-            // one := T(1)
-            uint32_t zero = 0;
-            uint32_t one = 0;
-            if (value_sem_type->Is<sem::I32>()) {
-                zero = GenerateConstantIfNeeded(ScalarConstant::I32(0u));
-                one = GenerateConstantIfNeeded(ScalarConstant::I32(1u));
-            } else if (value_sem_type->Is<sem::U32>()) {
-                zero = GenerateConstantIfNeeded(ScalarConstant::U32(0u));
-                one = GenerateConstantIfNeeded(ScalarConstant::U32(1u));
-            } else {
-                TINT_UNREACHABLE(Writer, builder_.Diagnostics())
-                    << "unsupported atomic type " << value_sem_type->TypeInfo().name;
-            }
-            if (zero == 0 || one == 0) {
-                return false;
-            }
-
-            // xchg_success := values_equal ? one : zero
-            auto xchg_success = result_op();
-            if (!push_function_inst(spv::Op::OpSelect, {
-                                                           Operand(value_type),
-                                                           xchg_success,
-                                                           values_equal,
-                                                           Operand(one),
-                                                           Operand(zero),
-                                                       })) {
-                return false;
-            }
-
-            // result := vec2<T>(original_value, xchg_success)
             return push_function_inst(spv::Op::OpCompositeConstruct, {
                                                                          result_type,
                                                                          result_id,
                                                                          original_value,
-                                                                         xchg_success,
+                                                                         values_equal,
                                                                      });
         }
         default:

src/tint/writer/spirv/builder_builtin_test.cc

@@ -2018,15 +2018,15 @@ OpReturn
 
 TEST_F(BuiltinBuilderTest, Call_AtomicCompareExchangeWeak) {
     // struct S {
-    //   u : atomic<u32>;
+    //   u : atomic<u32>,
-    //   i : atomic<i32>;
+    //   i : atomic<i32>,
     // }
     //
     // @binding(1) @group(2) var<storage, read_write> b : S;
     //
     // fn a_func() {
-    //   let u : vec2<u32> = atomicCompareExchangeWeak(&b.u, 10u);
+    //   let u = atomicCompareExchangeWeak(&b.u, 10u, 20u);
-    //   let i : vec2<i32> = atomicCompareExchangeWeak(&b.i, 10);
+    //   let i = atomicCompareExchangeWeak(&b.i, 10, 10);
     // }
     auto* s = Structure("S", {
                                  Member("u", ty.atomic<u32>()),
@@ -2040,10 +2040,10 @@ TEST_F(BuiltinBuilderTest, Call_AtomicCompareExchangeWeak) {
 
     Func("a_func", {}, ty.void_(),
          ast::StatementList{
-             Decl(Let("u", ty.vec2<u32>(),
+             Decl(Let("u", nullptr,
                       Call("atomicCompareExchangeWeak", AddressOf(MemberAccessor("b", "u")), 10_u,
                            20_u))),
-             Decl(Let("i", ty.vec2<i32>(),
+             Decl(Let("i", nullptr,
                       Call("atomicCompareExchangeWeak", AddressOf(MemberAccessor("b", "i")), 10_i,
                            20_i))),
          },
@@ -2062,33 +2062,29 @@ TEST_F(BuiltinBuilderTest, Call_AtomicCompareExchangeWeak) {
 %1 = OpVariable %2 StorageBuffer
 %7 = OpTypeVoid
 %6 = OpTypeFunction %7
-%11 = OpTypeVector %4 2
+%12 = OpTypeBool
-%12 = OpConstant %4 1
+%11 = OpTypeStruct %4 %12
-%13 = OpConstant %4 0
+%13 = OpConstant %4 1
-%15 = OpTypePointer StorageBuffer %4
+%14 = OpConstant %4 0
-%17 = OpConstant %4 20
+%16 = OpTypePointer StorageBuffer %4
-%18 = OpConstant %4 10
+%18 = OpConstant %4 20
-%19 = OpTypeBool
+%19 = OpConstant %4 10
-%24 = OpTypeVector %5 2
+%23 = OpTypeStruct %5 %12
-%26 = OpTypePointer StorageBuffer %5
+%25 = OpTypePointer StorageBuffer %5
-%28 = OpConstant %5 20
+%27 = OpConstant %5 20
-%29 = OpConstant %5 10
+%28 = OpConstant %5 10
-%32 = OpConstant %5 0
-%33 = OpConstant %5 1
 )";
     auto got_types = DumpInstructions(b.types());
     EXPECT_EQ(expected_types, got_types);
 
-    auto* expected_instructions = R"(%16 = OpAccessChain %15 %1 %13
+    auto* expected_instructions = R"(%17 = OpAccessChain %16 %1 %14
-%20 = OpAtomicCompareExchange %4 %16 %12 %13 %13 %17 %18
+%20 = OpAtomicCompareExchange %4 %17 %13 %14 %14 %18 %19
-%21 = OpIEqual %19 %20 %17
+%21 = OpIEqual %12 %20 %18
-%22 = OpSelect %4 %21 %12 %13
+%10 = OpCompositeConstruct %11 %20 %21
-%10 = OpCompositeConstruct %11 %20 %22
+%26 = OpAccessChain %25 %1 %13
-%27 = OpAccessChain %26 %1 %12
+%29 = OpAtomicCompareExchange %5 %26 %13 %14 %14 %27 %28
-%30 = OpAtomicCompareExchange %5 %27 %12 %13 %13 %28 %29
+%30 = OpIEqual %12 %29 %27
-%31 = OpIEqual %19 %30 %28
+%22 = OpCompositeConstruct %23 %29 %30
-%34 = OpSelect %5 %31 %33 %32
-%23 = OpCompositeConstruct %24 %30 %34
 OpReturn
 )";
     auto got_instructions = DumpInstructions(b.functions()[0].instructions());

test/tint/bug/chromium/1273230.wgsl.expected.hlsl

@@ -2,24 +2,6 @@ uint value_or_one_if_zero_uint(uint value) {
   return value == 0u ? 1u : value;
 }
 
-uint atomicLoad_1(RWByteAddressBuffer buffer, uint offset) {
-  uint value = 0;
-  buffer.InterlockedOr(offset, 0, value);
-  return value;
-}
-
-int atomicLoad_2(RWByteAddressBuffer buffer, uint offset) {
-  int value = 0;
-  buffer.InterlockedOr(offset, 0, value);
-  return value;
-}
-
-int atomicAdd_1(RWByteAddressBuffer buffer, uint offset, int value) {
-  int original_value = 0;
-  buffer.InterlockedAdd(offset, value, original_value);
-  return original_value;
-}
-
 void marg8uintin() {
 }
 
@@ -61,19 +43,40 @@ float3 loadPosition(uint vertexIndex) {
   return position;
 }
 
+uint tint_atomicLoad(RWByteAddressBuffer buffer, uint offset) {
+  uint value = 0;
+  buffer.InterlockedOr(offset, 0, value);
+  return value;
+}
+
+
+int tint_atomicLoad_1(RWByteAddressBuffer buffer, uint offset) {
+  int value = 0;
+  buffer.InterlockedOr(offset, 0, value);
+  return value;
+}
+
+
 void doIgnore() {
   uint g43 = uniforms[0].x;
   uint kj6 = dbg.Load(20u);
-  uint b53 = atomicLoad_1(counters, (4u * uint(0)));
+  uint b53 = tint_atomicLoad(counters, (4u * uint(0)));
   uint rwg = indices.Load((4u * uint(0)));
   float rb5 = asfloat(positions.Load((4u * uint(0))));
-  int g55 = atomicLoad_2(LUT, (4u * uint(0)));
+  int g55 = tint_atomicLoad_1(LUT, (4u * uint(0)));
 }
 
 struct tint_symbol_1 {
   uint3 GlobalInvocationID : SV_DispatchThreadID;
 };
 
+int tint_atomicAdd(RWByteAddressBuffer buffer, uint offset, int value) {
+  int original_value = 0;
+  buffer.InterlockedAdd(offset, value, original_value);
+  return original_value;
+}
+
+
 void main_count_inner(uint3 GlobalInvocationID) {
   uint triangleIndex = GlobalInvocationID.x;
   if ((triangleIndex >= uniforms[0].x)) {
@@ -89,7 +92,7 @@ void main_count_inner(uint3 GlobalInvocationID) {
   float3 center = (((p0 + p2) + p1) / 3.0f);
   float3 voxelPos = toVoxelPos(p1);
   uint lIndex = toIndex1D(uniforms[0].y, p0);
-  int triangleOffset = atomicAdd_1(LUT, (4u * i1), 1);
+  int triangleOffset = tint_atomicAdd(LUT, (4u * i1), 1);
 }
 
 [numthreads(128, 1, 1)]

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

@@ -2,35 +2,6 @@ uint value_or_one_if_zero_uint(uint value) {
   return value == 0u ? 1u : value;
 }
 
-uint atomicLoad_1(RWByteAddressBuffer buffer, uint offset) {
-  uint value = 0;
-  buffer.InterlockedOr(offset, 0, value);
-  return value;
-}
-
-int atomicLoad_2(RWByteAddressBuffer buffer, uint offset) {
-  int value = 0;
-  buffer.InterlockedOr(offset, 0, value);
-  return value;
-}
-
-uint atomicAdd_1(RWByteAddressBuffer buffer, uint offset, uint value) {
-  uint original_value = 0;
-  buffer.InterlockedAdd(offset, value, original_value);
-  return original_value;
-}
-
-void atomicStore_1(RWByteAddressBuffer buffer, uint offset, int value) {
-  int ignored;
-  buffer.InterlockedExchange(offset, value, ignored);
-}
-
-int atomicAdd_2(RWByteAddressBuffer buffer, uint offset, int value) {
-  int original_value = 0;
-  buffer.InterlockedAdd(offset, value, original_value);
-  return original_value;
-}
-
 cbuffer cbuffer_uniforms : register(b0, space0) {
   uint4 uniforms[3];
 };
@@ -69,19 +40,40 @@ float3 loadPosition(uint vertexIndex) {
   return position;
 }
 
+uint tint_atomicLoad(RWByteAddressBuffer buffer, uint offset) {
+  uint value = 0;
+  buffer.InterlockedOr(offset, 0, value);
+  return value;
+}
+
+
+int tint_atomicLoad_1(RWByteAddressBuffer buffer, uint offset) {
+  int value = 0;
+  buffer.InterlockedOr(offset, 0, value);
+  return value;
+}
+
+
 void doIgnore() {
   uint g42 = uniforms[0].x;
   uint kj6 = dbg.Load(20u);
-  uint b53 = atomicLoad_1(counters, (4u * uint(0)));
+  uint b53 = tint_atomicLoad(counters, (4u * uint(0)));
   uint rwg = indices.Load((4u * uint(0)));
   float rb5 = asfloat(positions.Load((4u * uint(0))));
-  int g55 = atomicLoad_2(LUT, (4u * uint(0)));
+  int g55 = tint_atomicLoad_1(LUT, (4u * uint(0)));
 }
 
 struct tint_symbol_1 {
   uint3 GlobalInvocationID : SV_DispatchThreadID;
 };
 
+uint tint_atomicAdd(RWByteAddressBuffer buffer, uint offset, uint value) {
+  uint original_value = 0;
+  buffer.InterlockedAdd(offset, value, original_value);
+  return original_value;
+}
+
+
 void main_count_inner(uint3 GlobalInvocationID) {
   uint triangleIndex = GlobalInvocationID.x;
   if ((triangleIndex >= uniforms[0].x)) {
@@ -97,7 +89,7 @@ void main_count_inner(uint3 GlobalInvocationID) {
   float3 center = (((p0 + p1) + p2) / 3.0f);
   float3 voxelPos = toVoxelPos(center);
   uint voxelIndex = toIndex1D(uniforms[0].y, voxelPos);
-  uint acefg = atomicAdd_1(counters, (4u * voxelIndex), 1u);
+  uint acefg = tint_atomicAdd(counters, (4u * voxelIndex), 1u);
   if ((triangleIndex == 0u)) {
     dbg.Store(16u, asuint(uniforms[0].y));
     dbg.Store(32u, asuint(center.x));
@@ -116,6 +108,19 @@ struct tint_symbol_3 {
   uint3 GlobalInvocationID : SV_DispatchThreadID;
 };
 
+uint tint_atomicAdd_1(RWByteAddressBuffer buffer, uint offset, uint value) {
+  uint original_value = 0;
+  buffer.InterlockedAdd(offset, value, original_value);
+  return original_value;
+}
+
+
+void tint_atomicStore(RWByteAddressBuffer buffer, uint offset, int value) {
+  int ignored;
+  buffer.InterlockedExchange(offset, value, ignored);
+}
+
+
 void main_create_lut_inner(uint3 GlobalInvocationID) {
   uint voxelIndex = GlobalInvocationID.x;
   doIgnore();
@@ -123,13 +128,13 @@ void main_create_lut_inner(uint3 GlobalInvocationID) {
   if ((voxelIndex >= maxVoxels)) {
     return;
   }
-  uint numTriangles = atomicLoad_1(counters, (4u * voxelIndex));
+  uint numTriangles = tint_atomicLoad(counters, (4u * voxelIndex));
   int offset = -1;
   if ((numTriangles > 0u)) {
-    const uint tint_symbol_6 = atomicAdd_1(dbg, 0u, numTriangles);
+    const uint tint_symbol_6 = tint_atomicAdd_1(dbg, 0u, numTriangles);
     offset = int(tint_symbol_6);
   }
-  atomicStore_1(LUT, (4u * voxelIndex), offset);
+  tint_atomicStore(LUT, (4u * voxelIndex), offset);
 }
 
 [numthreads(128, 1, 1)]
@@ -142,6 +147,13 @@ struct tint_symbol_5 {
   uint3 GlobalInvocationID : SV_DispatchThreadID;
 };
 
+int tint_atomicAdd_2(RWByteAddressBuffer buffer, uint offset, int value) {
+  int original_value = 0;
+  buffer.InterlockedAdd(offset, value, original_value);
+  return original_value;
+}
+
+
 void main_sort_triangles_inner(uint3 GlobalInvocationID) {
   uint triangleIndex = GlobalInvocationID.x;
   doIgnore();
@@ -157,7 +169,7 @@ void main_sort_triangles_inner(uint3 GlobalInvocationID) {
   float3 center = (((p0 + p1) + p2) / 3.0f);
   float3 voxelPos = toVoxelPos(center);
   uint voxelIndex = toIndex1D(uniforms[0].y, voxelPos);
-  int triangleOffset = atomicAdd_2(LUT, (4u * voxelIndex), 1);
+  int triangleOffset = tint_atomicAdd_2(LUT, (4u * voxelIndex), 1);
 }
 
 [numthreads(128, 1, 1)]

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

@@ -1,9 +1,3 @@
-uint atomicAdd_1(RWByteAddressBuffer buffer, uint offset, uint value) {
-  uint original_value = 0;
-  buffer.InterlockedAdd(offset, value, original_value);
-  return original_value;
-}
-
 RWByteAddressBuffer lightsBuffer : register(u0, space0);
 
 RWByteAddressBuffer tileLightId : register(u0, space1);
@@ -28,6 +22,13 @@ float4x4 tint_symbol_6(uint4 buffer[11], uint offset) {
   return float4x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4]), asfloat(buffer[scalar_offset_3 / 4]));
 }
 
+uint tint_atomicAdd(RWByteAddressBuffer buffer, uint offset, uint value) {
+  uint original_value = 0;
+  buffer.InterlockedAdd(offset, value, original_value);
+  return original_value;
+}
+
+
 void main_inner(uint3 GlobalInvocationID) {
   uint index = GlobalInvocationID.x;
   if ((index >= config[0].x)) {
@@ -96,7 +97,7 @@ void main_inner(uint3 GlobalInvocationID) {
             if ((tint_tmp)) {
               continue;
             }
-            uint offset = atomicAdd_1(tileLightId, (260u * tileId), 1u);
+            uint offset = tint_atomicAdd(tileLightId, (260u * tileId), 1u);
             if ((offset >= config[1].x)) {
               continue;
             }

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

@@ -1,9 +1,3 @@
-uint atomicAdd_1(RWByteAddressBuffer buffer, uint offset, uint value) {
-  uint original_value = 0;
-  buffer.InterlockedAdd(offset, value, original_value);
-  return original_value;
-}
-
 RWByteAddressBuffer drawOut : register(u5, space0);
 static uint cubeVerts = 0u;
 
@@ -11,8 +5,15 @@ struct tint_symbol_1 {
   uint3 global_id : SV_DispatchThreadID;
 };
 
+uint tint_atomicAdd(RWByteAddressBuffer buffer, uint offset, uint value) {
+  uint original_value = 0;
+  buffer.InterlockedAdd(offset, value, original_value);
+  return original_value;
+}
+
+
 void computeMain_inner(uint3 global_id) {
-  const uint firstVertex = atomicAdd_1(drawOut, 0u, cubeVerts);
+  const uint firstVertex = tint_atomicAdd(drawOut, 0u, cubeVerts);
 }
 
 [numthreads(1, 1, 1)]

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

@@ -1,9 +1,3 @@
-int atomicLoad_1(RWByteAddressBuffer buffer, uint offset) {
-  int value = 0;
-  buffer.InterlockedOr(offset, 0, value);
-  return value;
-}
-
 cbuffer cbuffer_constants : register(b0, space1) {
   uint4 constants[1];
 };
@@ -12,8 +6,15 @@ RWByteAddressBuffer result : register(u1, space1);
 
 RWByteAddressBuffer s : register(u0, space0);
 
+int tint_atomicLoad(RWByteAddressBuffer buffer, uint offset) {
+  int value = 0;
+  buffer.InterlockedOr(offset, 0, value);
+  return value;
+}
+
+
 int runTest() {
-  return atomicLoad_1(s, (4u * (0u + uint(constants[0].x))));
+  return tint_atomicLoad(s, (4u * (0u + uint(constants[0].x))));
 }
 
 [numthreads(1, 1, 1)]

test/tint/builtins/gen/atomicAdd/8a199a.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-uint atomicAdd_1(RWByteAddressBuffer buffer, uint offset, uint value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+uint tint_atomicAdd(RWByteAddressBuffer buffer, uint offset, uint value) {
   uint original_value = 0;
   buffer.InterlockedAdd(offset, value, original_value);
   return original_value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicAdd_8a199a() {
-  uint res = atomicAdd_1(sb_rw, 0u, 1u);
+  uint res = tint_atomicAdd(sb_rw, 0u, 1u);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicAdd/d32fe4.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-int atomicAdd_1(RWByteAddressBuffer buffer, uint offset, int value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+int tint_atomicAdd(RWByteAddressBuffer buffer, uint offset, int value) {
   int original_value = 0;
   buffer.InterlockedAdd(offset, value, original_value);
   return original_value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicAdd_d32fe4() {
-  int res = atomicAdd_1(sb_rw, 0u, 1);
+  int res = tint_atomicAdd(sb_rw, 0u, 1);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicAnd/152966.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-int atomicAnd_1(RWByteAddressBuffer buffer, uint offset, int value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+int tint_atomicAnd(RWByteAddressBuffer buffer, uint offset, int value) {
   int original_value = 0;
   buffer.InterlockedAnd(offset, value, original_value);
   return original_value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicAnd_152966() {
-  int res = atomicAnd_1(sb_rw, 0u, 1);
+  int res = tint_atomicAnd(sb_rw, 0u, 1);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicAnd/85a8d9.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-uint atomicAnd_1(RWByteAddressBuffer buffer, uint offset, uint value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+uint tint_atomicAnd(RWByteAddressBuffer buffer, uint offset, uint value) {
   uint original_value = 0;
   buffer.InterlockedAnd(offset, value, original_value);
   return original_value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicAnd_85a8d9() {
-  uint res = atomicAnd_1(sb_rw, 0u, 1u);
+  uint res = tint_atomicAnd(sb_rw, 0u, 1u);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.glsl

@@ -1,76 +0,0 @@
-SKIP: FAILED
-
-#version 310 es
-precision mediump float;
-
-ivec2 tint_atomicCompareExchangeWeak(inout int param_0, int param_1, int param_2) {
-  ivec2 result;
-  result.x = atomicCompSwap(param_0, param_1, param_2);
-  result.y = result.x == param_2 ? 1 : 0;
-  return result;
-}
-
-
-struct SB_RW {
-  int arg_0;
-};
-
-layout(binding = 0, std430) buffer SB_RW_1 {
-  int arg_0;
-} sb_rw;
-void atomicCompareExchangeWeak_12871c() {
-  ivec2 res = tint_atomicCompareExchangeWeak(sb_rw.arg_0, 1, 1);
-}
-
-void fragment_main() {
-  atomicCompareExchangeWeak_12871c();
-}
-
-void main() {
-  fragment_main();
-  return;
-}
-Error parsing GLSL shader:
-ERROR: 0:6: 'atomicCompSwap' : Atomic memory function can only be used for shader storage block member or shared variable. 
-ERROR: 0:6: '' : compilation terminated 
-ERROR: 2 compilation errors.  No code generated.
-
-
-
-#version 310 es
-
-ivec2 tint_atomicCompareExchangeWeak(inout int param_0, int param_1, int param_2) {
-  ivec2 result;
-  result.x = atomicCompSwap(param_0, param_1, param_2);
-  result.y = result.x == param_2 ? 1 : 0;
-  return result;
-}
-
-
-struct SB_RW {
-  int arg_0;
-};
-
-layout(binding = 0, std430) buffer SB_RW_1 {
-  int arg_0;
-} sb_rw;
-void atomicCompareExchangeWeak_12871c() {
-  ivec2 res = tint_atomicCompareExchangeWeak(sb_rw.arg_0, 1, 1);
-}
-
-void compute_main() {
-  atomicCompareExchangeWeak_12871c();
-}
-
-layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
-void main() {
-  compute_main();
-  return;
-}
-Error parsing GLSL shader:
-ERROR: 0:5: 'atomicCompSwap' : Atomic memory function can only be used for shader storage block member or shared variable. 
-ERROR: 0:5: '' : compilation terminated 
-ERROR: 2 compilation errors.  No code generated.
-
-
-

test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.hlsl

@@ -1,23 +0,0 @@
-int2 atomicCompareExchangeWeak_1(RWByteAddressBuffer buffer, uint offset, int compare, int value) {
-  int2 result = {0, 0};
-  buffer.InterlockedCompareExchange(offset, compare, value, result.x);
-  result.y = result.x == compare;
-  return result;
-}
-
-RWByteAddressBuffer sb_rw : register(u0, space0);
-
-void atomicCompareExchangeWeak_12871c() {
-  int2 res = atomicCompareExchangeWeak_1(sb_rw, 0u, 1, 1);
-}
-
-void fragment_main() {
-  atomicCompareExchangeWeak_12871c();
-  return;
-}
-
-[numthreads(1, 1, 1)]
-void compute_main() {
-  atomicCompareExchangeWeak_12871c();
-  return;
-}

test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl

@@ -1,29 +0,0 @@
-#include <metal_stdlib>
-
-using namespace metal;
-
-template <typename A, typename T>
-vec<T, 2> atomicCompareExchangeWeak_1(device A* atomic, T compare, T value) {
-  T prev_value = compare;
-  bool matched = atomic_compare_exchange_weak_explicit(atomic, &prev_value, value, memory_order_relaxed, memory_order_relaxed);
-  return {prev_value, matched};
-}
-
-struct SB_RW {
-  /* 0x0000 */ atomic_int arg_0;
-};
-
-void atomicCompareExchangeWeak_12871c(device SB_RW* const tint_symbol) {
-  int2 res = atomicCompareExchangeWeak_1(&((*(tint_symbol)).arg_0), 1, 1);
-}
-
-fragment void fragment_main(device SB_RW* tint_symbol_1 [[buffer(0)]]) {
-  atomicCompareExchangeWeak_12871c(tint_symbol_1);
-  return;
-}
-
-kernel void compute_main(device SB_RW* tint_symbol_2 [[buffer(0)]]) {
-  atomicCompareExchangeWeak_12871c(tint_symbol_2);
-  return;
-}
-

test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl

@@ -27,17 +27,17 @@ struct SB_RW {
 };
 @group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
 
-// fn atomicCompareExchangeWeak(ptr<storage, atomic<i32>, read_write>, i32, i32) -> vec2<i32>
+// fn atomicCompareExchangeWeak(ptr<storage, atomic<i32>, read_write>, i32, i32) -> __atomic_compare_exchange_result<i32>
-fn atomicCompareExchangeWeak_12871c() {
+fn atomicCompareExchangeWeak_1bd40a() {
-  var res: vec2<i32> = atomicCompareExchangeWeak(&sb_rw.arg_0, 1, 1);
+  var res = atomicCompareExchangeWeak(&sb_rw.arg_0, 1, 1);
 }
 
 @stage(fragment)
 fn fragment_main() {
-  atomicCompareExchangeWeak_12871c();
+  atomicCompareExchangeWeak_1bd40a();
 }
 
 @stage(compute) @workgroup_size(1)
 fn compute_main() {
-  atomicCompareExchangeWeak_12871c();
+  atomicCompareExchangeWeak_1bd40a();
 }

test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.glsl

@@ -0,0 +1,62 @@
+#version 310 es
+precision mediump float;
+
+struct atomic_compare_exchange_resulti32 {
+  int old_value;
+  bool exchanged;
+};
+
+
+struct SB_RW {
+  int arg_0;
+};
+
+layout(binding = 0, std430) buffer SB_RW_1 {
+  int arg_0;
+} sb_rw;
+void atomicCompareExchangeWeak_1bd40a() {
+  atomic_compare_exchange_resulti32 atomic_compare_result;
+  atomic_compare_result.old_value = atomicCompSwap(sb_rw.arg_0, 1, 1);
+  atomic_compare_result.exchanged = atomic_compare_result.old_value == 1;
+  atomic_compare_exchange_resulti32 res = atomic_compare_result;
+}
+
+void fragment_main() {
+  atomicCompareExchangeWeak_1bd40a();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+struct atomic_compare_exchange_resulti32 {
+  int old_value;
+  bool exchanged;
+};
+
+
+struct SB_RW {
+  int arg_0;
+};
+
+layout(binding = 0, std430) buffer SB_RW_1 {
+  int arg_0;
+} sb_rw;
+void atomicCompareExchangeWeak_1bd40a() {
+  atomic_compare_exchange_resulti32 atomic_compare_result;
+  atomic_compare_result.old_value = atomicCompSwap(sb_rw.arg_0, 1, 1);
+  atomic_compare_result.exchanged = atomic_compare_result.old_value == 1;
+  atomic_compare_exchange_resulti32 res = atomic_compare_result;
+}
+
+void compute_main() {
+  atomicCompareExchangeWeak_1bd40a();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.hlsl

@@ -0,0 +1,29 @@
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+struct atomic_compare_exchange_weak_ret_type {
+  int old_value;
+  bool exchanged;
+};
+
+atomic_compare_exchange_weak_ret_type tint_atomicCompareExchangeWeak(RWByteAddressBuffer buffer, uint offset, int compare, int 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;
+}
+
+
+void atomicCompareExchangeWeak_1bd40a() {
+  atomic_compare_exchange_weak_ret_type res = tint_atomicCompareExchangeWeak(sb_rw, 0u, 1, 1);
+}
+
+void fragment_main() {
+  atomicCompareExchangeWeak_1bd40a();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  atomicCompareExchangeWeak_1bd40a();
+  return;
+}

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

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct atomic_compare_exchange_resulti32 {
+  int old_value;
+  bool exchanged;
+};
+template <typename A, typename T>
+atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(device A* atomic, T compare, T value) {
+  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};
+}
+
+struct SB_RW {
+  /* 0x0000 */ atomic_int arg_0;
+};
+
+void atomicCompareExchangeWeak_1bd40a(device SB_RW* const tint_symbol) {
+  atomic_compare_exchange_resulti32 res = atomicCompareExchangeWeak_1(&((*(tint_symbol)).arg_0), 1, 1);
+}
+
+fragment void fragment_main(device SB_RW* tint_symbol_1 [[buffer(0)]]) {
+  atomicCompareExchangeWeak_1bd40a(tint_symbol_1);
+  return;
+}
+
+kernel void compute_main(device SB_RW* tint_symbol_2 [[buffer(0)]]) {
+  atomicCompareExchangeWeak_1bd40a(tint_symbol_2);
+  return;
+}
+

test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 32
+; Bound: 30
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -12,7 +12,10 @@
                OpName %SB_RW "SB_RW"
                OpMemberName %SB_RW 0 "arg_0"
                OpName %sb_rw "sb_rw"
-               OpName %atomicCompareExchangeWeak_12871c "atomicCompareExchangeWeak_12871c"
+               OpName %atomicCompareExchangeWeak_1bd40a "atomicCompareExchangeWeak_1bd40a"
+               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 %res "res"
                OpName %fragment_main "fragment_main"
                OpName %compute_main "compute_main"
@@ -20,40 +23,40 @@
                OpMemberDecorate %SB_RW 0 Offset 0
                OpDecorate %sb_rw DescriptorSet 0
                OpDecorate %sb_rw Binding 0
+               OpMemberDecorate %__atomic_compare_exchange_resulti32 0 Offset 0
+               OpMemberDecorate %__atomic_compare_exchange_resulti32 1 Offset 4
         %int = OpTypeInt 32 1
       %SB_RW = OpTypeStruct %int
 %_ptr_StorageBuffer_SB_RW = OpTypePointer StorageBuffer %SB_RW
       %sb_rw = OpVariable %_ptr_StorageBuffer_SB_RW StorageBuffer
        %void = OpTypeVoid
           %5 = OpTypeFunction %void
-      %v2int = OpTypeVector %int 2
+       %bool = OpTypeBool
+%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool
        %uint = OpTypeInt 32 0
      %uint_1 = OpConstant %uint 1
      %uint_0 = OpConstant %uint 0
 %_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
       %int_1 = OpConstant %int 1
-       %bool = OpTypeBool
+%_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32
-      %int_0 = OpConstant %int 0
+         %23 = OpConstantNull %__atomic_compare_exchange_resulti32
-%_ptr_Function_v2int = OpTypePointer Function %v2int
+%atomicCompareExchangeWeak_1bd40a = OpFunction %void None %5
-         %25 = OpConstantNull %v2int
-%atomicCompareExchangeWeak_12871c = OpFunction %void None %5
           %8 = OpLabel
-        %res = OpVariable %_ptr_Function_v2int Function %25
+        %res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %23
-         %16 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0
+         %17 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0
-         %19 = OpAtomicCompareExchange %int %16 %uint_1 %uint_0 %uint_0 %int_1 %int_1
+         %19 = OpAtomicCompareExchange %int %17 %uint_1 %uint_0 %uint_0 %int_1 %int_1
          %20 = OpIEqual %bool %19 %int_1
-         %22 = OpSelect %int %20 %int_1 %int_0
+          %9 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %19 %20
-          %9 = OpCompositeConstruct %v2int %19 %22
                OpStore %res %9
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %5
-         %27 = OpLabel
+         %25 = OpLabel
-         %28 = OpFunctionCall %void %atomicCompareExchangeWeak_12871c
+         %26 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %5
-         %30 = OpLabel
+         %28 = OpLabel
-         %31 = OpFunctionCall %void %atomicCompareExchangeWeak_12871c
+         %29 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
                OpReturn
                OpFunctionEnd

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

@@ -4,16 +4,16 @@ struct SB_RW {
 
 @group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
 
-fn atomicCompareExchangeWeak_12871c() {
+fn atomicCompareExchangeWeak_1bd40a() {
-  var res : vec2<i32> = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1, 1);
+  var res = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1, 1);
 }
 
 @stage(fragment)
 fn fragment_main() {
-  atomicCompareExchangeWeak_12871c();
+  atomicCompareExchangeWeak_1bd40a();
 }
 
 @stage(compute) @workgroup_size(1)
 fn compute_main() {
-  atomicCompareExchangeWeak_12871c();
+  atomicCompareExchangeWeak_1bd40a();
 }

test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl

@@ -27,17 +27,17 @@ struct SB_RW {
 };
 @group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
 
-// fn atomicCompareExchangeWeak(ptr<storage, atomic<u32>, read_write>, u32, u32) -> vec2<u32>
+// fn atomicCompareExchangeWeak(ptr<storage, atomic<u32>, read_write>, u32, u32) -> __atomic_compare_exchange_result<u32>
-fn atomicCompareExchangeWeak_6673da() {
+fn atomicCompareExchangeWeak_63d8e6() {
-  var res: vec2<u32> = atomicCompareExchangeWeak(&sb_rw.arg_0, 1u, 1u);
+  var res = atomicCompareExchangeWeak(&sb_rw.arg_0, 1u, 1u);
 }
 
 @stage(fragment)
 fn fragment_main() {
-  atomicCompareExchangeWeak_6673da();
+  atomicCompareExchangeWeak_63d8e6();
 }
 
 @stage(compute) @workgroup_size(1)
 fn compute_main() {
-  atomicCompareExchangeWeak_6673da();
+  atomicCompareExchangeWeak_63d8e6();
 }

test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.glsl

@@ -0,0 +1,62 @@
+#version 310 es
+precision mediump float;
+
+struct atomic_compare_exchange_resultu32 {
+  uint old_value;
+  bool exchanged;
+};
+
+
+struct SB_RW {
+  uint arg_0;
+};
+
+layout(binding = 0, std430) buffer SB_RW_1 {
+  uint arg_0;
+} sb_rw;
+void atomicCompareExchangeWeak_63d8e6() {
+  atomic_compare_exchange_resultu32 atomic_compare_result;
+  atomic_compare_result.old_value = atomicCompSwap(sb_rw.arg_0, 1u, 1u);
+  atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u;
+  atomic_compare_exchange_resultu32 res = atomic_compare_result;
+}
+
+void fragment_main() {
+  atomicCompareExchangeWeak_63d8e6();
+}
+
+void main() {
+  fragment_main();
+  return;
+}
+#version 310 es
+
+struct atomic_compare_exchange_resultu32 {
+  uint old_value;
+  bool exchanged;
+};
+
+
+struct SB_RW {
+  uint arg_0;
+};
+
+layout(binding = 0, std430) buffer SB_RW_1 {
+  uint arg_0;
+} sb_rw;
+void atomicCompareExchangeWeak_63d8e6() {
+  atomic_compare_exchange_resultu32 atomic_compare_result;
+  atomic_compare_result.old_value = atomicCompSwap(sb_rw.arg_0, 1u, 1u);
+  atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u;
+  atomic_compare_exchange_resultu32 res = atomic_compare_result;
+}
+
+void compute_main() {
+  atomicCompareExchangeWeak_63d8e6();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main();
+  return;
+}

test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.hlsl

@@ -0,0 +1,29 @@
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+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;
+}
+
+
+void atomicCompareExchangeWeak_63d8e6() {
+  atomic_compare_exchange_weak_ret_type res = tint_atomicCompareExchangeWeak(sb_rw, 0u, 1u, 1u);
+}
+
+void fragment_main() {
+  atomicCompareExchangeWeak_63d8e6();
+  return;
+}
+
+[numthreads(1, 1, 1)]
+void compute_main() {
+  atomicCompareExchangeWeak_63d8e6();
+  return;
+}

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

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct atomic_compare_exchange_resultu32 {
+  uint old_value;
+  bool exchanged;
+};
+template <typename A, typename T>
+atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(device A* atomic, T compare, T value) {
+  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};
+}
+
+struct SB_RW {
+  /* 0x0000 */ atomic_uint arg_0;
+};
+
+void atomicCompareExchangeWeak_63d8e6(device SB_RW* const tint_symbol) {
+  atomic_compare_exchange_resultu32 res = atomicCompareExchangeWeak_1(&((*(tint_symbol)).arg_0), 1u, 1u);
+}
+
+fragment void fragment_main(device SB_RW* tint_symbol_1 [[buffer(0)]]) {
+  atomicCompareExchangeWeak_63d8e6(tint_symbol_1);
+  return;
+}
+
+kernel void compute_main(device SB_RW* tint_symbol_2 [[buffer(0)]]) {
+  atomicCompareExchangeWeak_63d8e6(tint_symbol_2);
+  return;
+}
+

test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 29
+; Bound: 28
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -12,7 +12,10 @@
                OpName %SB_RW "SB_RW"
                OpMemberName %SB_RW 0 "arg_0"
                OpName %sb_rw "sb_rw"
-               OpName %atomicCompareExchangeWeak_6673da "atomicCompareExchangeWeak_6673da"
+               OpName %atomicCompareExchangeWeak_63d8e6 "atomicCompareExchangeWeak_63d8e6"
+               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 %res "res"
                OpName %fragment_main "fragment_main"
                OpName %compute_main "compute_main"
@@ -20,37 +23,38 @@
                OpMemberDecorate %SB_RW 0 Offset 0
                OpDecorate %sb_rw DescriptorSet 0
                OpDecorate %sb_rw Binding 0
+               OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0
+               OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4
        %uint = OpTypeInt 32 0
       %SB_RW = OpTypeStruct %uint
 %_ptr_StorageBuffer_SB_RW = OpTypePointer StorageBuffer %SB_RW
       %sb_rw = OpVariable %_ptr_StorageBuffer_SB_RW StorageBuffer
        %void = OpTypeVoid
           %5 = OpTypeFunction %void
-     %v2uint = OpTypeVector %uint 2
+       %bool = OpTypeBool
+%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool
      %uint_1 = OpConstant %uint 1
      %uint_0 = OpConstant %uint 0
 %_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
-       %bool = OpTypeBool
+%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32
-%_ptr_Function_v2uint = OpTypePointer Function %v2uint
+         %21 = OpConstantNull %__atomic_compare_exchange_resultu32
-         %22 = OpConstantNull %v2uint
+%atomicCompareExchangeWeak_63d8e6 = OpFunction %void None %5
-%atomicCompareExchangeWeak_6673da = OpFunction %void None %5
           %8 = OpLabel
-        %res = OpVariable %_ptr_Function_v2uint Function %22
+        %res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %21
-         %15 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0
+         %16 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0
-         %17 = OpAtomicCompareExchange %uint %15 %uint_1 %uint_0 %uint_0 %uint_1 %uint_1
+         %17 = OpAtomicCompareExchange %uint %16 %uint_1 %uint_0 %uint_0 %uint_1 %uint_1
          %18 = OpIEqual %bool %17 %uint_1
-         %19 = OpSelect %uint %18 %uint_1 %uint_0
+          %9 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %17 %18
-          %9 = OpCompositeConstruct %v2uint %17 %19
                OpStore %res %9
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %5
-         %24 = OpLabel
+         %23 = OpLabel
-         %25 = OpFunctionCall %void %atomicCompareExchangeWeak_6673da
+         %24 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %5
-         %27 = OpLabel
+         %26 = OpLabel
-         %28 = OpFunctionCall %void %atomicCompareExchangeWeak_6673da
+         %27 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
                OpReturn
                OpFunctionEnd

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

@@ -4,16 +4,16 @@ struct SB_RW {
 
 @group(0) @binding(0) var<storage, read_write> sb_rw : SB_RW;
 
-fn atomicCompareExchangeWeak_6673da() {
+fn atomicCompareExchangeWeak_63d8e6() {
-  var res : vec2<u32> = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1u, 1u);
+  var res = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1u, 1u);
 }
 
 @stage(fragment)
 fn fragment_main() {
-  atomicCompareExchangeWeak_6673da();
+  atomicCompareExchangeWeak_63d8e6();
 }
 
 @stage(compute) @workgroup_size(1)
 fn compute_main() {
-  atomicCompareExchangeWeak_6673da();
+  atomicCompareExchangeWeak_63d8e6();
 }

test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.glsl

@@ -1,76 +0,0 @@
-SKIP: FAILED
-
-#version 310 es
-precision mediump float;
-
-uvec2 tint_atomicCompareExchangeWeak(inout uint param_0, uint param_1, uint param_2) {
-  uvec2 result;
-  result.x = atomicCompSwap(param_0, param_1, param_2);
-  result.y = result.x == param_2 ? 1u : 0u;
-  return result;
-}
-
-
-struct SB_RW {
-  uint arg_0;
-};
-
-layout(binding = 0, std430) buffer SB_RW_1 {
-  uint arg_0;
-} sb_rw;
-void atomicCompareExchangeWeak_6673da() {
-  uvec2 res = tint_atomicCompareExchangeWeak(sb_rw.arg_0, 1u, 1u);
-}
-
-void fragment_main() {
-  atomicCompareExchangeWeak_6673da();
-}
-
-void main() {
-  fragment_main();
-  return;
-}
-Error parsing GLSL shader:
-ERROR: 0:6: 'atomicCompSwap' : Atomic memory function can only be used for shader storage block member or shared variable. 
-ERROR: 0:6: '' : compilation terminated 
-ERROR: 2 compilation errors.  No code generated.
-
-
-
-#version 310 es
-
-uvec2 tint_atomicCompareExchangeWeak(inout uint param_0, uint param_1, uint param_2) {
-  uvec2 result;
-  result.x = atomicCompSwap(param_0, param_1, param_2);
-  result.y = result.x == param_2 ? 1u : 0u;
-  return result;
-}
-
-
-struct SB_RW {
-  uint arg_0;
-};
-
-layout(binding = 0, std430) buffer SB_RW_1 {
-  uint arg_0;
-} sb_rw;
-void atomicCompareExchangeWeak_6673da() {
-  uvec2 res = tint_atomicCompareExchangeWeak(sb_rw.arg_0, 1u, 1u);
-}
-
-void compute_main() {
-  atomicCompareExchangeWeak_6673da();
-}
-
-layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
-void main() {
-  compute_main();
-  return;
-}
-Error parsing GLSL shader:
-ERROR: 0:5: 'atomicCompSwap' : Atomic memory function can only be used for shader storage block member or shared variable. 
-ERROR: 0:5: '' : compilation terminated 
-ERROR: 2 compilation errors.  No code generated.
-
-
-

test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.hlsl

@@ -1,23 +0,0 @@
-uint2 atomicCompareExchangeWeak_1(RWByteAddressBuffer buffer, uint offset, uint compare, uint value) {
-  uint2 result = {0, 0};
-  buffer.InterlockedCompareExchange(offset, compare, value, result.x);
-  result.y = result.x == compare;
-  return result;
-}
-
-RWByteAddressBuffer sb_rw : register(u0, space0);
-
-void atomicCompareExchangeWeak_6673da() {
-  uint2 res = atomicCompareExchangeWeak_1(sb_rw, 0u, 1u, 1u);
-}
-
-void fragment_main() {
-  atomicCompareExchangeWeak_6673da();
-  return;
-}
-
-[numthreads(1, 1, 1)]
-void compute_main() {
-  atomicCompareExchangeWeak_6673da();
-  return;
-}

test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl

@@ -1,29 +0,0 @@
-#include <metal_stdlib>
-
-using namespace metal;
-
-template <typename A, typename T>
-vec<T, 2> atomicCompareExchangeWeak_1(device A* atomic, T compare, T value) {
-  T prev_value = compare;
-  bool matched = atomic_compare_exchange_weak_explicit(atomic, &prev_value, value, memory_order_relaxed, memory_order_relaxed);
-  return {prev_value, matched};
-}
-
-struct SB_RW {
-  /* 0x0000 */ atomic_uint arg_0;
-};
-
-void atomicCompareExchangeWeak_6673da(device SB_RW* const tint_symbol) {
-  uint2 res = atomicCompareExchangeWeak_1(&((*(tint_symbol)).arg_0), 1u, 1u);
-}
-
-fragment void fragment_main(device SB_RW* tint_symbol_1 [[buffer(0)]]) {
-  atomicCompareExchangeWeak_6673da(tint_symbol_1);
-  return;
-}
-
-kernel void compute_main(device SB_RW* tint_symbol_2 [[buffer(0)]]) {
-  atomicCompareExchangeWeak_6673da(tint_symbol_2);
-  return;
-}
-

test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl

@@ -24,12 +24,12 @@
 
 var<workgroup> arg_0: atomic<u32>;
 
-// fn atomicCompareExchangeWeak(ptr<workgroup, atomic<u32>, read_write>, u32, u32) -> vec2<u32>
+// fn atomicCompareExchangeWeak(ptr<workgroup, atomic<u32>, read_write>, u32, u32) -> __atomic_compare_exchange_result<u32>
-fn atomicCompareExchangeWeak_b2ab2c() {
+fn atomicCompareExchangeWeak_83580d() {
-  var res: vec2<u32> = atomicCompareExchangeWeak(&arg_0, 1u, 1u);
+  var res = atomicCompareExchangeWeak(&arg_0, 1u, 1u);
 }
 
 @stage(compute) @workgroup_size(1)
 fn compute_main() {
-  atomicCompareExchangeWeak_b2ab2c();
+  atomicCompareExchangeWeak_83580d();
 }

test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.glsl

@@ -0,0 +1,29 @@
+#version 310 es
+
+struct atomic_compare_exchange_resultu32 {
+  uint old_value;
+  bool exchanged;
+};
+
+
+shared uint arg_0;
+void atomicCompareExchangeWeak_83580d() {
+  atomic_compare_exchange_resultu32 atomic_compare_result;
+  atomic_compare_result.old_value = atomicCompSwap(arg_0, 1u, 1u);
+  atomic_compare_result.exchanged = atomic_compare_result.old_value == 1u;
+  atomic_compare_exchange_resultu32 res = atomic_compare_result;
+}
+
+void compute_main(uint local_invocation_index) {
+  {
+    atomicExchange(arg_0, 0u);
+  }
+  barrier();
+  atomicCompareExchangeWeak_83580d();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}

test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.hlsl

@@ -1,11 +1,15 @@
+struct atomic_compare_exchange_resultu32 {
+  uint old_value;
+  bool exchanged;
+};
 groupshared uint arg_0;
 
-void atomicCompareExchangeWeak_b2ab2c() {
+void atomicCompareExchangeWeak_83580d() {
-  uint2 atomic_result = uint2(0u, 0u);
+  atomic_compare_exchange_resultu32 atomic_result = (atomic_compare_exchange_resultu32)0;
   uint atomic_compare_value = 1u;
-  InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.x);
+  InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.old_value);
-  atomic_result.y = atomic_result.x == atomic_compare_value;
+  atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
-  uint2 res = atomic_result;
+  atomic_compare_exchange_resultu32 res = atomic_result;
 }
 
 struct tint_symbol_1 {
@@ -18,7 +22,7 @@ void compute_main_inner(uint local_invocation_index) {
     InterlockedExchange(arg_0, 0u, atomic_result_1);
   }
   GroupMemoryBarrierWithGroupSync();
-  atomicCompareExchangeWeak_b2ab2c();
+  atomicCompareExchangeWeak_83580d();
 }
 
 [numthreads(1, 1, 1)]

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

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct atomic_compare_exchange_resultu32 {
+  uint old_value;
+  bool exchanged;
+};
+template <typename A, typename T>
+atomic_compare_exchange_resultu32 atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) {
+  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};
+}
+
+void atomicCompareExchangeWeak_83580d(threadgroup atomic_uint* const tint_symbol) {
+  atomic_compare_exchange_resultu32 res = atomicCompareExchangeWeak_1(tint_symbol, 1u, 1u);
+}
+
+void compute_main_inner(uint local_invocation_index, threadgroup atomic_uint* const tint_symbol_1) {
+  {
+    atomic_store_explicit(tint_symbol_1, uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicCompareExchangeWeak_83580d(tint_symbol_1);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  compute_main_inner(local_invocation_index, &(tint_symbol_2));
+  return;
+}
+

test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 37
+; Bound: 36
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -9,12 +9,17 @@
                OpExecutionMode %compute_main LocalSize 1 1 1
                OpName %local_invocation_index_1 "local_invocation_index_1"
                OpName %arg_0 "arg_0"
-               OpName %atomicCompareExchangeWeak_b2ab2c "atomicCompareExchangeWeak_b2ab2c"
+               OpName %atomicCompareExchangeWeak_83580d "atomicCompareExchangeWeak_83580d"
+               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 %res "res"
                OpName %compute_main_inner "compute_main_inner"
                OpName %local_invocation_index "local_invocation_index"
                OpName %compute_main "compute_main"
                OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               OpMemberDecorate %__atomic_compare_exchange_resultu32 0 Offset 0
+               OpMemberDecorate %__atomic_compare_exchange_resultu32 1 Offset 4
        %uint = OpTypeInt 32 0
 %_ptr_Input_uint = OpTypePointer Input %uint
 %local_invocation_index_1 = OpVariable %_ptr_Input_uint Input
@@ -22,37 +27,36 @@
       %arg_0 = OpVariable %_ptr_Workgroup_uint Workgroup
        %void = OpTypeVoid
           %6 = OpTypeFunction %void
-     %v2uint = OpTypeVector %uint 2
+       %bool = OpTypeBool
+%__atomic_compare_exchange_resultu32 = OpTypeStruct %uint %bool
      %uint_2 = OpConstant %uint 2
      %uint_0 = OpConstant %uint 0
      %uint_1 = OpConstant %uint 1
-       %bool = OpTypeBool
+%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32
-%_ptr_Function_v2uint = OpTypePointer Function %v2uint
+         %21 = OpConstantNull %__atomic_compare_exchange_resultu32
-         %22 = OpConstantNull %v2uint
+         %22 = OpTypeFunction %void %uint
-         %23 = OpTypeFunction %void %uint
+         %28 = OpConstantNull %uint
-         %29 = OpConstantNull %uint
    %uint_264 = OpConstant %uint 264
-%atomicCompareExchangeWeak_b2ab2c = OpFunction %void None %6
+%atomicCompareExchangeWeak_83580d = OpFunction %void None %6
           %9 = OpLabel
-        %res = OpVariable %_ptr_Function_v2uint Function %22
+        %res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %21
          %17 = OpAtomicCompareExchange %uint %arg_0 %uint_2 %uint_0 %uint_0 %uint_1 %uint_1
          %18 = OpIEqual %bool %17 %uint_1
-         %19 = OpSelect %uint %18 %uint_1 %uint_0
+         %10 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %17 %18
-         %10 = OpCompositeConstruct %v2uint %17 %19
                OpStore %res %10
                OpReturn
                OpFunctionEnd
-%compute_main_inner = OpFunction %void None %23
+%compute_main_inner = OpFunction %void None %22
 %local_invocation_index = OpFunctionParameter %uint
-         %26 = OpLabel
+         %25 = OpLabel
-               OpAtomicStore %arg_0 %uint_2 %uint_0 %29
+               OpAtomicStore %arg_0 %uint_2 %uint_0 %28
                OpControlBarrier %uint_2 %uint_2 %uint_264
-         %32 = OpFunctionCall %void %atomicCompareExchangeWeak_b2ab2c
+         %31 = OpFunctionCall %void %atomicCompareExchangeWeak_83580d
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %6
-         %34 = OpLabel
+         %33 = OpLabel
-         %36 = OpLoad %uint %local_invocation_index_1
+         %35 = OpLoad %uint %local_invocation_index_1
-         %35 = OpFunctionCall %void %compute_main_inner %36
+         %34 = OpFunctionCall %void %compute_main_inner %35
                OpReturn
                OpFunctionEnd

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

@@ -0,0 +1,10 @@
+var<workgroup> arg_0 : atomic<u32>;
+
+fn atomicCompareExchangeWeak_83580d() {
+  var res = atomicCompareExchangeWeak(&(arg_0), 1u, 1u);
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  atomicCompareExchangeWeak_83580d();
+}

test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.glsl

@@ -1,37 +0,0 @@
-SKIP: FAILED
-
-#version 310 es
-
-ivec2 tint_atomicCompareExchangeWeak(inout int param_0, int param_1, int param_2) {
-  ivec2 result;
-  result.x = atomicCompSwap(param_0, param_1, param_2);
-  result.y = result.x == param_2 ? 1 : 0;
-  return result;
-}
-
-
-shared int arg_0;
-void atomicCompareExchangeWeak_89ea3b() {
-  ivec2 res = tint_atomicCompareExchangeWeak(arg_0, 1, 1);
-}
-
-void compute_main(uint local_invocation_index) {
-  {
-    atomicExchange(arg_0, 0);
-  }
-  barrier();
-  atomicCompareExchangeWeak_89ea3b();
-}
-
-layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
-void main() {
-  compute_main(gl_LocalInvocationIndex);
-  return;
-}
-Error parsing GLSL shader:
-ERROR: 0:5: 'atomicCompSwap' : Atomic memory function can only be used for shader storage block member or shared variable. 
-ERROR: 0:5: '' : compilation terminated 
-ERROR: 2 compilation errors.  No code generated.
-
-
-

test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl

@@ -1,29 +0,0 @@
-#include <metal_stdlib>
-
-using namespace metal;
-
-template <typename A, typename T>
-vec<T, 2> atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) {
-  T prev_value = compare;
-  bool matched = atomic_compare_exchange_weak_explicit(atomic, &prev_value, value, memory_order_relaxed, memory_order_relaxed);
-  return {prev_value, matched};
-}
-
-void atomicCompareExchangeWeak_89ea3b(threadgroup atomic_int* const tint_symbol) {
-  int2 res = atomicCompareExchangeWeak_1(tint_symbol, 1, 1);
-}
-
-void compute_main_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_1) {
-  {
-    atomic_store_explicit(tint_symbol_1, int(), memory_order_relaxed);
-  }
-  threadgroup_barrier(mem_flags::mem_threadgroup);
-  atomicCompareExchangeWeak_89ea3b(tint_symbol_1);
-}
-
-kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
-  threadgroup atomic_int tint_symbol_2;
-  compute_main_inner(local_invocation_index, &(tint_symbol_2));
-  return;
-}
-

test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.wgsl

@@ -1,10 +0,0 @@
-var<workgroup> arg_0 : atomic<i32>;
-
-fn atomicCompareExchangeWeak_89ea3b() {
-  var res : vec2<i32> = atomicCompareExchangeWeak(&(arg_0), 1, 1);
-}
-
-@stage(compute) @workgroup_size(1)
-fn compute_main() {
-  atomicCompareExchangeWeak_89ea3b();
-}

test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.glsl

@@ -1,37 +0,0 @@
-SKIP: FAILED
-
-#version 310 es
-
-uvec2 tint_atomicCompareExchangeWeak(inout uint param_0, uint param_1, uint param_2) {
-  uvec2 result;
-  result.x = atomicCompSwap(param_0, param_1, param_2);
-  result.y = result.x == param_2 ? 1u : 0u;
-  return result;
-}
-
-
-shared uint arg_0;
-void atomicCompareExchangeWeak_b2ab2c() {
-  uvec2 res = tint_atomicCompareExchangeWeak(arg_0, 1u, 1u);
-}
-
-void compute_main(uint local_invocation_index) {
-  {
-    atomicExchange(arg_0, 0u);
-  }
-  barrier();
-  atomicCompareExchangeWeak_b2ab2c();
-}
-
-layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
-void main() {
-  compute_main(gl_LocalInvocationIndex);
-  return;
-}
-Error parsing GLSL shader:
-ERROR: 0:5: 'atomicCompSwap' : Atomic memory function can only be used for shader storage block member or shared variable. 
-ERROR: 0:5: '' : compilation terminated 
-ERROR: 2 compilation errors.  No code generated.
-
-
-

test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl

@@ -1,29 +0,0 @@
-#include <metal_stdlib>
-
-using namespace metal;
-
-template <typename A, typename T>
-vec<T, 2> atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) {
-  T prev_value = compare;
-  bool matched = atomic_compare_exchange_weak_explicit(atomic, &prev_value, value, memory_order_relaxed, memory_order_relaxed);
-  return {prev_value, matched};
-}
-
-void atomicCompareExchangeWeak_b2ab2c(threadgroup atomic_uint* const tint_symbol) {
-  uint2 res = atomicCompareExchangeWeak_1(tint_symbol, 1u, 1u);
-}
-
-void compute_main_inner(uint local_invocation_index, threadgroup atomic_uint* const tint_symbol_1) {
-  {
-    atomic_store_explicit(tint_symbol_1, uint(), memory_order_relaxed);
-  }
-  threadgroup_barrier(mem_flags::mem_threadgroup);
-  atomicCompareExchangeWeak_b2ab2c(tint_symbol_1);
-}
-
-kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
-  threadgroup atomic_uint tint_symbol_2;
-  compute_main_inner(local_invocation_index, &(tint_symbol_2));
-  return;
-}
-

test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.wgsl

@@ -1,10 +0,0 @@
-var<workgroup> arg_0 : atomic<u32>;
-
-fn atomicCompareExchangeWeak_b2ab2c() {
-  var res : vec2<u32> = atomicCompareExchangeWeak(&(arg_0), 1u, 1u);
-}
-
-@stage(compute) @workgroup_size(1)
-fn compute_main() {
-  atomicCompareExchangeWeak_b2ab2c();
-}

test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl

@@ -24,12 +24,12 @@
 
 var<workgroup> arg_0: atomic<i32>;
 
-// fn atomicCompareExchangeWeak(ptr<workgroup, atomic<i32>, read_write>, i32, i32) -> vec2<i32>
+// fn atomicCompareExchangeWeak(ptr<workgroup, atomic<i32>, read_write>, i32, i32) -> __atomic_compare_exchange_result<i32>
-fn atomicCompareExchangeWeak_89ea3b() {
+fn atomicCompareExchangeWeak_e88938() {
-  var res: vec2<i32> = atomicCompareExchangeWeak(&arg_0, 1, 1);
+  var res = atomicCompareExchangeWeak(&arg_0, 1, 1);
 }
 
 @stage(compute) @workgroup_size(1)
 fn compute_main() {
-  atomicCompareExchangeWeak_89ea3b();
+  atomicCompareExchangeWeak_e88938();
 }

test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.glsl

@@ -0,0 +1,29 @@
+#version 310 es
+
+struct atomic_compare_exchange_resulti32 {
+  int old_value;
+  bool exchanged;
+};
+
+
+shared int arg_0;
+void atomicCompareExchangeWeak_e88938() {
+  atomic_compare_exchange_resulti32 atomic_compare_result;
+  atomic_compare_result.old_value = atomicCompSwap(arg_0, 1, 1);
+  atomic_compare_result.exchanged = atomic_compare_result.old_value == 1;
+  atomic_compare_exchange_resulti32 res = atomic_compare_result;
+}
+
+void compute_main(uint local_invocation_index) {
+  {
+    atomicExchange(arg_0, 0);
+  }
+  barrier();
+  atomicCompareExchangeWeak_e88938();
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  compute_main(gl_LocalInvocationIndex);
+  return;
+}

test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.hlsl

@@ -1,11 +1,15 @@
+struct atomic_compare_exchange_resulti32 {
+  int old_value;
+  bool exchanged;
+};
 groupshared int arg_0;
 
-void atomicCompareExchangeWeak_89ea3b() {
+void atomicCompareExchangeWeak_e88938() {
-  int2 atomic_result = int2(0, 0);
+  atomic_compare_exchange_resulti32 atomic_result = (atomic_compare_exchange_resulti32)0;
   int atomic_compare_value = 1;
-  InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.x);
+  InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.old_value);
-  atomic_result.y = atomic_result.x == atomic_compare_value;
+  atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
-  int2 res = atomic_result;
+  atomic_compare_exchange_resulti32 res = atomic_result;
 }
 
 struct tint_symbol_1 {
@@ -18,7 +22,7 @@ void compute_main_inner(uint local_invocation_index) {
     InterlockedExchange(arg_0, 0, atomic_result_1);
   }
   GroupMemoryBarrierWithGroupSync();
-  atomicCompareExchangeWeak_89ea3b();
+  atomicCompareExchangeWeak_e88938();
 }
 
 [numthreads(1, 1, 1)]

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

@@ -0,0 +1,33 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+struct atomic_compare_exchange_resulti32 {
+  int old_value;
+  bool exchanged;
+};
+template <typename A, typename T>
+atomic_compare_exchange_resulti32 atomicCompareExchangeWeak_1(threadgroup A* atomic, T compare, T value) {
+  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};
+}
+
+void atomicCompareExchangeWeak_e88938(threadgroup atomic_int* const tint_symbol) {
+  atomic_compare_exchange_resulti32 res = atomicCompareExchangeWeak_1(tint_symbol, 1, 1);
+}
+
+void compute_main_inner(uint local_invocation_index, threadgroup atomic_int* const tint_symbol_1) {
+  {
+    atomic_store_explicit(tint_symbol_1, int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicCompareExchangeWeak_e88938(tint_symbol_1);
+}
+
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  compute_main_inner(local_invocation_index, &(tint_symbol_2));
+  return;
+}
+

test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm

@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 39
+; Bound: 37
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -9,12 +9,17 @@
                OpExecutionMode %compute_main LocalSize 1 1 1
                OpName %local_invocation_index_1 "local_invocation_index_1"
                OpName %arg_0 "arg_0"
-               OpName %atomicCompareExchangeWeak_89ea3b "atomicCompareExchangeWeak_89ea3b"
+               OpName %atomicCompareExchangeWeak_e88938 "atomicCompareExchangeWeak_e88938"
+               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 %res "res"
                OpName %compute_main_inner "compute_main_inner"
                OpName %local_invocation_index "local_invocation_index"
                OpName %compute_main "compute_main"
                OpDecorate %local_invocation_index_1 BuiltIn LocalInvocationIndex
+               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
@@ -23,38 +28,36 @@
       %arg_0 = OpVariable %_ptr_Workgroup_int Workgroup
        %void = OpTypeVoid
           %7 = OpTypeFunction %void
-      %v2int = OpTypeVector %int 2
+       %bool = OpTypeBool
+%__atomic_compare_exchange_resulti32 = OpTypeStruct %int %bool
      %uint_2 = OpConstant %uint 2
      %uint_0 = OpConstant %uint 0
       %int_1 = OpConstant %int 1
-       %bool = OpTypeBool
+%_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32
-      %int_0 = OpConstant %int 0
+         %22 = OpConstantNull %__atomic_compare_exchange_resulti32
-%_ptr_Function_v2int = OpTypePointer Function %v2int
+         %23 = OpTypeFunction %void %uint
-         %24 = OpConstantNull %v2int
+         %29 = OpConstantNull %int
-         %25 = OpTypeFunction %void %uint
-         %31 = OpConstantNull %int
    %uint_264 = OpConstant %uint 264
-%atomicCompareExchangeWeak_89ea3b = OpFunction %void None %7
+%atomicCompareExchangeWeak_e88938 = OpFunction %void None %7
          %10 = OpLabel
-        %res = OpVariable %_ptr_Function_v2int Function %24
+        %res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %22
          %18 = OpAtomicCompareExchange %int %arg_0 %uint_2 %uint_0 %uint_0 %int_1 %int_1
          %19 = OpIEqual %bool %18 %int_1
-         %21 = OpSelect %int %19 %int_1 %int_0
+         %11 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %18 %19
-         %11 = OpCompositeConstruct %v2int %18 %21
                OpStore %res %11
                OpReturn
                OpFunctionEnd
-%compute_main_inner = OpFunction %void None %25
+%compute_main_inner = OpFunction %void None %23
 %local_invocation_index = OpFunctionParameter %uint
-         %28 = OpLabel
+         %26 = OpLabel
-               OpAtomicStore %arg_0 %uint_2 %uint_0 %31
+               OpAtomicStore %arg_0 %uint_2 %uint_0 %29
                OpControlBarrier %uint_2 %uint_2 %uint_264
-         %34 = OpFunctionCall %void %atomicCompareExchangeWeak_89ea3b
+         %32 = OpFunctionCall %void %atomicCompareExchangeWeak_e88938
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %7
-         %36 = OpLabel
+         %34 = OpLabel
-         %38 = OpLoad %uint %local_invocation_index_1
+         %36 = OpLoad %uint %local_invocation_index_1
-         %37 = OpFunctionCall %void %compute_main_inner %38
+         %35 = OpFunctionCall %void %compute_main_inner %36
                OpReturn
                OpFunctionEnd

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

@@ -0,0 +1,10 @@
+var<workgroup> arg_0 : atomic<i32>;
+
+fn atomicCompareExchangeWeak_e88938() {
+  var res = atomicCompareExchangeWeak(&(arg_0), 1, 1);
+}
+
+@stage(compute) @workgroup_size(1)
+fn compute_main() {
+  atomicCompareExchangeWeak_e88938();
+}

test/tint/builtins/gen/atomicExchange/d59712.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-uint atomicExchange_1(RWByteAddressBuffer buffer, uint offset, uint value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+uint tint_atomicExchange(RWByteAddressBuffer buffer, uint offset, uint value) {
   uint original_value = 0;
   buffer.InterlockedExchange(offset, value, original_value);
   return original_value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicExchange_d59712() {
-  uint res = atomicExchange_1(sb_rw, 0u, 1u);
+  uint res = tint_atomicExchange(sb_rw, 0u, 1u);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicExchange/f2e22f.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-int atomicExchange_1(RWByteAddressBuffer buffer, uint offset, int value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+int tint_atomicExchange(RWByteAddressBuffer buffer, uint offset, int value) {
   int original_value = 0;
   buffer.InterlockedExchange(offset, value, original_value);
   return original_value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicExchange_f2e22f() {
-  int res = atomicExchange_1(sb_rw, 0u, 1);
+  int res = tint_atomicExchange(sb_rw, 0u, 1);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicLoad/0806ad.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-int atomicLoad_1(RWByteAddressBuffer buffer, uint offset) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+int tint_atomicLoad(RWByteAddressBuffer buffer, uint offset) {
   int value = 0;
   buffer.InterlockedOr(offset, 0, value);
   return value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicLoad_0806ad() {
-  int res = atomicLoad_1(sb_rw, 0u);
+  int res = tint_atomicLoad(sb_rw, 0u);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicLoad/fe6cc3.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-uint atomicLoad_1(RWByteAddressBuffer buffer, uint offset) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+uint tint_atomicLoad(RWByteAddressBuffer buffer, uint offset) {
   uint value = 0;
   buffer.InterlockedOr(offset, 0, value);
   return value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicLoad_fe6cc3() {
-  uint res = atomicLoad_1(sb_rw, 0u);
+  uint res = tint_atomicLoad(sb_rw, 0u);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicMax/51b9be.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-uint atomicMax_1(RWByteAddressBuffer buffer, uint offset, uint value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+uint tint_atomicMax(RWByteAddressBuffer buffer, uint offset, uint value) {
   uint original_value = 0;
   buffer.InterlockedMax(offset, value, original_value);
   return original_value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicMax_51b9be() {
-  uint res = atomicMax_1(sb_rw, 0u, 1u);
+  uint res = tint_atomicMax(sb_rw, 0u, 1u);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicMax/92aa72.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-int atomicMax_1(RWByteAddressBuffer buffer, uint offset, int value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+int tint_atomicMax(RWByteAddressBuffer buffer, uint offset, int value) {
   int original_value = 0;
   buffer.InterlockedMax(offset, value, original_value);
   return original_value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicMax_92aa72() {
-  int res = atomicMax_1(sb_rw, 0u, 1);
+  int res = tint_atomicMax(sb_rw, 0u, 1);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicMin/8e38dc.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-int atomicMin_1(RWByteAddressBuffer buffer, uint offset, int value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+int tint_atomicMin(RWByteAddressBuffer buffer, uint offset, int value) {
   int original_value = 0;
   buffer.InterlockedMin(offset, value, original_value);
   return original_value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicMin_8e38dc() {
-  int res = atomicMin_1(sb_rw, 0u, 1);
+  int res = tint_atomicMin(sb_rw, 0u, 1);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicMin/c67a74.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-uint atomicMin_1(RWByteAddressBuffer buffer, uint offset, uint value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+uint tint_atomicMin(RWByteAddressBuffer buffer, uint offset, uint value) {
   uint original_value = 0;
   buffer.InterlockedMin(offset, value, original_value);
   return original_value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicMin_c67a74() {
-  uint res = atomicMin_1(sb_rw, 0u, 1u);
+  uint res = tint_atomicMin(sb_rw, 0u, 1u);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicOr/5e95d4.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-uint atomicOr_1(RWByteAddressBuffer buffer, uint offset, uint value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+uint tint_atomicOr(RWByteAddressBuffer buffer, uint offset, uint value) {
   uint original_value = 0;
   buffer.InterlockedOr(offset, value, original_value);
   return original_value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicOr_5e95d4() {
-  uint res = atomicOr_1(sb_rw, 0u, 1u);
+  uint res = tint_atomicOr(sb_rw, 0u, 1u);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicOr/8d96a0.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-int atomicOr_1(RWByteAddressBuffer buffer, uint offset, int value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+int tint_atomicOr(RWByteAddressBuffer buffer, uint offset, int value) {
   int original_value = 0;
   buffer.InterlockedOr(offset, value, original_value);
   return original_value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicOr_8d96a0() {
-  int res = atomicOr_1(sb_rw, 0u, 1);
+  int res = tint_atomicOr(sb_rw, 0u, 1);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicStore/cdc29e.wgsl.expected.hlsl

@@ -1,12 +1,13 @@
-void atomicStore_1(RWByteAddressBuffer buffer, uint offset, uint value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+void tint_atomicStore(RWByteAddressBuffer buffer, uint offset, uint value) {
   uint ignored;
   buffer.InterlockedExchange(offset, value, ignored);
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicStore_cdc29e() {
-  atomicStore_1(sb_rw, 0u, 1u);
+  tint_atomicStore(sb_rw, 0u, 1u);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicStore/d1e9a6.wgsl.expected.hlsl

@@ -1,12 +1,13 @@
-void atomicStore_1(RWByteAddressBuffer buffer, uint offset, int value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+void tint_atomicStore(RWByteAddressBuffer buffer, uint offset, int value) {
   int ignored;
   buffer.InterlockedExchange(offset, value, ignored);
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicStore_d1e9a6() {
-  atomicStore_1(sb_rw, 0u, 1);
+  tint_atomicStore(sb_rw, 0u, 1);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicSub/051100.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-int atomicSub_1(RWByteAddressBuffer buffer, uint offset, int value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+int tint_atomicSub(RWByteAddressBuffer buffer, uint offset, int value) {
   int original_value = 0;
   buffer.InterlockedAdd(offset, -value, original_value);
   return original_value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicSub_051100() {
-  int res = atomicSub_1(sb_rw, 0u, 1);
+  int res = tint_atomicSub(sb_rw, 0u, 1);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicSub/15bfc9.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-uint atomicSub_1(RWByteAddressBuffer buffer, uint offset, uint value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+uint tint_atomicSub(RWByteAddressBuffer buffer, uint offset, uint value) {
   uint original_value = 0;
   buffer.InterlockedAdd(offset, -value, original_value);
   return original_value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicSub_15bfc9() {
-  uint res = atomicSub_1(sb_rw, 0u, 1u);
+  uint res = tint_atomicSub(sb_rw, 0u, 1u);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicXor/54510e.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-uint atomicXor_1(RWByteAddressBuffer buffer, uint offset, uint value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+uint tint_atomicXor(RWByteAddressBuffer buffer, uint offset, uint value) {
   uint original_value = 0;
   buffer.InterlockedXor(offset, value, original_value);
   return original_value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicXor_54510e() {
-  uint res = atomicXor_1(sb_rw, 0u, 1u);
+  uint res = tint_atomicXor(sb_rw, 0u, 1u);
 }
 
 void fragment_main() {

test/tint/builtins/gen/atomicXor/c1b78c.wgsl.expected.hlsl

@@ -1,13 +1,14 @@
-int atomicXor_1(RWByteAddressBuffer buffer, uint offset, int value) {
+RWByteAddressBuffer sb_rw : register(u0, space0);
+
+int tint_atomicXor(RWByteAddressBuffer buffer, uint offset, int value) {
   int original_value = 0;
   buffer.InterlockedXor(offset, value, original_value);
   return original_value;
 }
 
-RWByteAddressBuffer sb_rw : register(u0, space0);
 
 void atomicXor_c1b78c() {
-  int res = atomicXor_1(sb_rw, 0u, 1);
+  int res = tint_atomicXor(sb_rw, 0u, 1);
 }
 
 void fragment_main() {