diff --git a/src/tint/intrinsics.def b/src/tint/intrinsics.def
index bc4bf28dfd..a6792b13ae 100644
--- a/src/tint/intrinsics.def
+++ b/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                                                          //
diff --git a/src/tint/resolver/intrinsic_table.cc b/src/tint/resolver/intrinsic_table.cc
index 6c19f7cd94..85711769f1 100644
--- a/src/tint/resolver/intrinsic_table.cc
+++ b/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)
diff --git a/src/tint/resolver/intrinsic_table.inl b/src/tint/resolver/intrinsic_table.inl
index c518e58f7d..7a422fb6f7 100644
--- a/src/tint/resolver/intrinsic_table.inl
+++ b/src/tint/resolver/intrinsic_table.inl
@@ -1512,8 +1512,41 @@ std::string FrexpResultVec::String(MatchState* state) const {
   return ss.str();
 }
 
+/// TypeMatcher for 'type __atomic_compare_exchange_result'
+/// @see src/tint/intrinsics.def:121:6
+class AtomicCompareExchangeResult : public TypeMatcher {
+ public:
+  /// Checks whether the given type matches the matcher rules.
+  /// Match may define and refine the template types and numbers in state.
+  /// @param state the MatchState
+  /// @param type the type to match
+  /// @returns the canonicalized type on match, otherwise nullptr
+  const sem::Type* Match(MatchState& state,
+                         const sem::Type* type) const override;
+  /// @param state the MatchState
+  /// @return a string representation of the matcher.
+  std::string String(MatchState* state) const override;
+};
+
+const sem::Type* AtomicCompareExchangeResult::Match(MatchState& state, const sem::Type* ty) const {
+  const sem::Type* T = nullptr;
+  if (!match_atomic_compare_exchange_result(ty, T)) {
+    return nullptr;
+  }
+  T = state.Type(T);
+  if (T == nullptr) {
+    return nullptr;
+  }
+  return build_atomic_compare_exchange_result(state, T);
+}
+
+std::string AtomicCompareExchangeResult::String(MatchState* state) const {
+  const std::string T = state->TypeName();
+  return "__atomic_compare_exchange_result<" + T + ">";
+}
+
 /// TypeMatcher for 'match fiu32'
-/// @see src/tint/intrinsics.def:127:7
+/// @see src/tint/intrinsics.def:129:7
 class Fiu32 : public TypeMatcher {
  public:
   /// Checks whether the given type matches the matcher rules, and returns the
@@ -1551,7 +1584,7 @@ std::string Fiu32::String(MatchState*) const {
 }
 
 /// TypeMatcher for 'match fi32'
-/// @see src/tint/intrinsics.def:128:7
+/// @see src/tint/intrinsics.def:130:7
 class Fi32 : public TypeMatcher {
  public:
   /// Checks whether the given type matches the matcher rules, and returns the
@@ -1586,7 +1619,7 @@ std::string Fi32::String(MatchState*) const {
 }
 
 /// TypeMatcher for 'match iu32'
-/// @see src/tint/intrinsics.def:129:7
+/// @see src/tint/intrinsics.def:131:7
 class Iu32 : public TypeMatcher {
  public:
   /// Checks whether the given type matches the matcher rules, and returns the
@@ -1621,7 +1654,7 @@ std::string Iu32::String(MatchState*) const {
 }
 
 /// TypeMatcher for 'match scalar'
-/// @see src/tint/intrinsics.def:130:7
+/// @see src/tint/intrinsics.def:132:7
 class Scalar : public TypeMatcher {
  public:
   /// Checks whether the given type matches the matcher rules, and returns the
@@ -1662,7 +1695,7 @@ std::string Scalar::String(MatchState*) const {
 }
 
 /// TypeMatcher for 'match abstract_or_scalar'
-/// @see src/tint/intrinsics.def:131:7
+/// @see src/tint/intrinsics.def:133:7
 class AbstractOrScalar : public TypeMatcher {
  public:
   /// Checks whether the given type matches the matcher rules, and returns the
@@ -1709,7 +1742,7 @@ std::string AbstractOrScalar::String(MatchState*) const {
 }
 
 /// TypeMatcher for 'match af_f32'
-/// @see src/tint/intrinsics.def:132:7
+/// @see src/tint/intrinsics.def:134:7
 class AfF32 : public TypeMatcher {
  public:
   /// Checks whether the given type matches the matcher rules, and returns the
@@ -1744,7 +1777,7 @@ std::string AfF32::String(MatchState*) const {
 }
 
 /// TypeMatcher for 'match scalar_no_f32'
-/// @see src/tint/intrinsics.def:133:7
+/// @see src/tint/intrinsics.def:135:7
 class ScalarNoF32 : public TypeMatcher {
  public:
   /// Checks whether the given type matches the matcher rules, and returns the
@@ -1782,7 +1815,7 @@ std::string ScalarNoF32::String(MatchState*) const {
 }
 
 /// TypeMatcher for 'match scalar_no_i32'
-/// @see src/tint/intrinsics.def:134:7
+/// @see src/tint/intrinsics.def:136:7
 class ScalarNoI32 : public TypeMatcher {
  public:
   /// Checks whether the given type matches the matcher rules, and returns the
@@ -1820,7 +1853,7 @@ std::string ScalarNoI32::String(MatchState*) const {
 }
 
 /// TypeMatcher for 'match scalar_no_u32'
-/// @see src/tint/intrinsics.def:135:7
+/// @see src/tint/intrinsics.def:137:7
 class ScalarNoU32 : public TypeMatcher {
  public:
   /// Checks whether the given type matches the matcher rules, and returns the
@@ -1858,7 +1891,7 @@ std::string ScalarNoU32::String(MatchState*) const {
 }
 
 /// TypeMatcher for 'match scalar_no_bool'
-/// @see src/tint/intrinsics.def:136:7
+/// @see src/tint/intrinsics.def:138:7
 class ScalarNoBool : public TypeMatcher {
  public:
   /// Checks whether the given type matches the matcher rules, and returns the
@@ -1896,7 +1929,7 @@ std::string ScalarNoBool::String(MatchState*) const {
 }
 
 /// EnumMatcher for 'match f32_texel_format'
-/// @see src/tint/intrinsics.def:147:7
+/// @see src/tint/intrinsics.def:149:7
 class F32TexelFormat : public NumberMatcher {
  public:
   /// Checks whether the given number matches the enum matcher rules.
@@ -1929,7 +1962,7 @@ std::string F32TexelFormat::String(MatchState*) const {
 }
 
 /// EnumMatcher for 'match i32_texel_format'
-/// @see src/tint/intrinsics.def:149:7
+/// @see src/tint/intrinsics.def:151:7
 class I32TexelFormat : public NumberMatcher {
  public:
   /// Checks whether the given number matches the enum matcher rules.
@@ -1961,7 +1994,7 @@ std::string I32TexelFormat::String(MatchState*) const {
 }
 
 /// EnumMatcher for 'match u32_texel_format'
-/// @see src/tint/intrinsics.def:151:7
+/// @see src/tint/intrinsics.def:153:7
 class U32TexelFormat : public NumberMatcher {
  public:
   /// Checks whether the given number matches the enum matcher rules.
@@ -1993,7 +2026,7 @@ std::string U32TexelFormat::String(MatchState*) const {
 }
 
 /// EnumMatcher for 'match write_only'
-/// @see src/tint/intrinsics.def:154:7
+/// @see src/tint/intrinsics.def:156:7
 class WriteOnly : public NumberMatcher {
  public:
   /// Checks whether the given number matches the enum matcher rules.
@@ -2019,7 +2052,7 @@ std::string WriteOnly::String(MatchState*) const {
 }
 
 /// EnumMatcher for 'match function_private_workgroup'
-/// @see src/tint/intrinsics.def:156:7
+/// @see src/tint/intrinsics.def:158:7
 class FunctionPrivateWorkgroup : public NumberMatcher {
  public:
   /// Checks whether the given number matches the enum matcher rules.
@@ -2049,7 +2082,7 @@ std::string FunctionPrivateWorkgroup::String(MatchState*) const {
 }
 
 /// EnumMatcher for 'match workgroup_or_storage'
-/// @see src/tint/intrinsics.def:157:7
+/// @see src/tint/intrinsics.def:159:7
 class WorkgroupOrStorage : public NumberMatcher {
  public:
   /// Checks whether the given number matches the enum matcher rules.
@@ -2206,6 +2239,7 @@ class Matchers {
   ModfResultVec ModfResultVec_;
   FrexpResult FrexpResult_;
   FrexpResultVec FrexpResultVec_;
+  AtomicCompareExchangeResult AtomicCompareExchangeResult_;
   Fiu32 Fiu32_;
   Fi32 Fi32_;
   Iu32 Iu32_;
@@ -2233,7 +2267,7 @@ class Matchers {
   ~Matchers();
 
   /// The template types, types, and type matchers
-  TypeMatcher const* const type[58] = {
+  TypeMatcher const* const type[59] = {
     /* [0] */ &template_type_0_,
     /* [1] */ &template_type_1_,
     /* [2] */ &Bool_,
@@ -2282,16 +2316,17 @@ class Matchers {
     /* [45] */ &ModfResultVec_,
     /* [46] */ &FrexpResult_,
     /* [47] */ &FrexpResultVec_,
-    /* [48] */ &Fiu32_,
-    /* [49] */ &Fi32_,
-    /* [50] */ &Iu32_,
-    /* [51] */ &Scalar_,
-    /* [52] */ &AbstractOrScalar_,
-    /* [53] */ &AfF32_,
-    /* [54] */ &ScalarNoF32_,
-    /* [55] */ &ScalarNoI32_,
-    /* [56] */ &ScalarNoU32_,
-    /* [57] */ &ScalarNoBool_,
+    /* [48] */ &AtomicCompareExchangeResult_,
+    /* [49] */ &Fiu32_,
+    /* [50] */ &Fi32_,
+    /* [51] */ &Iu32_,
+    /* [52] */ &Scalar_,
+    /* [53] */ &AbstractOrScalar_,
+    /* [54] */ &AfF32_,
+    /* [55] */ &ScalarNoF32_,
+    /* [56] */ &ScalarNoI32_,
+    /* [57] */ &ScalarNoU32_,
+    /* [58] */ &ScalarNoBool_,
   };
 
   /// The template numbers, and number matchers
@@ -2488,34 +2523,36 @@ constexpr MatcherIndex kMatcherIndices[] = {
   /* [170] */ 7,
   /* [171] */ 17,
   /* [172] */ 0,
-  /* [173] */ 18,
-  /* [174] */ 7,
+  /* [173] */ 48,
+  /* [174] */ 0,
   /* [175] */ 18,
-  /* [176] */ 0,
-  /* [177] */ 27,
-  /* [178] */ 7,
-  /* [179] */ 28,
+  /* [176] */ 7,
+  /* [177] */ 18,
+  /* [178] */ 0,
+  /* [179] */ 27,
   /* [180] */ 7,
-  /* [181] */ 29,
+  /* [181] */ 28,
   /* [182] */ 7,
-  /* [183] */ 19,
+  /* [183] */ 29,
   /* [184] */ 7,
-  /* [185] */ 30,
+  /* [185] */ 19,
   /* [186] */ 7,
-  /* [187] */ 31,
+  /* [187] */ 30,
   /* [188] */ 7,
-  /* [189] */ 32,
+  /* [189] */ 31,
   /* [190] */ 7,
-  /* [191] */ 25,
-  /* [192] */ 26,
-  /* [193] */ 37,
-  /* [194] */ 36,
-  /* [195] */ 35,
-  /* [196] */ 34,
-  /* [197] */ 43,
-  /* [198] */ 38,
-  /* [199] */ 44,
-  /* [200] */ 46,
+  /* [191] */ 32,
+  /* [192] */ 7,
+  /* [193] */ 25,
+  /* [194] */ 26,
+  /* [195] */ 37,
+  /* [196] */ 36,
+  /* [197] */ 35,
+  /* [198] */ 34,
+  /* [199] */ 43,
+  /* [200] */ 38,
+  /* [201] */ 44,
+  /* [202] */ 46,
 };
 
 // Assert that the MatcherIndex is big enough to index all the matchers, plus
@@ -2853,12 +2890,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [65] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[181],
+    /* matcher indices */ &kMatcherIndices[183],
   },
   {
     /* [66] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [67] */
@@ -2888,12 +2925,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [72] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[195],
+    /* matcher indices */ &kMatcherIndices[197],
   },
   {
     /* [73] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [74] */
@@ -2948,12 +2985,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [84] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[179],
+    /* matcher indices */ &kMatcherIndices[181],
   },
   {
     /* [85] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [86] */
@@ -3018,7 +3055,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [98] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [99] */
@@ -3038,12 +3075,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [102] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[181],
+    /* matcher indices */ &kMatcherIndices[183],
   },
   {
     /* [103] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [104] */
@@ -3068,12 +3105,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [108] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[195],
+    /* matcher indices */ &kMatcherIndices[197],
   },
   {
     /* [109] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [110] */
@@ -3098,12 +3135,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [114] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[185],
+    /* matcher indices */ &kMatcherIndices[187],
   },
   {
     /* [115] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [116] */
@@ -3128,12 +3165,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [120] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[189],
+    /* matcher indices */ &kMatcherIndices[191],
   },
   {
     /* [121] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [122] */
@@ -3158,12 +3195,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [126] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[181],
+    /* matcher indices */ &kMatcherIndices[183],
   },
   {
     /* [127] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [128] */
@@ -3188,12 +3225,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [132] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[181],
+    /* matcher indices */ &kMatcherIndices[183],
   },
   {
     /* [133] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [134] */
@@ -3218,12 +3255,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [138] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[195],
+    /* matcher indices */ &kMatcherIndices[197],
   },
   {
     /* [139] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [140] */
@@ -3248,12 +3285,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [144] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[195],
+    /* matcher indices */ &kMatcherIndices[197],
   },
   {
     /* [145] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [146] */
@@ -3278,12 +3315,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [150] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[195],
+    /* matcher indices */ &kMatcherIndices[197],
   },
   {
     /* [151] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [152] */
@@ -3303,12 +3340,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [155] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[196],
+    /* matcher indices */ &kMatcherIndices[198],
   },
   {
     /* [156] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [157] */
@@ -3328,12 +3365,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [160] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[193],
+    /* matcher indices */ &kMatcherIndices[195],
   },
   {
     /* [161] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [162] */
@@ -3353,12 +3390,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [165] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[189],
+    /* matcher indices */ &kMatcherIndices[191],
   },
   {
     /* [166] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [167] */
@@ -3378,12 +3415,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [170] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[193],
+    /* matcher indices */ &kMatcherIndices[195],
   },
   {
     /* [171] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [172] */
@@ -3403,12 +3440,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [175] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[185],
+    /* matcher indices */ &kMatcherIndices[187],
   },
   {
     /* [176] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [177] */
@@ -3428,12 +3465,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [180] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[179],
+    /* matcher indices */ &kMatcherIndices[181],
   },
   {
     /* [181] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [182] */
@@ -3453,12 +3490,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [185] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[195],
+    /* matcher indices */ &kMatcherIndices[197],
   },
   {
     /* [186] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [187] */
@@ -3478,12 +3515,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [190] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[181],
+    /* matcher indices */ &kMatcherIndices[183],
   },
   {
     /* [191] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [192] */
@@ -3503,12 +3540,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [195] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[179],
+    /* matcher indices */ &kMatcherIndices[181],
   },
   {
     /* [196] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [197] */
@@ -3528,12 +3565,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [200] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[195],
+    /* matcher indices */ &kMatcherIndices[197],
   },
   {
     /* [201] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [202] */
@@ -3553,12 +3590,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [205] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[196],
+    /* matcher indices */ &kMatcherIndices[198],
   },
   {
     /* [206] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [207] */
@@ -3578,12 +3615,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [210] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[187],
+    /* matcher indices */ &kMatcherIndices[189],
   },
   {
     /* [211] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [212] */
@@ -3603,12 +3640,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [215] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[181],
+    /* matcher indices */ &kMatcherIndices[183],
   },
   {
     /* [216] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [217] */
@@ -3628,12 +3665,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [220] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[185],
+    /* matcher indices */ &kMatcherIndices[187],
   },
   {
     /* [221] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [222] */
@@ -3653,12 +3690,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [225] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[195],
+    /* matcher indices */ &kMatcherIndices[197],
   },
   {
     /* [226] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [227] */
@@ -3688,7 +3725,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [232] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [233] */
@@ -3713,7 +3750,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [237] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [238] */
@@ -3728,12 +3765,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [240] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[179],
+    /* matcher indices */ &kMatcherIndices[181],
   },
   {
     /* [241] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [242] */
@@ -3753,12 +3790,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [245] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[193],
+    /* matcher indices */ &kMatcherIndices[195],
   },
   {
     /* [246] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [247] */
@@ -3778,12 +3815,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [250] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[185],
+    /* matcher indices */ &kMatcherIndices[187],
   },
   {
     /* [251] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [252] */
@@ -3813,7 +3850,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [257] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [258] */
@@ -3828,12 +3865,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [260] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[195],
+    /* matcher indices */ &kMatcherIndices[197],
   },
   {
     /* [261] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [262] */
@@ -3853,12 +3890,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [265] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[196],
+    /* matcher indices */ &kMatcherIndices[198],
   },
   {
     /* [266] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [267] */
@@ -3878,12 +3915,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [270] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[189],
+    /* matcher indices */ &kMatcherIndices[191],
   },
   {
     /* [271] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [272] */
@@ -3903,12 +3940,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [275] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[193],
+    /* matcher indices */ &kMatcherIndices[195],
   },
   {
     /* [276] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [277] */
@@ -3928,12 +3965,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [280] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[196],
+    /* matcher indices */ &kMatcherIndices[198],
   },
   {
     /* [281] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [282] */
@@ -3953,12 +3990,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [285] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[195],
+    /* matcher indices */ &kMatcherIndices[197],
   },
   {
     /* [286] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [287] */
@@ -3978,12 +4015,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [290] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[181],
+    /* matcher indices */ &kMatcherIndices[183],
   },
   {
     /* [291] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [292] */
@@ -4003,12 +4040,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [295] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[194],
+    /* matcher indices */ &kMatcherIndices[196],
   },
   {
     /* [296] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [297] */
@@ -4023,12 +4060,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [299] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[196],
+    /* matcher indices */ &kMatcherIndices[198],
   },
   {
     /* [300] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [301] */
@@ -4043,12 +4080,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [303] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[196],
+    /* matcher indices */ &kMatcherIndices[198],
   },
   {
     /* [304] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [305] */
@@ -4063,12 +4100,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [307] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[187],
+    /* matcher indices */ &kMatcherIndices[189],
   },
   {
     /* [308] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [309] */
@@ -4083,12 +4120,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [311] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[194],
+    /* matcher indices */ &kMatcherIndices[196],
   },
   {
     /* [312] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [313] */
@@ -4103,12 +4140,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [315] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[185],
+    /* matcher indices */ &kMatcherIndices[187],
   },
   {
     /* [316] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [317] */
@@ -4123,12 +4160,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [319] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[179],
+    /* matcher indices */ &kMatcherIndices[181],
   },
   {
     /* [320] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [321] */
@@ -4163,12 +4200,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [327] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[185],
+    /* matcher indices */ &kMatcherIndices[187],
   },
   {
     /* [328] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [329] */
@@ -4183,12 +4220,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [331] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[179],
+    /* matcher indices */ &kMatcherIndices[181],
   },
   {
     /* [332] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [333] */
@@ -4203,12 +4240,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [335] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[193],
+    /* matcher indices */ &kMatcherIndices[195],
   },
   {
     /* [336] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [337] */
@@ -4223,12 +4260,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [339] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[187],
+    /* matcher indices */ &kMatcherIndices[189],
   },
   {
     /* [340] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [341] */
@@ -4243,12 +4280,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [343] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[195],
+    /* matcher indices */ &kMatcherIndices[197],
   },
   {
     /* [344] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [345] */
@@ -4263,12 +4300,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [347] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[196],
+    /* matcher indices */ &kMatcherIndices[198],
   },
   {
     /* [348] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [349] */
@@ -4283,12 +4320,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [351] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[189],
+    /* matcher indices */ &kMatcherIndices[191],
   },
   {
     /* [352] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [353] */
@@ -4303,12 +4340,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [355] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[185],
+    /* matcher indices */ &kMatcherIndices[187],
   },
   {
     /* [356] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [357] */
@@ -4333,7 +4370,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [361] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [362] */
@@ -4343,12 +4380,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [363] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[181],
+    /* matcher indices */ &kMatcherIndices[183],
   },
   {
     /* [364] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [365] */
@@ -4363,12 +4400,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [367] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[179],
+    /* matcher indices */ &kMatcherIndices[181],
   },
   {
     /* [368] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [369] */
@@ -4383,12 +4420,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [371] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[196],
+    /* matcher indices */ &kMatcherIndices[198],
   },
   {
     /* [372] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [373] */
@@ -4403,12 +4440,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [375] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[194],
+    /* matcher indices */ &kMatcherIndices[196],
   },
   {
     /* [376] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [377] */
@@ -4423,12 +4460,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [379] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[194],
+    /* matcher indices */ &kMatcherIndices[196],
   },
   {
     /* [380] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [381] */
@@ -4483,12 +4520,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [391] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[196],
+    /* matcher indices */ &kMatcherIndices[198],
   },
   {
     /* [392] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[192],
+    /* matcher indices */ &kMatcherIndices[194],
   },
   {
     /* [393] */
@@ -4503,12 +4540,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [395] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[193],
+    /* matcher indices */ &kMatcherIndices[195],
   },
   {
     /* [396] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [397] */
@@ -4563,7 +4600,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [407] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[195],
+    /* matcher indices */ &kMatcherIndices[197],
   },
   {
     /* [408] */
@@ -4583,12 +4620,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [411] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[195],
+    /* matcher indices */ &kMatcherIndices[197],
   },
   {
     /* [412] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [413] */
@@ -4603,12 +4640,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [415] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[196],
+    /* matcher indices */ &kMatcherIndices[198],
   },
   {
     /* [416] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [417] */
@@ -4653,7 +4690,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [425] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [426] */
@@ -4763,12 +4800,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [447] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[196],
+    /* matcher indices */ &kMatcherIndices[198],
   },
   {
     /* [448] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [449] */
@@ -4808,12 +4845,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [456] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[197],
+    /* matcher indices */ &kMatcherIndices[199],
   },
   {
     /* [457] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [458] */
@@ -5198,12 +5235,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [534] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[194],
+    /* matcher indices */ &kMatcherIndices[196],
   },
   {
     /* [535] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [536] */
@@ -5258,12 +5295,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [546] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[196],
+    /* matcher indices */ &kMatcherIndices[198],
   },
   {
     /* [547] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [548] */
@@ -5348,12 +5385,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [564] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[187],
+    /* matcher indices */ &kMatcherIndices[189],
   },
   {
     /* [565] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [566] */
@@ -5378,12 +5415,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [570] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[185],
+    /* matcher indices */ &kMatcherIndices[187],
   },
   {
     /* [571] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [572] */
@@ -5423,7 +5460,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [579] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[196],
+    /* matcher indices */ &kMatcherIndices[198],
   },
   {
     /* [580] */
@@ -5438,12 +5475,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [582] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[179],
+    /* matcher indices */ &kMatcherIndices[181],
   },
   {
     /* [583] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [584] */
@@ -5453,12 +5490,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [585] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[177],
+    /* matcher indices */ &kMatcherIndices[179],
   },
   {
     /* [586] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [587] */
@@ -5468,7 +5505,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [588] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[198],
+    /* matcher indices */ &kMatcherIndices[200],
   },
   {
     /* [589] */
@@ -5498,12 +5535,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [594] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[194],
+    /* matcher indices */ &kMatcherIndices[196],
   },
   {
     /* [595] */
     /* usage */ ParameterUsage::kSampler,
-    /* matcher indices */ &kMatcherIndices[191],
+    /* matcher indices */ &kMatcherIndices[193],
   },
   {
     /* [596] */
@@ -5713,7 +5750,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [637] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[196],
+    /* matcher indices */ &kMatcherIndices[198],
   },
   {
     /* [638] */
@@ -5733,7 +5770,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [641] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[195],
+    /* matcher indices */ &kMatcherIndices[197],
   },
   {
     /* [642] */
@@ -5753,7 +5790,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [645] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[194],
+    /* matcher indices */ &kMatcherIndices[196],
   },
   {
     /* [646] */
@@ -5773,7 +5810,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [649] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[193],
+    /* matcher indices */ &kMatcherIndices[195],
   },
   {
     /* [650] */
@@ -6163,7 +6200,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [727] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[197],
+    /* matcher indices */ &kMatcherIndices[199],
   },
   {
     /* [728] */
@@ -6748,7 +6785,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [844] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[198],
+    /* matcher indices */ &kMatcherIndices[200],
   },
   {
     /* [845] */
@@ -6758,7 +6795,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [846] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[193],
+    /* matcher indices */ &kMatcherIndices[195],
   },
   {
     /* [847] */
@@ -6768,17 +6805,17 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [848] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[194],
+    /* matcher indices */ &kMatcherIndices[196],
   },
   {
     /* [849] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[195],
+    /* matcher indices */ &kMatcherIndices[197],
   },
   {
     /* [850] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[196],
+    /* matcher indices */ &kMatcherIndices[198],
   },
   {
     /* [851] */
@@ -6848,12 +6885,12 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [864] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[193],
+    /* matcher indices */ &kMatcherIndices[195],
   },
   {
     /* [865] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[195],
+    /* matcher indices */ &kMatcherIndices[197],
   },
   {
     /* [866] */
@@ -6943,7 +6980,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [883] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[197],
+    /* matcher indices */ &kMatcherIndices[199],
   },
   {
     /* [884] */
@@ -6968,27 +7005,27 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [888] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[198],
+    /* matcher indices */ &kMatcherIndices[200],
   },
   {
     /* [889] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[193],
+    /* matcher indices */ &kMatcherIndices[195],
   },
   {
     /* [890] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[194],
+    /* matcher indices */ &kMatcherIndices[196],
   },
   {
     /* [891] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[195],
+    /* matcher indices */ &kMatcherIndices[197],
   },
   {
     /* [892] */
     /* usage */ ParameterUsage::kTexture,
-    /* matcher indices */ &kMatcherIndices[196],
+    /* matcher indices */ &kMatcherIndices[198],
   },
   {
     /* [893] */
@@ -7123,7 +7160,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [919] */
     /* usage */ ParameterUsage::kNone,
-    /* matcher indices */ &kMatcherIndices[183],
+    /* matcher indices */ &kMatcherIndices[185],
   },
   {
     /* [920] */
@@ -7183,7 +7220,7 @@ constexpr ParameterInfo kParameters[] = {
   {
     /* [931] */
     /* usage */ ParameterUsage::kNone,
-    /* matcher indices */ &kMatcherIndices[173],
+    /* matcher indices */ &kMatcherIndices[175],
   },
   {
     /* [932] */
@@ -7486,7 +7523,7 @@ constexpr TemplateTypeInfo kTemplateTypes[] = {
   {
     /* [1] */
     /* name */ "U",
-    /* matcher index */ 57,
+    /* matcher index */ 58,
   },
   {
     /* [2] */
@@ -7496,7 +7533,7 @@ constexpr TemplateTypeInfo kTemplateTypes[] = {
   {
     /* [3] */
     /* name */ "U",
-    /* matcher index */ 54,
+    /* matcher index */ 55,
   },
   {
     /* [4] */
@@ -7506,7 +7543,7 @@ constexpr TemplateTypeInfo kTemplateTypes[] = {
   {
     /* [5] */
     /* name */ "U",
-    /* matcher index */ 55,
+    /* matcher index */ 56,
   },
   {
     /* [6] */
@@ -7516,12 +7553,12 @@ constexpr TemplateTypeInfo kTemplateTypes[] = {
   {
     /* [7] */
     /* name */ "U",
-    /* matcher index */ 56,
+    /* matcher index */ 57,
   },
   {
     /* [8] */
     /* name */ "T",
-    /* matcher index */ 48,
+    /* matcher index */ 49,
   },
   {
     /* [9] */
@@ -7531,22 +7568,22 @@ constexpr TemplateTypeInfo kTemplateTypes[] = {
   {
     /* [10] */
     /* name */ "T",
-    /* matcher index */ 53,
+    /* matcher index */ 54,
   },
   {
     /* [11] */
     /* name */ "T",
-    /* matcher index */ 50,
+    /* matcher index */ 51,
   },
   {
     /* [12] */
     /* name */ "T",
-    /* matcher index */ 52,
+    /* matcher index */ 53,
   },
   {
     /* [13] */
     /* name */ "T",
-    /* matcher index */ 51,
+    /* matcher index */ 52,
   },
   {
     /* [14] */
@@ -7556,27 +7593,27 @@ constexpr TemplateTypeInfo kTemplateTypes[] = {
   {
     /* [15] */
     /* name */ "T",
-    /* matcher index */ 57,
+    /* matcher index */ 58,
   },
   {
     /* [16] */
     /* name */ "T",
-    /* matcher index */ 54,
+    /* matcher index */ 55,
   },
   {
     /* [17] */
     /* name */ "T",
-    /* matcher index */ 56,
+    /* matcher index */ 57,
   },
   {
     /* [18] */
     /* name */ "T",
-    /* matcher index */ 55,
+    /* matcher index */ 56,
   },
   {
     /* [19] */
     /* name */ "T",
-    /* matcher index */ 49,
+    /* matcher index */ 50,
   },
 };
 
@@ -9952,7 +9989,7 @@ constexpr OverloadInfo kOverloads[] = {
     /* template types */ &kTemplateTypes[20],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[990],
-    /* return matcher indices */ &kMatcherIndices[173],
+    /* return matcher indices */ &kMatcherIndices[175],
     /* flags */ OverloadFlags(OverloadFlag::kIsConstructor, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
   },
   {
@@ -9963,7 +10000,7 @@ constexpr OverloadInfo kOverloads[] = {
     /* template types */ &kTemplateTypes[9],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[931],
-    /* return matcher indices */ &kMatcherIndices[173],
+    /* return matcher indices */ &kMatcherIndices[175],
     /* flags */ OverloadFlags(OverloadFlag::kIsConstructor, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
   },
   {
@@ -9974,7 +10011,7 @@ constexpr OverloadInfo kOverloads[] = {
     /* template types */ &kTemplateTypes[10],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[928],
-    /* return matcher indices */ &kMatcherIndices[175],
+    /* return matcher indices */ &kMatcherIndices[177],
     /* flags */ OverloadFlags(OverloadFlag::kIsConstructor, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
   },
   {
@@ -9985,7 +10022,7 @@ constexpr OverloadInfo kOverloads[] = {
     /* template types */ &kTemplateTypes[10],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[28],
-    /* return matcher indices */ &kMatcherIndices[175],
+    /* return matcher indices */ &kMatcherIndices[177],
     /* flags */ OverloadFlags(OverloadFlag::kIsConstructor, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
   },
   {
@@ -9996,7 +10033,7 @@ constexpr OverloadInfo kOverloads[] = {
     /* template types */ &kTemplateTypes[10],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[435],
-    /* return matcher indices */ &kMatcherIndices[175],
+    /* return matcher indices */ &kMatcherIndices[177],
     /* flags */ OverloadFlags(OverloadFlag::kIsConstructor, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
   },
   {
@@ -10117,7 +10154,7 @@ constexpr OverloadInfo kOverloads[] = {
     /* template types */ &kTemplateTypes[20],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[990],
-    /* return matcher indices */ &kMatcherIndices[183],
+    /* return matcher indices */ &kMatcherIndices[185],
     /* flags */ OverloadFlags(OverloadFlag::kIsConstructor, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
   },
   {
@@ -10128,7 +10165,7 @@ constexpr OverloadInfo kOverloads[] = {
     /* template types */ &kTemplateTypes[9],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[919],
-    /* return matcher indices */ &kMatcherIndices[183],
+    /* return matcher indices */ &kMatcherIndices[185],
     /* flags */ OverloadFlags(OverloadFlag::kIsConstructor, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
   },
   {
@@ -10843,7 +10880,7 @@ constexpr OverloadInfo kOverloads[] = {
     /* template types */ &kTemplateTypes[20],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[843],
-    /* return matcher indices */ &kMatcherIndices[200],
+    /* return matcher indices */ &kMatcherIndices[202],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
   },
   {
@@ -10865,7 +10902,7 @@ constexpr OverloadInfo kOverloads[] = {
     /* template types */ &kTemplateTypes[20],
     /* template numbers */ &kTemplateNumbers[10],
     /* parameters */ &kParameters[973],
-    /* return matcher indices */ &kMatcherIndices[199],
+    /* return matcher indices */ &kMatcherIndices[201],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
   },
   {
@@ -12526,7 +12563,7 @@ constexpr OverloadInfo kOverloads[] = {
     /* template types */ &kTemplateTypes[11],
     /* template numbers */ &kTemplateNumbers[9],
     /* parameters */ &kParameters[591],
-    /* return matcher indices */ &kMatcherIndices[105],
+    /* return matcher indices */ &kMatcherIndices[173],
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
   },
 };
@@ -13358,7 +13395,7 @@ constexpr IntrinsicInfo kBuiltins[] = {
   },
   {
     /* [106] */
-    /* fn atomicCompareExchangeWeak<T : iu32, S : workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T, T) -> vec2<T> */
+    /* fn atomicCompareExchangeWeak<T : iu32, S : workgroup_or_storage>(ptr<S, atomic<T>, read_write>, T, T) -> __atomic_compare_exchange_result<T> */
     /* num overloads */ 1,
     /* overloads */ &kOverloads[444],
   },
diff --git a/src/tint/transform/decompose_memory_access.cc b/src/tint/transform/decompose_memory_access.cc
index 775cc05f63..a90a6e2711 100644
--- a/src/tint/transform/decompose_memory_access.cc
+++ b/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());
-            auto* func =
-                b.create<ast::Function>(b.Sym(), params, ret_ty, nullptr,
-                                        ast::AttributeList{
-                                            atomic,
-                                            b.Disable(ast::DisabledValidation::kFunctionHasNoBody),
-                                        },
-                                        ast::AttributeList{});
+            const ast::Type* ret_ty = nullptr;
+
+            // For intrinsics that return a struct, there is no AST node for it, so create one now.
+            if (intrinsic->Type() == sem::BuiltinType::kAtomicCompareExchangeWeak) {
+                auto* str = intrinsic->ReturnType()->As<sem::Struct>();
+                TINT_ASSERT(Transform, str && str->Declaration() == nullptr);
+
+                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;
 
diff --git a/src/tint/transform/decompose_memory_access.h b/src/tint/transform/decompose_memory_access.h
index 7a7b783d83..76cb23e2ff 100644
--- a/src/tint/transform/decompose_memory_access.h
+++ b/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;
 
diff --git a/src/tint/transform/decompose_memory_access_test.cc b/src/tint/transform/decompose_memory_access_test.cc
index 22b5da43cb..19f8b2e690 100644
--- a/src/tint/transform/decompose_memory_access_test.cc
+++ b/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_symbol_1(sb, 16u);
-  tint_symbol_2(sb, 16u, 123);
-  tint_symbol_3(sb, 16u, 123);
-  tint_symbol_4(sb, 16u, 123);
-  tint_symbol_5(sb, 16u, 123);
-  tint_symbol_6(sb, 16u, 123);
-  tint_symbol_7(sb, 16u, 123);
-  tint_symbol_8(sb, 16u, 123);
-  tint_symbol_9(sb, 16u, 123);
-  tint_symbol_10(sb, 16u, 123, 345);
-  tint_symbol_11(sb, 20u, 123u);
-  tint_symbol_12(sb, 20u);
-  tint_symbol_13(sb, 20u, 123u);
-  tint_symbol_14(sb, 20u, 123u);
-  tint_symbol_15(sb, 20u, 123u);
-  tint_symbol_16(sb, 20u, 123u);
-  tint_symbol_17(sb, 20u, 123u);
-  tint_symbol_18(sb, 20u, 123u);
-  tint_symbol_19(sb, 20u, 123u);
-  tint_symbol_20(sb, 20u, 123u);
-  tint_symbol_21(sb, 20u, 123u, 345u);
+  tint_atomicStore(sb, 16u, 123);
+  tint_atomicLoad(sb, 16u);
+  tint_atomicAdd(sb, 16u, 123);
+  tint_atomicSub(sb, 16u, 123);
+  tint_atomicMax(sb, 16u, 123);
+  tint_atomicMin(sb, 16u, 123);
+  tint_atomicAnd(sb, 16u, 123);
+  tint_atomicOr(sb, 16u, 123);
+  tint_atomicXor(sb, 16u, 123);
+  tint_atomicExchange(sb, 16u, 123);
+  tint_atomicCompareExchangeWeak(sb, 16u, 123, 345);
+  tint_atomicStore_1(sb, 20u, 123u);
+  tint_atomicLoad_1(sb, 20u);
+  tint_atomicAdd_1(sb, 20u, 123u);
+  tint_atomicSub_1(sb, 20u, 123u);
+  tint_atomicMax_1(sb, 20u, 123u);
+  tint_atomicMin_1(sb, 20u, 123u);
+  tint_atomicAnd_1(sb, 20u, 123u);
+  tint_atomicOr_1(sb, 20u, 123u);
+  tint_atomicXor_1(sb, 20u, 123u);
+  tint_atomicExchange_1(sb, 20u, 123u);
+  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_symbol_1(sb, 16u);
-  tint_symbol_2(sb, 16u, 123);
-  tint_symbol_3(sb, 16u, 123);
-  tint_symbol_4(sb, 16u, 123);
-  tint_symbol_5(sb, 16u, 123);
-  tint_symbol_6(sb, 16u, 123);
-  tint_symbol_7(sb, 16u, 123);
-  tint_symbol_8(sb, 16u, 123);
-  tint_symbol_9(sb, 16u, 123);
-  tint_symbol_10(sb, 16u, 123, 345);
-  tint_symbol_11(sb, 20u, 123u);
-  tint_symbol_12(sb, 20u);
-  tint_symbol_13(sb, 20u, 123u);
-  tint_symbol_14(sb, 20u, 123u);
-  tint_symbol_15(sb, 20u, 123u);
-  tint_symbol_16(sb, 20u, 123u);
-  tint_symbol_17(sb, 20u, 123u);
-  tint_symbol_18(sb, 20u, 123u);
-  tint_symbol_19(sb, 20u, 123u);
-  tint_symbol_20(sb, 20u, 123u);
-  tint_symbol_21(sb, 20u, 123u, 345u);
+  tint_atomicStore(sb, 16u, 123);
+  tint_atomicLoad(sb, 16u);
+  tint_atomicAdd(sb, 16u, 123);
+  tint_atomicSub(sb, 16u, 123);
+  tint_atomicMax(sb, 16u, 123);
+  tint_atomicMin(sb, 16u, 123);
+  tint_atomicAnd(sb, 16u, 123);
+  tint_atomicOr(sb, 16u, 123);
+  tint_atomicXor(sb, 16u, 123);
+  tint_atomicExchange(sb, 16u, 123);
+  tint_atomicCompareExchangeWeak(sb, 16u, 123, 345);
+  tint_atomicStore_1(sb, 20u, 123u);
+  tint_atomicLoad_1(sb, 20u);
+  tint_atomicAdd_1(sb, 20u, 123u);
+  tint_atomicSub_1(sb, 20u, 123u);
+  tint_atomicMax_1(sb, 20u, 123u);
+  tint_atomicMin_1(sb, 20u, 123u);
+  tint_atomicAnd_1(sb, 20u, 123u);
+  tint_atomicOr_1(sb, 20u, 123u);
+  tint_atomicXor_1(sb, 20u, 123u);
+  tint_atomicExchange_1(sb, 20u, 123u);
+  tint_atomicCompareExchangeWeak_1(sb, 20u, 123u, 345u);
 }
 
 @group(0) @binding(0) var<storage, read_write> sb : SB;
diff --git a/src/tint/transform/manager.cc b/src/tint/transform/manager.cc
index 823474c63b..e5f7682ead 100644
--- a/src/tint/transform/manager.cc
+++ b/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()));
diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc
index ccce06a89a..8a48156fe0 100644
--- a/src/tint/writer/glsl/generator_impl.cc
+++ b/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(
-                out, expr, builtin, [&](TextBuffer* b, const std::vector<std::string>& params) {
-                    {
-                        auto pre = line(b);
-                        if (!EmitTypeAndName(pre, builtin->ReturnType(), ast::StorageClass::kNone,
-                                             ast::Access::kUndefined, "result")) {
-                            return false;
-                        }
-                        pre << ";";
+            // 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];
+
+            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;
                     }
-                    {
-                        auto pre = line(b);
-                        pre << "result.x = atomicCompSwap";
-                        {
-                            ScopedParen sp(pre);
-                            pre << params[0];
-                            pre << ", " << params[1];
-                            pre << ", " << params[2];
-                        }
-                        pre << ";";
+                    pre << ", ";
+                    if (!EmitExpression(pre, compare_value)) {
+                        return false;
                     }
-                    {
-                        auto pre = line(b);
-                        pre << "result.y = result.x == " << params[2] << " ? ";
-                        if (TypeOf(expr->args[2])->Is<sem::U32>()) {
-                            pre << "1u : 0u;";
-                        } else {
-                            pre << "1 : 0;";
-                        }
+                    pre << ", ";
+                    if (!EmitExpression(pre, value)) {
+                        return false;
                     }
-                    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:
diff --git a/src/tint/writer/glsl/generator_impl.h b/src/tint/writer/glsl/generator_impl.h
index 72def34c62..1d566ca509 100644
--- a/src/tint/writer/glsl/generator_impl.h
+++ b/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
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index d6a5fa7549..a2cac0f8a9 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/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;
         }
-
-        case Op::kAtomicLoad:
-        case Op::kAtomicStore:
-        case Op::kAtomicAdd:
-        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);
+        default:
+            // Break out to error case below/
+            // Note that atomic intrinsics are generated as functions.
+            break;
     }
 
     TINT_UNREACHABLE(Writer, diagnostics_)
@@ -1465,32 +1459,127 @@ bool GeneratorImpl::EmitStorageBufferAccess(
     return false;
 }
 
-bool GeneratorImpl::EmitStorageAtomicCall(
-    std::ostream& out,
-    const ast::CallExpression* expr,
+bool GeneratorImpl::EmitStorageAtomicIntrinsic(
+    const ast::Function* func,
     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
-    // DecomposeMemoryAccess::Intrinsic call into the corresponding HLSL
-    // atomic intrinsic function.
-    auto generate_helper = [&]() -> std::string {
-        auto rmw = [&](const char* wgsl, const char* hlsl) -> std::string {
-            auto name = UniqueIdentifier(wgsl);
+        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,
+                                 "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, ";
-                if (!EmitTypeAndName(fn, result_ty, ast::StorageClass::kNone,
+                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 "";
+                    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,
-                                     ast::Access::kUndefined, "original_value")) {
-                    return "";
+                if (!EmitTypeAndName(l, value_ty, ast::StorageClass::kNone, ast::Access::kUndefined,
+                                     "ignored")) {
+                    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 (intrinsic->op == Op::kAtomicSub) {
-                    l << "-";
+                if (!EmitTypeAndName(l, result_ty, ast::StorageClass::kNone,
+                                     ast::Access::kUndefined, "result")) {
+                    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) {
-            case Op::kAtomicAdd:
-                return rmw("atomicAdd", "InterlockedAdd");
+            line(&buf) << "buffer.InterlockedCompareExchange(offset, compare, value, "
+                          "result.old_value);";
+            line(&buf) << "result.exchanged = result.old_value == compare;";
+            line(&buf) << "return result;";
 
-            case Op::kAtomicSub:
-                // 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;
+            return true;
         }
-        TINT_UNREACHABLE(Writer, diagnostics_)
-            << "unsupported atomic DecomposeMemoryAccess::Intrinsic::Op: "
-            << static_cast<int>(intrinsic->op);
-        return "";
-    };
-
-    auto func = utils::GetOrCreate(dma_intrinsics_, DMAIntrinsic{intrinsic->op, intrinsic->type},
-                                   generate_helper);
-    if (func.empty()) {
-        return false;
+        default:
+            break;
     }
 
-    out << func;
-    {
-        ScopedParen sp(out);
-        bool first = true;
-        for (auto* arg : expr->args) {
-            if (!first) {
-                out << ", ";
-            }
-            first = false;
-            if (!EmitExpression(out, arg)) {
-                return false;
-            }
-        }
-    }
-
-    return true;
+    TINT_UNREACHABLE(Writer, diagnostics_)
+        << "unsupported atomic DecomposeMemoryAccess::Intrinsic::Op: "
+        << static_cast<int>(intrinsic->op);
+    return false;
 }
 
 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;
-                line() << result << ".y = " << result << ".x == " << compare << ";";
-            }
+            // result.exchanged = result.old_value == 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;
 }
 
diff --git a/src/tint/writer/hlsl/generator_impl.h b/src/tint/writer/hlsl/generator_impl.h
index 86bbd7dedc..0e8ca4c2be 100644
--- a/src/tint/writer/hlsl/generator_impl.h
+++ b/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_;
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index 54d9164823..578e78df8d 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/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) << "bool matched = "
+                line(&buf) << "T old_value = compare;";
+                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;
             });
 
diff --git a/src/tint/writer/spirv/builder.cc b/src/tint/writer/spirv/builder.cc
index e2d1fac831..eb8795f714 100644
--- a/src/tint/writer/spirv/builder.cc
+++ b/src/tint/writer/spirv/builder.cc
@@ -3201,42 +3201,12 @@ bool Builder::GenerateAtomicBuiltin(const sem::Call* call,
                 return false;
             }
 
-            // zero := T(0)
-            // 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)
+            // result := __atomic_compare_exchange_result<T>(original_value, values_equal)
             return push_function_inst(spv::Op::OpCompositeConstruct, {
                                                                          result_type,
                                                                          result_id,
                                                                          original_value,
-                                                                         xchg_success,
+                                                                         values_equal,
                                                                      });
         }
         default:
diff --git a/src/tint/writer/spirv/builder_builtin_test.cc b/src/tint/writer/spirv/builder_builtin_test.cc
index 59a567aae1..734007429c 100644
--- a/src/tint/writer/spirv/builder_builtin_test.cc
+++ b/src/tint/writer/spirv/builder_builtin_test.cc
@@ -2018,15 +2018,15 @@ OpReturn
 
 TEST_F(BuiltinBuilderTest, Call_AtomicCompareExchangeWeak) {
     // struct S {
-    //   u : atomic<u32>;
-    //   i : atomic<i32>;
+    //   u : atomic<u32>,
+    //   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 i : vec2<i32> = atomicCompareExchangeWeak(&b.i, 10);
+    //   let u = atomicCompareExchangeWeak(&b.u, 10u, 20u);
+    //   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 = OpConstant %4 1
-%13 = OpConstant %4 0
-%15 = OpTypePointer StorageBuffer %4
-%17 = OpConstant %4 20
-%18 = OpConstant %4 10
-%19 = OpTypeBool
-%24 = OpTypeVector %5 2
-%26 = OpTypePointer StorageBuffer %5
-%28 = OpConstant %5 20
-%29 = OpConstant %5 10
-%32 = OpConstant %5 0
-%33 = OpConstant %5 1
+%12 = OpTypeBool
+%11 = OpTypeStruct %4 %12
+%13 = OpConstant %4 1
+%14 = OpConstant %4 0
+%16 = OpTypePointer StorageBuffer %4
+%18 = OpConstant %4 20
+%19 = OpConstant %4 10
+%23 = OpTypeStruct %5 %12
+%25 = OpTypePointer StorageBuffer %5
+%27 = OpConstant %5 20
+%28 = OpConstant %5 10
 )";
     auto got_types = DumpInstructions(b.types());
     EXPECT_EQ(expected_types, got_types);
 
-    auto* expected_instructions = R"(%16 = OpAccessChain %15 %1 %13
-%20 = OpAtomicCompareExchange %4 %16 %12 %13 %13 %17 %18
-%21 = OpIEqual %19 %20 %17
-%22 = OpSelect %4 %21 %12 %13
-%10 = OpCompositeConstruct %11 %20 %22
-%27 = OpAccessChain %26 %1 %12
-%30 = OpAtomicCompareExchange %5 %27 %12 %13 %13 %28 %29
-%31 = OpIEqual %19 %30 %28
-%34 = OpSelect %5 %31 %33 %32
-%23 = OpCompositeConstruct %24 %30 %34
+    auto* expected_instructions = R"(%17 = OpAccessChain %16 %1 %14
+%20 = OpAtomicCompareExchange %4 %17 %13 %14 %14 %18 %19
+%21 = OpIEqual %12 %20 %18
+%10 = OpCompositeConstruct %11 %20 %21
+%26 = OpAccessChain %25 %1 %13
+%29 = OpAtomicCompareExchange %5 %26 %13 %14 %14 %27 %28
+%30 = OpIEqual %12 %29 %27
+%22 = OpCompositeConstruct %23 %29 %30
 OpReturn
 )";
     auto got_instructions = DumpInstructions(b.functions()[0].instructions());
diff --git a/test/tint/bug/chromium/1273230.wgsl.expected.hlsl b/test/tint/bug/chromium/1273230.wgsl.expected.hlsl
index 5b6f132b18..05d4fe035b 100644
--- a/test/tint/bug/chromium/1273230.wgsl.expected.hlsl
+++ b/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)]
diff --git a/test/tint/bug/tint/1113.wgsl.expected.hlsl b/test/tint/bug/tint/1113.wgsl.expected.hlsl
index fa8e7e281e..e201d4a83c 100644
--- a/test/tint/bug/tint/1113.wgsl.expected.hlsl
+++ b/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)]
diff --git a/test/tint/bug/tint/1121.wgsl.expected.hlsl b/test/tint/bug/tint/1121.wgsl.expected.hlsl
index a5ab12d742..82a2ef64ea 100644
--- a/test/tint/bug/tint/1121.wgsl.expected.hlsl
+++ b/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;
             }
diff --git a/test/tint/bug/tint/926.wgsl.expected.hlsl b/test/tint/bug/tint/926.wgsl.expected.hlsl
index 9f7e997c6f..f5ed4b6d4a 100644
--- a/test/tint/bug/tint/926.wgsl.expected.hlsl
+++ b/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)]
diff --git a/test/tint/bug/tint/993.wgsl.expected.hlsl b/test/tint/bug/tint/993.wgsl.expected.hlsl
index d1c5d71291..6ec8aed1a5 100644
--- a/test/tint/bug/tint/993.wgsl.expected.hlsl
+++ b/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)]
diff --git a/test/tint/builtins/gen/atomicAdd/8a199a.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicAdd/8a199a.wgsl.expected.hlsl
index 35e6ef35c5..f31fb5cf7f 100644
--- a/test/tint/builtins/gen/atomicAdd/8a199a.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicAdd/d32fe4.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicAdd/d32fe4.wgsl.expected.hlsl
index f87e491c2c..62ae701abd 100644
--- a/test/tint/builtins/gen/atomicAdd/d32fe4.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicAnd/152966.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicAnd/152966.wgsl.expected.hlsl
index fa6c15f837..281ec13b3d 100644
--- a/test/tint/builtins/gen/atomicAnd/152966.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicAnd/85a8d9.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicAnd/85a8d9.wgsl.expected.hlsl
index 9d4eb2f1a1..1c89d149e9 100644
--- a/test/tint/builtins/gen/atomicAnd/85a8d9.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.glsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.glsl
deleted file mode 100644
index 6aa8f5a9f7..0000000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.glsl
+++ /dev/null
@@ -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.
-
-
-
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.hlsl
deleted file mode 100644
index 9bd884cb28..0000000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.hlsl
+++ /dev/null
@@ -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;
-}
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl b/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl
deleted file mode 100644
index a7bb20c970..0000000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl
+++ /dev/null
@@ -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;
-}
-
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl
similarity index 84%
rename from test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl
rename to test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl
index 1f1c2d1e47..f3c62f75cb 100644
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl
+++ b/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_12871c() {
-  var res: vec2<i32> = atomicCompareExchangeWeak(&sb_rw.arg_0, 1, 1);
+// fn atomicCompareExchangeWeak(ptr<storage, atomic<i32>, read_write>, i32, i32) -> __atomic_compare_exchange_result<i32>
+fn atomicCompareExchangeWeak_1bd40a() {
+  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();
 }
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.glsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.glsl
new file mode 100644
index 0000000000..001110fd65
--- /dev/null
+++ b/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;
+}
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.hlsl
new file mode 100644
index 0000000000..b1e30e0e79
--- /dev/null
+++ b/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;
+}
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.msl
new file mode 100644
index 0000000000..6cab275aa3
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.spvasm b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm
similarity index 56%
rename from test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.spvasm
rename to test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.spvasm
index f4b46af60c..454abb1a57 100644
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.spvasm
+++ b/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
-      %int_0 = OpConstant %int 0
-%_ptr_Function_v2int = OpTypePointer Function %v2int
-         %25 = OpConstantNull %v2int
-%atomicCompareExchangeWeak_12871c = OpFunction %void None %5
+%_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32
+         %23 = OpConstantNull %__atomic_compare_exchange_resulti32
+%atomicCompareExchangeWeak_1bd40a = OpFunction %void None %5
           %8 = OpLabel
-        %res = OpVariable %_ptr_Function_v2int Function %25
-         %16 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0
-         %19 = OpAtomicCompareExchange %int %16 %uint_1 %uint_0 %uint_0 %int_1 %int_1
+        %res = OpVariable %_ptr_Function___atomic_compare_exchange_resulti32 Function %23
+         %17 = OpAccessChain %_ptr_StorageBuffer_int %sb_rw %uint_0
+         %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 %v2int %19 %22
+          %9 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %19 %20
                OpStore %res %9
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %5
-         %27 = OpLabel
-         %28 = OpFunctionCall %void %atomicCompareExchangeWeak_12871c
+         %25 = OpLabel
+         %26 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %5
-         %30 = OpLabel
-         %31 = OpFunctionCall %void %atomicCompareExchangeWeak_12871c
+         %28 = OpLabel
+         %29 = OpFunctionCall %void %atomicCompareExchangeWeak_1bd40a
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.wgsl
similarity index 52%
rename from test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.wgsl
rename to test/tint/builtins/gen/atomicCompareExchangeWeak/1bd40a.wgsl.expected.wgsl
index ba2ab03c7e..7e22177963 100644
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.wgsl
+++ b/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() {
-  var res : vec2<i32> = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1, 1);
+fn atomicCompareExchangeWeak_1bd40a() {
+  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();
 }
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl
similarity index 84%
rename from test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl
rename to test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl
index d4473253e9..2fde0d8537 100644
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl
+++ b/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_6673da() {
-  var res: vec2<u32> = atomicCompareExchangeWeak(&sb_rw.arg_0, 1u, 1u);
+// fn atomicCompareExchangeWeak(ptr<storage, atomic<u32>, read_write>, u32, u32) -> __atomic_compare_exchange_result<u32>
+fn atomicCompareExchangeWeak_63d8e6() {
+  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();
 }
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.glsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.glsl
new file mode 100644
index 0000000000..e5738b7cb4
--- /dev/null
+++ b/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;
+}
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.hlsl
new file mode 100644
index 0000000000..9dc72b3c46
--- /dev/null
+++ b/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;
+}
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.msl
new file mode 100644
index 0000000000..94166ca79b
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.spvasm b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm
similarity index 55%
rename from test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.spvasm
rename to test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.spvasm
index 7edaed7b1f..09e83a0ae1 100644
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.spvasm
+++ b/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_v2uint = OpTypePointer Function %v2uint
-         %22 = OpConstantNull %v2uint
-%atomicCompareExchangeWeak_6673da = OpFunction %void None %5
+%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32
+         %21 = OpConstantNull %__atomic_compare_exchange_resultu32
+%atomicCompareExchangeWeak_63d8e6 = OpFunction %void None %5
           %8 = OpLabel
-        %res = OpVariable %_ptr_Function_v2uint Function %22
-         %15 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0
-         %17 = OpAtomicCompareExchange %uint %15 %uint_1 %uint_0 %uint_0 %uint_1 %uint_1
+        %res = OpVariable %_ptr_Function___atomic_compare_exchange_resultu32 Function %21
+         %16 = OpAccessChain %_ptr_StorageBuffer_uint %sb_rw %uint_0
+         %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 %v2uint %17 %19
+          %9 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %17 %18
                OpStore %res %9
                OpReturn
                OpFunctionEnd
 %fragment_main = OpFunction %void None %5
-         %24 = OpLabel
-         %25 = OpFunctionCall %void %atomicCompareExchangeWeak_6673da
+         %23 = OpLabel
+         %24 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
                OpReturn
                OpFunctionEnd
 %compute_main = OpFunction %void None %5
-         %27 = OpLabel
-         %28 = OpFunctionCall %void %atomicCompareExchangeWeak_6673da
+         %26 = OpLabel
+         %27 = OpFunctionCall %void %atomicCompareExchangeWeak_63d8e6
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.wgsl
similarity index 51%
rename from test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.wgsl
rename to test/tint/builtins/gen/atomicCompareExchangeWeak/63d8e6.wgsl.expected.wgsl
index dff8dcad30..3ecac33dbf 100644
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.wgsl
+++ b/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() {
-  var res : vec2<u32> = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1u, 1u);
+fn atomicCompareExchangeWeak_63d8e6() {
+  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();
 }
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.glsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.glsl
deleted file mode 100644
index 65d12ed736..0000000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.glsl
+++ /dev/null
@@ -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.
-
-
-
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.hlsl
deleted file mode 100644
index 430f1323cb..0000000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.hlsl
+++ /dev/null
@@ -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;
-}
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl b/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl
deleted file mode 100644
index b3a827f788..0000000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl
+++ /dev/null
@@ -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;
-}
-
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl
similarity index 85%
rename from test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl
rename to test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl
index d75b725e85..599f9be635 100644
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl
+++ b/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_b2ab2c() {
-  var res: vec2<u32> = atomicCompareExchangeWeak(&arg_0, 1u, 1u);
+// fn atomicCompareExchangeWeak(ptr<workgroup, atomic<u32>, read_write>, u32, u32) -> __atomic_compare_exchange_result<u32>
+fn atomicCompareExchangeWeak_83580d() {
+  var res = atomicCompareExchangeWeak(&arg_0, 1u, 1u);
 }
 
 @stage(compute) @workgroup_size(1)
 fn compute_main() {
-  atomicCompareExchangeWeak_b2ab2c();
+  atomicCompareExchangeWeak_83580d();
 }
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.glsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.glsl
new file mode 100644
index 0000000000..589a98f416
--- /dev/null
+++ b/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;
+}
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.hlsl
similarity index 56%
rename from test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.hlsl
rename to test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.hlsl
index 05cfbc2e74..3f46597640 100644
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.hlsl
+++ b/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() {
-  uint2 atomic_result = uint2(0u, 0u);
+void atomicCompareExchangeWeak_83580d() {
+  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);
-  atomic_result.y = atomic_result.x == atomic_compare_value;
-  uint2 res = atomic_result;
+  InterlockedCompareExchange(arg_0, atomic_compare_value, 1u, atomic_result.old_value);
+  atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
+  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)]
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.msl b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.msl
new file mode 100644
index 0000000000..7a3443b90d
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.spvasm b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm
similarity index 54%
rename from test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.spvasm
rename to test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.spvasm
index 6098500d00..826498596a 100644
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.spvasm
+++ b/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_v2uint = OpTypePointer Function %v2uint
-         %22 = OpConstantNull %v2uint
-         %23 = OpTypeFunction %void %uint
-         %29 = OpConstantNull %uint
+%_ptr_Function___atomic_compare_exchange_resultu32 = OpTypePointer Function %__atomic_compare_exchange_resultu32
+         %21 = OpConstantNull %__atomic_compare_exchange_resultu32
+         %22 = OpTypeFunction %void %uint
+         %28 = 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 %v2uint %17 %19
+         %10 = OpCompositeConstruct %__atomic_compare_exchange_resultu32 %17 %18
                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
-               OpAtomicStore %arg_0 %uint_2 %uint_0 %29
+         %25 = OpLabel
+               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
-         %36 = OpLoad %uint %local_invocation_index_1
-         %35 = OpFunctionCall %void %compute_main_inner %36
+         %33 = OpLabel
+         %35 = OpLoad %uint %local_invocation_index_1
+         %34 = OpFunctionCall %void %compute_main_inner %35
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/83580d.wgsl.expected.wgsl
new file mode 100644
index 0000000000..77a8862a5c
--- /dev/null
+++ b/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();
+}
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.glsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.glsl
deleted file mode 100644
index 2f649cd762..0000000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.glsl
+++ /dev/null
@@ -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.
-
-
-
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl b/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl
deleted file mode 100644
index b002e501d4..0000000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl
+++ /dev/null
@@ -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;
-}
-
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.wgsl
deleted file mode 100644
index 4357511711..0000000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.wgsl
+++ /dev/null
@@ -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();
-}
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.glsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.glsl
deleted file mode 100644
index cbb201a116..0000000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.glsl
+++ /dev/null
@@ -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.
-
-
-
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl b/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl
deleted file mode 100644
index 6a9485892d..0000000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl
+++ /dev/null
@@ -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;
-}
-
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.wgsl
deleted file mode 100644
index d53e099785..0000000000
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.wgsl
+++ /dev/null
@@ -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();
-}
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl
similarity index 85%
rename from test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl
rename to test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl
index c81fc384cd..2147f98f2f 100644
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl
+++ b/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_89ea3b() {
-  var res: vec2<i32> = atomicCompareExchangeWeak(&arg_0, 1, 1);
+// fn atomicCompareExchangeWeak(ptr<workgroup, atomic<i32>, read_write>, i32, i32) -> __atomic_compare_exchange_result<i32>
+fn atomicCompareExchangeWeak_e88938() {
+  var res = atomicCompareExchangeWeak(&arg_0, 1, 1);
 }
 
 @stage(compute) @workgroup_size(1)
 fn compute_main() {
-  atomicCompareExchangeWeak_89ea3b();
+  atomicCompareExchangeWeak_e88938();
 }
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.glsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.glsl
new file mode 100644
index 0000000000..ff5e7a1b5e
--- /dev/null
+++ b/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;
+}
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.hlsl
similarity index 56%
rename from test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.hlsl
rename to test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.hlsl
index 97cc6c4c1e..4d201c49d5 100644
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.hlsl
+++ b/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() {
-  int2 atomic_result = int2(0, 0);
+void atomicCompareExchangeWeak_e88938() {
+  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);
-  atomic_result.y = atomic_result.x == atomic_compare_value;
-  int2 res = atomic_result;
+  InterlockedCompareExchange(arg_0, atomic_compare_value, 1, atomic_result.old_value);
+  atomic_result.exchanged = atomic_result.old_value == atomic_compare_value;
+  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)]
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.msl b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.msl
new file mode 100644
index 0000000000..9815b6d0d6
--- /dev/null
+++ b/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;
+}
+
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.spvasm b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm
similarity index 54%
rename from test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.spvasm
rename to test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.spvasm
index 40f727b99d..a0f338c664 100644
--- a/test/tint/builtins/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.spvasm
+++ b/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
-      %int_0 = OpConstant %int 0
-%_ptr_Function_v2int = OpTypePointer Function %v2int
-         %24 = OpConstantNull %v2int
-         %25 = OpTypeFunction %void %uint
-         %31 = OpConstantNull %int
+%_ptr_Function___atomic_compare_exchange_resulti32 = OpTypePointer Function %__atomic_compare_exchange_resulti32
+         %22 = OpConstantNull %__atomic_compare_exchange_resulti32
+         %23 = OpTypeFunction %void %uint
+         %29 = 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 %v2int %18 %21
+         %11 = OpCompositeConstruct %__atomic_compare_exchange_resulti32 %18 %19
                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
-               OpAtomicStore %arg_0 %uint_2 %uint_0 %31
+         %26 = OpLabel
+               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
-         %38 = OpLoad %uint %local_invocation_index_1
-         %37 = OpFunctionCall %void %compute_main_inner %38
+         %34 = OpLabel
+         %36 = OpLoad %uint %local_invocation_index_1
+         %35 = OpFunctionCall %void %compute_main_inner %36
                OpReturn
                OpFunctionEnd
diff --git a/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.wgsl b/test/tint/builtins/gen/atomicCompareExchangeWeak/e88938.wgsl.expected.wgsl
new file mode 100644
index 0000000000..e882b01620
--- /dev/null
+++ b/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();
+}
diff --git a/test/tint/builtins/gen/atomicExchange/d59712.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicExchange/d59712.wgsl.expected.hlsl
index 506c6fee25..b40e1463d9 100644
--- a/test/tint/builtins/gen/atomicExchange/d59712.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicExchange/f2e22f.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicExchange/f2e22f.wgsl.expected.hlsl
index 1ab0d6cf4d..ea1abc5d03 100644
--- a/test/tint/builtins/gen/atomicExchange/f2e22f.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicLoad/0806ad.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicLoad/0806ad.wgsl.expected.hlsl
index 8c58350d1c..a278b2c16c 100644
--- a/test/tint/builtins/gen/atomicLoad/0806ad.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicLoad/fe6cc3.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicLoad/fe6cc3.wgsl.expected.hlsl
index 967d4c0cad..77edab3db2 100644
--- a/test/tint/builtins/gen/atomicLoad/fe6cc3.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicMax/51b9be.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicMax/51b9be.wgsl.expected.hlsl
index 621944f603..c8f0893154 100644
--- a/test/tint/builtins/gen/atomicMax/51b9be.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicMax/92aa72.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicMax/92aa72.wgsl.expected.hlsl
index 4ac6cd8f3c..f3d398cc33 100644
--- a/test/tint/builtins/gen/atomicMax/92aa72.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicMin/8e38dc.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicMin/8e38dc.wgsl.expected.hlsl
index 5c550171a3..b8d48b8247 100644
--- a/test/tint/builtins/gen/atomicMin/8e38dc.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicMin/c67a74.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicMin/c67a74.wgsl.expected.hlsl
index 4b350c308f..3ae5176288 100644
--- a/test/tint/builtins/gen/atomicMin/c67a74.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicOr/5e95d4.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicOr/5e95d4.wgsl.expected.hlsl
index 21059e60ee..97883cbd8c 100644
--- a/test/tint/builtins/gen/atomicOr/5e95d4.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicOr/8d96a0.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicOr/8d96a0.wgsl.expected.hlsl
index 1504decac9..04be9d7304 100644
--- a/test/tint/builtins/gen/atomicOr/8d96a0.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicStore/cdc29e.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicStore/cdc29e.wgsl.expected.hlsl
index d78cdbec5e..cb8cfddaab 100644
--- a/test/tint/builtins/gen/atomicStore/cdc29e.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicStore/d1e9a6.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicStore/d1e9a6.wgsl.expected.hlsl
index afac632ba9..599f575c7f 100644
--- a/test/tint/builtins/gen/atomicStore/d1e9a6.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicSub/051100.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicSub/051100.wgsl.expected.hlsl
index 26ad745c6d..29d8e0413c 100644
--- a/test/tint/builtins/gen/atomicSub/051100.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicSub/15bfc9.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicSub/15bfc9.wgsl.expected.hlsl
index 21e149f65d..e5d0027b60 100644
--- a/test/tint/builtins/gen/atomicSub/15bfc9.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicXor/54510e.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicXor/54510e.wgsl.expected.hlsl
index 9cc11db1f3..9ed582e64e 100644
--- a/test/tint/builtins/gen/atomicXor/54510e.wgsl.expected.hlsl
+++ b/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() {
diff --git a/test/tint/builtins/gen/atomicXor/c1b78c.wgsl.expected.hlsl b/test/tint/builtins/gen/atomicXor/c1b78c.wgsl.expected.hlsl
index 7d483b70f8..bb38f4211e 100644
--- a/test/tint/builtins/gen/atomicXor/c1b78c.wgsl.expected.hlsl
+++ b/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() {