diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 14a78f9f65..c808b4397e 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -669,6 +669,7 @@ if(${TINT_BUILD_TESTS})
     utils/unique_vector_test.cc
     writer/append_vector_test.cc
     writer/float_to_string_test.cc
+    writer/text_generator_test.cc
   )
 
   if(${TINT_BUILD_SPV_READER})
diff --git a/src/resolver/atomics_validation_test.cc b/src/resolver/atomics_validation_test.cc
index 767a648ebc..4788d457a6 100644
--- a/src/resolver/atomics_validation_test.cc
+++ b/src/resolver/atomics_validation_test.cc
@@ -61,7 +61,7 @@ TEST_F(ResolverAtomicValidationTest, Local) {
 
   EXPECT_FALSE(r()->Resolve());
   EXPECT_EQ(r()->error(),
-            "12:34 error: cannot declare an atomic var in a function scope");
+            "12:34 error: atomic var requires workgroup storage");
 }
 
 TEST_F(ResolverAtomicValidationTest, NoAtomicExpr) {
diff --git a/src/resolver/resolver.cc b/src/resolver/resolver.cc
index c6462df3ed..f613790f57 100644
--- a/src/resolver/resolver.cc
+++ b/src/resolver/resolver.cc
@@ -909,21 +909,13 @@ bool Resolver::ValidateVariable(const VariableInfo* info) {
   // https://gpuweb.github.io/gpuweb/wgsl/#atomic-types
   // Atomic types may only be instantiated by variables in the workgroup storage
   // class or by storage buffer variables with a read_write access mode.
-  if (info->type->UnwrapRef()->Is<sem::Atomic>()) {
-    if (info->kind != VariableKind::kGlobal) {
-      // Neither storage nor workgroup storage classes can be used in function
-      // scopes.
-      AddError("cannot declare an atomic var in a function scope",
-               info->declaration->type()->source());
-      return false;
-    }
-    if (info->storage_class != ast::StorageClass::kWorkgroup) {
-      // Storage buffers require a structure, so just check for workgroup
-      // storage here.
-      AddError("atomic var requires workgroup storage",
-               info->declaration->type()->source());
-      return false;
-    }
+  if (info->type->UnwrapRef()->Is<sem::Atomic>() &&
+      info->storage_class != ast::StorageClass::kWorkgroup) {
+    // Storage buffers require a structure, so just check for workgroup
+    // storage here.
+    AddError("atomic var requires workgroup storage",
+             info->declaration->type()->source());
+    return false;
   }
 
   return true;
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 430f05f94c..9e81adb5da 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -104,8 +104,7 @@ std::ostream& operator<<(std::ostream& s, const RegisterAndSpace& rs) {
 
 }  // namespace
 
-GeneratorImpl::GeneratorImpl(const Program* program)
-    : builder_(ProgramBuilder::Wrap(program)) {}
+GeneratorImpl::GeneratorImpl(const Program* program) : TextGenerator(program) {}
 
 GeneratorImpl::~GeneratorImpl() = default;
 
@@ -165,10 +164,6 @@ bool GeneratorImpl::Generate() {
   return true;
 }
 
-std::string GeneratorImpl::generate_name(const std::string& prefix) {
-  return builder_.Symbols().NameFor(builder_.Symbols().New(prefix));
-}
-
 bool GeneratorImpl::EmitArrayAccessor(std::ostream& out,
                                       ast::ArrayAccessorExpression* expr) {
   if (!EmitExpression(out, expr->array())) {
@@ -222,7 +217,7 @@ bool GeneratorImpl::EmitAssign(ast::AssignmentStatement* stmt) {
 bool GeneratorImpl::EmitBinary(std::ostream& out, ast::BinaryExpression* expr) {
   if (expr->op() == ast::BinaryOp::kLogicalAnd ||
       expr->op() == ast::BinaryOp::kLogicalOr) {
-    auto name = generate_name(kTempNamePrefix);
+    auto name = UniqueIdentifier(kTempNamePrefix);
 
     {
       auto pre = line();
@@ -505,7 +500,7 @@ bool GeneratorImpl::EmitUniformBufferAccess(
     const transform::DecomposeMemoryAccess::Intrinsic* intrinsic) {
   const auto& params = expr->params();
 
-  std::string scalar_offset = generate_name("scalar_offset");
+  std::string scalar_offset = UniqueIdentifier("scalar_offset");
   {
     auto pre = line();
     pre << "const int " << scalar_offset << " = (";
@@ -534,7 +529,7 @@ bool GeneratorImpl::EmitUniformBufferAccess(
       };
       // Has a minimum alignment of 8 bytes, so is either .xy or .zw
       auto load_vec2 = [&] {
-        std::string ubo_load = generate_name("ubo_load");
+        std::string ubo_load = UniqueIdentifier("ubo_load");
 
         {
           auto pre = line();
@@ -744,7 +739,7 @@ bool GeneratorImpl::EmitStorageAtomicCall(
     transform::DecomposeMemoryAccess::Intrinsic::Op op) {
   using Op = transform::DecomposeMemoryAccess::Intrinsic::Op;
 
-  std::string result = generate_name("atomic_result");
+  std::string result = UniqueIdentifier("atomic_result");
 
   auto* result_ty = TypeOf(expr);
   if (!result_ty->Is<sem::Void>()) {
@@ -849,7 +844,7 @@ bool GeneratorImpl::EmitStorageAtomicCall(
       auto* compare_value = expr->params()[2];
       auto* value = expr->params()[3];
 
-      std::string compare = generate_name("atomic_compare_value");
+      std::string compare = UniqueIdentifier("atomic_compare_value");
       {  // T atomic_compare_value = compare_value;
         auto pre = line();
         if (!EmitTypeAndName(pre, TypeOf(compare_value),
@@ -924,7 +919,7 @@ bool GeneratorImpl::EmitStorageAtomicCall(
 bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
                                             ast::CallExpression* expr,
                                             const sem::Intrinsic* intrinsic) {
-  std::string result = generate_name("atomic_result");
+  std::string result = UniqueIdentifier("atomic_result");
 
   if (!intrinsic->ReturnType()->Is<sem::Void>()) {
     auto pre = line();
@@ -1018,7 +1013,7 @@ bool GeneratorImpl::EmitWorkgroupAtomicCall(std::ostream& out,
       auto* compare_value = expr->params()[1];
       auto* value = expr->params()[2];
 
-      std::string compare = generate_name("atomic_compare_value");
+      std::string compare = UniqueIdentifier("atomic_compare_value");
 
       {  // T compare_value = <compare_value>;
         auto pre = line();
@@ -1130,8 +1125,8 @@ bool GeneratorImpl::EmitFrexpCall(std::ostream& out,
 
   // Exponent is an integer, which HLSL does not have an overload for.
   // We need to cast from a float.
-  auto float_exp = generate_name(kTempNamePrefix);
-  auto significand = generate_name(kTempNamePrefix);
+  auto float_exp = UniqueIdentifier(kTempNamePrefix);
+  auto significand = UniqueIdentifier(kTempNamePrefix);
   line() << "float" << width << " " << float_exp << ";";
   {
     auto pre = line();
@@ -1173,8 +1168,8 @@ bool GeneratorImpl::EmitIsNormalCall(std::ostream& out,
   constexpr auto* kMinNormalExponent = "0x0080000";
   constexpr auto* kMaxNormalExponent = "0x7f00000";
 
-  auto exponent = generate_name("tint_isnormal_exponent");
-  auto clamped = generate_name("tint_isnormal_clamped");
+  auto exponent = UniqueIdentifier("tint_isnormal_exponent");
+  auto clamped = UniqueIdentifier("tint_isnormal_clamped");
 
   {
     auto pre = line();
@@ -1196,7 +1191,7 @@ bool GeneratorImpl::EmitDataPackingCall(std::ostream& out,
                                         ast::CallExpression* expr,
                                         const sem::Intrinsic* intrinsic) {
   auto* param = expr->params()[0];
-  auto tmp_name = generate_name(kTempNamePrefix);
+  auto tmp_name = UniqueIdentifier(kTempNamePrefix);
   std::ostringstream expr_out;
   if (!EmitExpression(expr_out, param)) {
     return false;
@@ -1261,7 +1256,7 @@ bool GeneratorImpl::EmitDataUnpackingCall(std::ostream& out,
                                           ast::CallExpression* expr,
                                           const sem::Intrinsic* intrinsic) {
   auto* param = expr->params()[0];
-  auto tmp_name = generate_name(kTempNamePrefix);
+  auto tmp_name = UniqueIdentifier(kTempNamePrefix);
   std::ostringstream expr_out;
   if (!EmitExpression(expr_out, param)) {
     return false;
@@ -1282,7 +1277,7 @@ bool GeneratorImpl::EmitDataUnpackingCall(std::ostream& out,
   switch (intrinsic->Type()) {
     case sem::IntrinsicType::kUnpack4x8snorm:
     case sem::IntrinsicType::kUnpack2x16snorm: {
-      auto tmp_name2 = generate_name(kTempNamePrefix);
+      auto tmp_name2 = UniqueIdentifier(kTempNamePrefix);
       line() << "int " << tmp_name2 << " = int(" << expr_out.str() << ");";
       {  // Perform sign extension on the converted values.
         auto pre = line();
@@ -1302,7 +1297,7 @@ bool GeneratorImpl::EmitDataUnpackingCall(std::ostream& out,
     }
     case sem::IntrinsicType::kUnpack4x8unorm:
     case sem::IntrinsicType::kUnpack2x16unorm: {
-      auto tmp_name2 = generate_name(kTempNamePrefix);
+      auto tmp_name2 = UniqueIdentifier(kTempNamePrefix);
       line() << "uint " << tmp_name2 << " = " << expr_out.str() << ";";
       {
         auto pre = line();
@@ -1492,7 +1487,7 @@ bool GeneratorImpl::EmitTextureCall(std::ostream& out,
       }
 
       // Declare a variable to hold the queried texture info
-      auto dims = generate_name(kTempNamePrefix);
+      auto dims = UniqueIdentifier(kTempNamePrefix);
       if (num_dimensions == 1) {
         line() << "int " << dims << ";";
       } else {
diff --git a/src/writer/hlsl/generator_impl.h b/src/writer/hlsl/generator_impl.h
index 77f2bcfa82..58b791df96 100644
--- a/src/writer/hlsl/generator_impl.h
+++ b/src/writer/hlsl/generator_impl.h
@@ -361,11 +361,6 @@ class GeneratorImpl : public TextGenerator {
       ast::InterpolationType type,
       ast::InterpolationSampling sampling) const;
 
-  /// Generate a unique name
-  /// @param prefix the name prefix
-  /// @returns a unique name
-  std::string generate_name(const std::string& prefix);
-
  private:
   enum class VarType { kIn, kOut };
 
@@ -376,25 +371,6 @@ class GeneratorImpl : public TextGenerator {
 
   std::string get_buffer_name(ast::Expression* expr);
 
-  /// @returns the resolved type of the ast::Expression `expr`
-  /// @param expr the expression
-  sem::Type* TypeOf(ast::Expression* expr) const {
-    return builder_.TypeOf(expr);
-  }
-
-  /// @returns the resolved type of the ast::Type `type`
-  /// @param type the type
-  const sem::Type* TypeOf(const ast::Type* type) const {
-    return builder_.TypeOf(type);
-  }
-
-  /// @returns the resolved type of the ast::TypeDecl `type_decl`
-  /// @param type_decl the type
-  const sem::Type* TypeOf(const ast::TypeDecl* type_decl) const {
-    return builder_.TypeOf(type_decl);
-  }
-
-  ProgramBuilder builder_;
   std::function<bool()> emit_continuing_;
   std::unordered_map<const sem::Struct*, std::string> structure_builders_;
 };
diff --git a/src/writer/hlsl/generator_impl_test.cc b/src/writer/hlsl/generator_impl_test.cc
index 30cf6593bf..cd01730cb4 100644
--- a/src/writer/hlsl/generator_impl_test.cc
+++ b/src/writer/hlsl/generator_impl_test.cc
@@ -43,24 +43,6 @@ TEST_F(HlslGeneratorImplTest, Generate) {
 )");
 }
 
-TEST_F(HlslGeneratorImplTest, InputStructName) {
-  GeneratorImpl& gen = Build();
-
-  ASSERT_EQ(gen.generate_name("func_main_in"), "func_main_in");
-}
-
-TEST_F(HlslGeneratorImplTest, InputStructName_ConflictWithExisting) {
-  Symbols().Register("func_main_out_1");
-  Symbols().Register("func_main_out_2");
-
-  GeneratorImpl& gen = Build();
-
-  ASSERT_EQ(gen.generate_name("func_main_out"), "func_main_out");
-  ASSERT_EQ(gen.generate_name("func_main_out"), "func_main_out_3");
-  ASSERT_EQ(gen.generate_name("func_main_out"), "func_main_out_4");
-  ASSERT_EQ(gen.generate_name("func_main_out"), "func_main_out_5");
-}
-
 struct HlslBuiltinData {
   ast::Builtin builtin;
   const char* attribute_name;
diff --git a/src/writer/msl/generator.cc b/src/writer/msl/generator.cc
index 9f08f561ac..3c2190fc3d 100644
--- a/src/writer/msl/generator.cc
+++ b/src/writer/msl/generator.cc
@@ -13,6 +13,7 @@
 // limitations under the License.
 
 #include "src/writer/msl/generator.h"
+#include "src/writer/msl/generator_impl.h"
 
 namespace tint {
 namespace writer {
diff --git a/src/writer/msl/generator.h b/src/writer/msl/generator.h
index 3844875713..d63f755e79 100644
--- a/src/writer/msl/generator.h
+++ b/src/writer/msl/generator.h
@@ -18,13 +18,14 @@
 #include <memory>
 #include <string>
 
-#include "src/writer/msl/generator_impl.h"
 #include "src/writer/text.h"
 
 namespace tint {
 namespace writer {
 namespace msl {
 
+class GeneratorImpl;
+
 /// Class to generate MSL source
 class Generator : public Text {
  public:
@@ -46,6 +47,9 @@ class Generator : public Text {
   std::string error() const;
 
  private:
+  Generator(const Generator&) = delete;
+  Generator& operator=(const Generator&) = delete;
+
   std::unique_ptr<GeneratorImpl> impl_;
 };
 
diff --git a/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index a12c44d686..bf7ffabc7c 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -33,6 +33,7 @@
 #include "src/ast/variable_decl_statement.h"
 #include "src/ast/void.h"
 #include "src/sem/array.h"
+#include "src/sem/atomic_type.h"
 #include "src/sem/bool_type.h"
 #include "src/sem/call.h"
 #include "src/sem/depth_texture_type.h"
@@ -71,8 +72,7 @@ bool last_is_break_or_fallthrough(const ast::BlockStatement* stmts) {
 
 }  // namespace
 
-GeneratorImpl::GeneratorImpl(const Program* program)
-    : TextGenerator(), program_(program) {}
+GeneratorImpl::GeneratorImpl(const Program* program) : TextGenerator(program) {}
 
 GeneratorImpl::~GeneratorImpl() = default;
 
@@ -359,6 +359,9 @@ bool GeneratorImpl::EmitCall(std::ostream& out, ast::CallExpression* expr) {
 bool GeneratorImpl::EmitIntrinsicCall(std::ostream& out,
                                       ast::CallExpression* expr,
                                       const sem::Intrinsic* intrinsic) {
+  if (intrinsic->IsAtomic()) {
+    return EmitAtomicCall(out, expr, intrinsic);
+  }
   if (intrinsic->IsTexture()) {
     return EmitTextureCall(out, expr, intrinsic);
   }
@@ -422,6 +425,111 @@ bool GeneratorImpl::EmitIntrinsicCall(std::ostream& out,
   return true;
 }
 
+bool GeneratorImpl::EmitAtomicCall(std::ostream& out,
+                                   ast::CallExpression* expr,
+                                   const sem::Intrinsic* intrinsic) {
+  auto call = [&](const char* name) {
+    out << name;
+    {
+      ScopedParen sp(out);
+      for (size_t i = 0; i < expr->params().size(); i++) {
+        auto* arg = expr->params()[i];
+        if (i > 0) {
+          out << ", ";
+        }
+        if (!EmitExpression(out, arg)) {
+          return false;
+        }
+      }
+      out << ", memory_order_relaxed";
+    }
+    return true;
+  };
+
+  switch (intrinsic->Type()) {
+    case sem::IntrinsicType::kAtomicLoad:
+      return call("atomic_load_explicit");
+
+    case sem::IntrinsicType::kAtomicStore:
+      return call("atomic_store_explicit");
+
+    case sem::IntrinsicType::kAtomicAdd:
+      return call("atomic_fetch_add_explicit");
+
+    case sem::IntrinsicType::kAtomicMax:
+      return call("atomic_fetch_max_explicit");
+
+    case sem::IntrinsicType::kAtomicMin:
+      return call("atomic_fetch_min_explicit");
+
+    case sem::IntrinsicType::kAtomicAnd:
+      return call("atomic_fetch_and_explicit");
+
+    case sem::IntrinsicType::kAtomicOr:
+      return call("atomic_fetch_or_explicit");
+
+    case sem::IntrinsicType::kAtomicXor:
+      return call("atomic_fetch_xor_explicit");
+
+    case sem::IntrinsicType::kAtomicExchange:
+      return call("atomic_exchange_explicit");
+
+    case sem::IntrinsicType::kAtomicCompareExchangeWeak: {
+      auto* target = expr->params()[0];
+      auto* compare_value = expr->params()[1];
+      auto* value = expr->params()[2];
+
+      auto prev_value = UniqueIdentifier("prev_value");
+      auto matched = UniqueIdentifier("matched");
+
+      {  // prev_value = <compare_value>;
+        auto pre = line();
+        if (!EmitType(pre, TypeOf(value), "")) {
+          return false;
+        }
+        pre << " " << prev_value << " = ";
+        if (!EmitExpression(pre, compare_value)) {
+          return false;
+        }
+        pre << ";";
+      }
+
+      {  // bool matched = atomic_compare_exchange_weak_explicit(
+         //   target, &got, <value>, memory_order_relaxed, memory_order_relaxed)
+        auto pre = line();
+        pre << "bool " << matched << " = atomic_compare_exchange_weak_explicit";
+        {
+          ScopedParen sp(pre);
+          if (!EmitExpression(pre, target)) {
+            return false;
+          }
+          pre << ", &" << prev_value << ", ";
+          if (!EmitExpression(pre, value)) {
+            return false;
+          }
+          pre << ", memory_order_relaxed, memory_order_relaxed";
+        }
+        pre << ";";
+      }
+
+      {  // [u]int2(got, matched)
+        if (!EmitType(out, TypeOf(expr), "")) {
+          return false;
+        }
+        out << "(" << prev_value << ", " << matched << ")";
+      }
+      return true;
+    }
+
+    default:
+      break;
+  }
+
+  TINT_UNREACHABLE(Writer, diagnostics_)
+      << "unsupported atomic intrinsic: " << intrinsic->Type();
+  return false;
+}
+
 bool GeneratorImpl::EmitTextureCall(std::ostream& out,
                                     ast::CallExpression* expr,
                                     const sem::Intrinsic* intrinsic) {
@@ -1550,6 +1658,20 @@ bool GeneratorImpl::EmitSwitch(ast::SwitchStatement* stmt) {
 bool GeneratorImpl::EmitType(std::ostream& out,
                              const sem::Type* type,
                              const std::string& name) {
+  if (auto* atomic = type->As<sem::Atomic>()) {
+    if (atomic->Type()->Is<sem::I32>()) {
+      out << "atomic_int";
+      return true;
+    }
+    if (atomic->Type()->Is<sem::U32>()) {
+      out << "atomic_uint";
+      return true;
+    }
+    TINT_ICE(Writer, diagnostics_)
+        << "unhandled atomic type " << atomic->Type()->type_name();
+    return false;
+  }
+
   if (auto* ary = type->As<sem::Array>()) {
     const sem::Type* base_type = ary;
     std::vector<uint32_t> sizes;
@@ -1570,18 +1692,33 @@ bool GeneratorImpl::EmitType(std::ostream& out,
     for (uint32_t size : sizes) {
       out << "[" << size << "]";
     }
-  } else if (type->Is<sem::Bool>()) {
+    return true;
+  }
+
+  if (type->Is<sem::Bool>()) {
     out << "bool";
-  } else if (type->Is<sem::F32>()) {
+    return true;
+  }
+
+  if (type->Is<sem::F32>()) {
     out << "float";
-  } else if (type->Is<sem::I32>()) {
+    return true;
+  }
+
+  if (type->Is<sem::I32>()) {
     out << "int";
-  } else if (auto* mat = type->As<sem::Matrix>()) {
+    return true;
+  }
+
+  if (auto* mat = type->As<sem::Matrix>()) {
     if (!EmitType(out, mat->type(), "")) {
       return false;
     }
     out << mat->columns() << "x" << mat->rows();
-  } else if (auto* ptr = type->As<sem::Pointer>()) {
+    return true;
+  }
+
+  if (auto* ptr = type->As<sem::Pointer>()) {
     switch (ptr->StorageClass()) {
       case ast::StorageClass::kFunction:
       case ast::StorageClass::kPrivate:
@@ -1611,13 +1748,22 @@ bool GeneratorImpl::EmitType(std::ostream& out,
       }
       out << "* " << name;
     }
-  } else if (type->Is<sem::Sampler>()) {
+    return true;
+  }
+
+  if (type->Is<sem::Sampler>()) {
     out << "sampler";
-  } else if (auto* str = type->As<sem::Struct>()) {
+    return true;
+  }
+
+  if (auto* str = type->As<sem::Struct>()) {
     // The struct type emits as just the name. The declaration would be emitted
     // as part of emitting the declared types.
     out << program_->Symbols().NameFor(str->Declaration()->name());
-  } else if (auto* tex = type->As<sem::Texture>()) {
+    return true;
+  }
+
+  if (auto* tex = type->As<sem::Texture>()) {
     if (tex->Is<sem::DepthTexture>()) {
       out << "depth";
     } else {
@@ -1684,23 +1830,30 @@ bool GeneratorImpl::EmitType(std::ostream& out,
       return false;
     }
     out << ">";
+    return true;
+  }
 
-  } else if (type->Is<sem::U32>()) {
+  if (type->Is<sem::U32>()) {
     out << "uint";
-  } else if (auto* vec = type->As<sem::Vector>()) {
+    return true;
+  }
+
+  if (auto* vec = type->As<sem::Vector>()) {
     if (!EmitType(out, vec->type(), "")) {
       return false;
     }
     out << vec->size();
-  } else if (type->Is<sem::Void>()) {
-    out << "void";
-  } else {
-    diagnostics_.add_error(diag::System::Writer,
-                           "unknown type in EmitType: " + type->type_name());
-    return false;
+    return true;
   }
 
-  return true;
+  if (type->Is<sem::Void>()) {
+    out << "void";
+    return true;
+  }
+
+  diagnostics_.add_error(diag::System::Writer,
+                         "unknown type in EmitType: " + type->type_name());
+  return false;
 }
 
 bool GeneratorImpl::EmitPackedType(std::ostream& out,
@@ -2039,6 +2192,10 @@ GeneratorImpl::SizeAndAlign GeneratorImpl::MslPackedTypeSizeAndAlign(
     return SizeAndAlign{str->Size(), str->Align()};
   }
 
+  if (auto* atomic = ty->As<sem::Atomic>()) {
+    return MslPackedTypeSizeAndAlign(atomic->Type());
+  }
+
   TINT_UNREACHABLE(Writer, diagnostics_)
       << "Unhandled type " << ty->TypeInfo().name;
   return {};
diff --git a/src/writer/msl/generator_impl.h b/src/writer/msl/generator_impl.h
index a527eb948e..0f5bf1a7e3 100644
--- a/src/writer/msl/generator_impl.h
+++ b/src/writer/msl/generator_impl.h
@@ -104,6 +104,15 @@ class GeneratorImpl : public TextGenerator {
   bool EmitIntrinsicCall(std::ostream& out,
                          ast::CallExpression* expr,
                          const sem::Intrinsic* intrinsic);
+  /// Handles generating a call to an atomic function (`atomicAdd`,
+  /// `atomicMax`, etc)
+  /// @param out the output of the expression stream
+  /// @param expr the call expression
+  /// @param intrinsic the semantic information for the atomic intrinsic
+  /// @returns true if the call expression is emitted
+  bool EmitAtomicCall(std::ostream& out,
+                      ast::CallExpression* expr,
+                      const sem::Intrinsic* intrinsic);
   /// Handles generating a call to a texture function (`textureSample`,
   /// `textureSampleGrad`, etc)
   /// @param out the output of the expression stream
@@ -263,24 +272,6 @@ class GeneratorImpl : public TextGenerator {
       ast::InterpolationSampling sampling) const;
 
  private:
-  /// @returns the resolved type of the ast::Expression `expr`
-  /// @param expr the expression
-  sem::Type* TypeOf(ast::Expression* expr) const {
-    return program_->TypeOf(expr);
-  }
-
-  /// @returns the resolved type of the ast::Type `type`
-  /// @param type the type
-  const sem::Type* TypeOf(const ast::Type* type) const {
-    return program_->TypeOf(type);
-  }
-
-  /// @returns the resolved type of the ast::TypeDecl `type_decl`
-  /// @param type_decl the type declaration
-  const sem::Type* TypeOf(const ast::TypeDecl* type_decl) const {
-    return program_->TypeOf(type_decl);
-  }
-
   // A pair of byte size and alignment `uint32_t`s.
   struct SizeAndAlign {
     uint32_t size;
@@ -291,7 +282,6 @@ class GeneratorImpl : public TextGenerator {
   /// type.
   SizeAndAlign MslPackedTypeSizeAndAlign(const sem::Type* ty);
 
-  const Program* program_ = nullptr;
   std::function<bool()> emit_continuing_;
 };
 
diff --git a/src/writer/text_generator.cc b/src/writer/text_generator.cc
index a3ec05c5f6..6a6c36423f 100644
--- a/src/writer/text_generator.cc
+++ b/src/writer/text_generator.cc
@@ -14,10 +14,13 @@
 
 #include "src/writer/text_generator.h"
 
+#include <limits>
+
 namespace tint {
 namespace writer {
 
-TextGenerator::TextGenerator() = default;
+TextGenerator::TextGenerator(const Program* program)
+    : program_(program), builder_(ProgramBuilder::Wrap(program)) {}
 
 TextGenerator::~TextGenerator() = default;
 
@@ -31,6 +34,10 @@ void TextGenerator::make_indent(std::ostream& out) const {
   }
 }
 
+std::string TextGenerator::UniqueIdentifier(const std::string& prefix) {
+  return builder_.Symbols().NameFor(builder_.Symbols().New(prefix));
+}
+
 TextGenerator::LineWriter::LineWriter(TextGenerator* generator)
     : gen(generator) {}
 
diff --git a/src/writer/text_generator.h b/src/writer/text_generator.h
index 8f68385afb..57011d9fa1 100644
--- a/src/writer/text_generator.h
+++ b/src/writer/text_generator.h
@@ -20,6 +20,7 @@
 #include <utility>
 
 #include "src/diagnostic/diagnostic.h"
+#include "src/program_builder.h"
 
 namespace tint {
 namespace writer {
@@ -28,7 +29,8 @@ namespace writer {
 class TextGenerator {
  public:
   /// Constructor
-  TextGenerator();
+  /// @param program the program used by the generator
+  explicit TextGenerator(const Program* program);
   ~TextGenerator();
 
   /// Increment the emitter indent level
@@ -58,6 +60,11 @@ class TextGenerator {
   /// @returns the error
   std::string error() const { return diagnostics_.str(); }
 
+  /// @return a new, unique identifier with the given prefix.
+  /// @param prefix optional prefix to apply to the generated identifier. If
+  /// empty "tint" will be used.
+  std::string UniqueIdentifier(const std::string& prefix = "");
+
  protected:
   /// LineWriter is a helper that acts as a string buffer, who's content is
   /// emitted to the TextGenerator as a single line on destruction.
@@ -122,9 +129,31 @@ class TextGenerator {
     TextGenerator* gen;
   };
 
+  /// @returns the resolved type of the ast::Expression `expr`
+  /// @param expr the expression
+  sem::Type* TypeOf(ast::Expression* expr) const {
+    return builder_.TypeOf(expr);
+  }
+
+  /// @returns the resolved type of the ast::Type `type`
+  /// @param type the type
+  const sem::Type* TypeOf(const ast::Type* type) const {
+    return builder_.TypeOf(type);
+  }
+
+  /// @returns the resolved type of the ast::TypeDecl `type_decl`
+  /// @param type_decl the type
+  const sem::Type* TypeOf(const ast::TypeDecl* type_decl) const {
+    return builder_.TypeOf(type_decl);
+  }
+
   /// @returns a new LineWriter, used for buffering and writing a line to out_
   LineWriter line() { return LineWriter(this); }
 
+  /// The program
+  Program const* const program_;
+  /// A ProgramBuilder that thinly wraps program_
+  ProgramBuilder builder_;
   /// The text output stream
   std::ostringstream out_;
   /// Diagnostics generated by the generator
diff --git a/src/writer/text_generator_test.cc b/src/writer/text_generator_test.cc
new file mode 100644
index 0000000000..10c70860a5
--- /dev/null
+++ b/src/writer/text_generator_test.cc
@@ -0,0 +1,48 @@
+// Copyright 2021 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/writer/text_generator.h"
+
+#include "gtest/gtest.h"
+
+namespace tint {
+namespace writer {
+namespace {
+
+TEST(TextGeneratorTest, UniqueIdentifier) {
+  Program program(ProgramBuilder{});
+
+  TextGenerator gen(&program);
+
+  ASSERT_EQ(gen.UniqueIdentifier("ident"), "ident");
+  ASSERT_EQ(gen.UniqueIdentifier("ident"), "ident_1");
+}
+
+TEST(TextGeneratorTest, UniqueIdentifier_ConflictWithExisting) {
+  ProgramBuilder builder;
+  builder.Symbols().Register("ident_1");
+  builder.Symbols().Register("ident_2");
+  Program program(std::move(builder));
+
+  TextGenerator gen(&program);
+
+  ASSERT_EQ(gen.UniqueIdentifier("ident"), "ident");
+  ASSERT_EQ(gen.UniqueIdentifier("ident"), "ident_3");
+  ASSERT_EQ(gen.UniqueIdentifier("ident"), "ident_4");
+  ASSERT_EQ(gen.UniqueIdentifier("ident"), "ident_5");
+}
+
+}  // namespace
+}  // namespace writer
+}  // namespace tint
diff --git a/src/writer/wgsl/generator.cc b/src/writer/wgsl/generator.cc
index f19046375c..36eeabc58d 100644
--- a/src/writer/wgsl/generator.cc
+++ b/src/writer/wgsl/generator.cc
@@ -13,6 +13,7 @@
 // limitations under the License.
 
 #include "src/writer/wgsl/generator.h"
+#include "src/writer/wgsl/generator_impl.h"
 
 namespace tint {
 namespace writer {
diff --git a/src/writer/wgsl/generator.h b/src/writer/wgsl/generator.h
index 6085ecf099..6b0b358905 100644
--- a/src/writer/wgsl/generator.h
+++ b/src/writer/wgsl/generator.h
@@ -19,12 +19,13 @@
 #include <string>
 
 #include "src/writer/text.h"
-#include "src/writer/wgsl/generator_impl.h"
 
 namespace tint {
 namespace writer {
 namespace wgsl {
 
+class GeneratorImpl;
+
 /// Class to generate WGSL source
 class Generator : public Text {
  public:
@@ -46,6 +47,9 @@ class Generator : public Text {
   std::string error() const;
 
  private:
+  Generator(const Generator&) = delete;
+  Generator& operator=(const Generator&) = delete;
+
   std::unique_ptr<GeneratorImpl> impl_;
 };
 
diff --git a/src/writer/wgsl/generator_impl.cc b/src/writer/wgsl/generator_impl.cc
index 42785a827b..077e990e51 100644
--- a/src/writer/wgsl/generator_impl.cc
+++ b/src/writer/wgsl/generator_impl.cc
@@ -15,7 +15,6 @@
 #include "src/writer/wgsl/generator_impl.h"
 
 #include <algorithm>
-#include <limits>
 
 #include "src/ast/access.h"
 #include "src/ast/alias.h"
@@ -60,8 +59,7 @@ namespace tint {
 namespace writer {
 namespace wgsl {
 
-GeneratorImpl::GeneratorImpl(const Program* program)
-    : TextGenerator(), program_(program) {}
+GeneratorImpl::GeneratorImpl(const Program* program) : TextGenerator(program) {}
 
 GeneratorImpl::~GeneratorImpl() = default;
 
@@ -1059,24 +1057,6 @@ bool GeneratorImpl::EmitSwitch(ast::SwitchStatement* stmt) {
   return true;
 }
 
-std::string GeneratorImpl::UniqueIdentifier(const std::string& suffix) {
-  auto const limit =
-      std::numeric_limits<decltype(next_unique_identifier_suffix)>::max();
-  while (next_unique_identifier_suffix < limit) {
-    auto ident = "tint_" + std::to_string(next_unique_identifier_suffix);
-    if (!suffix.empty()) {
-      ident += "_" + suffix;
-    }
-    next_unique_identifier_suffix++;
-    if (!program_->Symbols().Get(ident).IsValid()) {
-      return ident;
-    }
-  }
-  diagnostics_.add_error(diag::System::Writer,
-                         "Unable to generate a unique WGSL identifier");
-  return "<invalid-ident>";
-}
-
 }  // namespace wgsl
 }  // namespace writer
 }  // namespace tint
diff --git a/src/writer/wgsl/generator_impl.h b/src/writer/wgsl/generator_impl.h
index ad9f535abc..7d7c732152 100644
--- a/src/writer/wgsl/generator_impl.h
+++ b/src/writer/wgsl/generator_impl.h
@@ -194,13 +194,6 @@ class GeneratorImpl : public TextGenerator {
   /// @param decos the decoration list
   /// @returns true if the decorations were emitted
   bool EmitDecorations(const ast::DecorationList& decos);
-
- private:
-  /// @return a new, unique, valid WGSL identifier with the given suffix.
-  std::string UniqueIdentifier(const std::string& suffix = "");
-
-  Program const* const program_;
-  uint32_t next_unique_identifier_suffix = 0;
 };
 
 }  // namespace wgsl
diff --git a/src/writer/wgsl/generator_impl_type_test.cc b/src/writer/wgsl/generator_impl_type_test.cc
index 991dc6991d..3273e9206a 100644
--- a/src/writer/wgsl/generator_impl_type_test.cc
+++ b/src/writer/wgsl/generator_impl_type_test.cc
@@ -141,10 +141,10 @@ TEST_F(WgslGeneratorImplTest, EmitType_StructOffsetDecl) {
   ASSERT_TRUE(gen.EmitStructType(s)) << gen.error();
   EXPECT_EQ(gen.result(), R"(struct S {
   [[size(8)]]
-  tint_0_padding : u32;
+  padding : u32;
   a : i32;
   [[size(4)]]
-  tint_1_padding : u32;
+  padding_1 : u32;
   b : f32;
 };
 )");
@@ -162,10 +162,10 @@ TEST_F(WgslGeneratorImplTest, EmitType_StructOffsetDecl_WithSymbolCollisions) {
   ASSERT_TRUE(gen.EmitStructType(s)) << gen.error();
   EXPECT_EQ(gen.result(), R"(struct S {
   [[size(8)]]
-  tint_1_padding : u32;
+  padding : u32;
   tint_0_padding : i32;
   [[size(4)]]
-  tint_3_padding : u32;
+  padding_1 : u32;
   tint_2_padding : f32;
 };
 )");
diff --git a/test/BUILD.gn b/test/BUILD.gn
index e120d8bc63..fba82209de 100644
--- a/test/BUILD.gn
+++ b/test/BUILD.gn
@@ -305,6 +305,7 @@ tint_unittests_source_set("tint_unittests_core_src") {
     "../src/utils/unique_vector_test.cc",
     "../src/writer/append_vector_test.cc",
     "../src/writer/float_to_string_test.cc",
+    "../src/writer/text_generator_test.cc",
   ]
 
   deps = [
diff --git a/test/intrinsics/gen/atomicAdd/794055.wgsl.expected.msl b/test/intrinsics/gen/atomicAdd/794055.wgsl.expected.msl
index 9cf781e6e0..825608bc58 100644
--- a/test/intrinsics/gen/atomicAdd/794055.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicAdd/794055.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicAdd_794055(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : i32 = atomicAdd(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicAdd_794055(threadgroup atomic_int* const tint_symbol_1) {
+  int res = atomic_fetch_add_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicAdd_794055(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicAdd_794055(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicAdd/8a199a.wgsl.expected.msl b/test/intrinsics/gen/atomicAdd/8a199a.wgsl.expected.msl
index 86b5492fb0..4d1bb52ea2 100644
--- a/test/intrinsics/gen/atomicAdd/8a199a.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicAdd/8a199a.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicAdd_8a199a() {
-  var res : u32 = atomicAdd(&(sb_rw.arg_0), 1u);
+void atomicAdd_8a199a(device SB_RW& sb_rw) {
+  uint res = atomic_fetch_add_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicAdd_8a199a();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicAdd_8a199a(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicAdd_8a199a();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicAdd_8a199a(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicAdd/d32fe4.wgsl.expected.msl b/test/intrinsics/gen/atomicAdd/d32fe4.wgsl.expected.msl
index 85b0ee5491..a1e838a57e 100644
--- a/test/intrinsics/gen/atomicAdd/d32fe4.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicAdd/d32fe4.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicAdd_d32fe4() {
-  var res : i32 = atomicAdd(&(sb_rw.arg_0), 1);
+void atomicAdd_d32fe4(device SB_RW& sb_rw) {
+  int res = atomic_fetch_add_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicAdd_d32fe4();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicAdd_d32fe4(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicAdd_d32fe4();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicAdd_d32fe4(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicAdd/d5db1d.wgsl.expected.msl b/test/intrinsics/gen/atomicAdd/d5db1d.wgsl.expected.msl
index 3e13cfdcc1..ee515459e2 100644
--- a/test/intrinsics/gen/atomicAdd/d5db1d.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicAdd/d5db1d.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicAdd_d5db1d(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : u32 = atomicAdd(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicAdd_d5db1d(threadgroup atomic_uint* const tint_symbol_1) {
+  uint res = atomic_fetch_add_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicAdd_d5db1d(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicAdd_d5db1d(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicAnd/152966.wgsl.expected.msl b/test/intrinsics/gen/atomicAnd/152966.wgsl.expected.msl
index efcd326bcd..d095ebb5e2 100644
--- a/test/intrinsics/gen/atomicAnd/152966.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicAnd/152966.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicAnd_152966() {
-  var res : i32 = atomicAnd(&(sb_rw.arg_0), 1);
+void atomicAnd_152966(device SB_RW& sb_rw) {
+  int res = atomic_fetch_and_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicAnd_152966();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicAnd_152966(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicAnd_152966();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicAnd_152966(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicAnd/34edd3.wgsl.expected.msl b/test/intrinsics/gen/atomicAnd/34edd3.wgsl.expected.msl
index 6a5bb56aa7..dd88dd39d0 100644
--- a/test/intrinsics/gen/atomicAnd/34edd3.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicAnd/34edd3.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicAnd_34edd3(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : u32 = atomicAnd(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicAnd_34edd3(threadgroup atomic_uint* const tint_symbol_1) {
+  uint res = atomic_fetch_and_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicAnd_34edd3(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicAnd_34edd3(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicAnd/45a819.wgsl.expected.msl b/test/intrinsics/gen/atomicAnd/45a819.wgsl.expected.msl
index 47755f6231..7a700e9a90 100644
--- a/test/intrinsics/gen/atomicAnd/45a819.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicAnd/45a819.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicAnd_45a819(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : i32 = atomicAnd(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicAnd_45a819(threadgroup atomic_int* const tint_symbol_1) {
+  int res = atomic_fetch_and_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicAnd_45a819(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicAnd_45a819(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicAnd/85a8d9.wgsl.expected.msl b/test/intrinsics/gen/atomicAnd/85a8d9.wgsl.expected.msl
index 31f25e8d54..b9ca1b0b6e 100644
--- a/test/intrinsics/gen/atomicAnd/85a8d9.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicAnd/85a8d9.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicAnd_85a8d9() {
-  var res : u32 = atomicAnd(&(sb_rw.arg_0), 1u);
+void atomicAnd_85a8d9(device SB_RW& sb_rw) {
+  uint res = atomic_fetch_and_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicAnd_85a8d9();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicAnd_85a8d9(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicAnd_85a8d9();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicAnd_85a8d9(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl b/test/intrinsics/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl
index 7135dea0c4..56d8090e48 100644
--- a/test/intrinsics/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicCompareExchangeWeak/12871c.wgsl.expected.msl
@@ -1,25 +1,23 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicCompareExchangeWeak_12871c() {
-  var res : vec2<i32> = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1, 1);
+void atomicCompareExchangeWeak_12871c(device SB_RW& sb_rw) {
+  int prev_value = 1;
+  bool matched = atomic_compare_exchange_weak_explicit(&(sb_rw.arg_0), &prev_value, 1, memory_order_relaxed, memory_order_relaxed);
+  int2 res = int2(prev_value, matched);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicCompareExchangeWeak_12871c();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicCompareExchangeWeak_12871c(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicCompareExchangeWeak_12871c();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicCompareExchangeWeak_12871c(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl b/test/intrinsics/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl
index a6c99cee60..c4999ee58d 100644
--- a/test/intrinsics/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicCompareExchangeWeak/6673da.wgsl.expected.msl
@@ -1,25 +1,23 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicCompareExchangeWeak_6673da() {
-  var res : vec2<u32> = atomicCompareExchangeWeak(&(sb_rw.arg_0), 1u, 1u);
+void atomicCompareExchangeWeak_6673da(device SB_RW& sb_rw) {
+  uint prev_value = 1u;
+  bool matched = atomic_compare_exchange_weak_explicit(&(sb_rw.arg_0), &prev_value, 1u, memory_order_relaxed, memory_order_relaxed);
+  uint2 res = uint2(prev_value, matched);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicCompareExchangeWeak_6673da();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicCompareExchangeWeak_6673da(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicCompareExchangeWeak_6673da();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicCompareExchangeWeak_6673da(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl b/test/intrinsics/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl
index c6431fb097..036e9bf29f 100644
--- a/test/intrinsics/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicCompareExchangeWeak/89ea3b.wgsl.expected.msl
@@ -1,14 +1,19 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicCompareExchangeWeak_89ea3b(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : vec2<i32> = atomicCompareExchangeWeak(&(*(tint_symbol)), 1, 1);
+using namespace metal;
+void atomicCompareExchangeWeak_89ea3b(threadgroup atomic_int* const tint_symbol_1) {
+  int prev_value = 1;
+  bool matched = atomic_compare_exchange_weak_explicit(&(*(tint_symbol_1)), &prev_value, 1, memory_order_relaxed, memory_order_relaxed);
+  int2 res = int2(prev_value, matched);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicCompareExchangeWeak_89ea3b(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicCompareExchangeWeak_89ea3b(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl b/test/intrinsics/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl
index 1b632df073..45d921cb00 100644
--- a/test/intrinsics/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicCompareExchangeWeak/b2ab2c.wgsl.expected.msl
@@ -1,14 +1,19 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicCompareExchangeWeak_b2ab2c(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : vec2<u32> = atomicCompareExchangeWeak(&(*(tint_symbol)), 1u, 1u);
+using namespace metal;
+void atomicCompareExchangeWeak_b2ab2c(threadgroup atomic_uint* const tint_symbol_1) {
+  uint prev_value = 1u;
+  bool matched = atomic_compare_exchange_weak_explicit(&(*(tint_symbol_1)), &prev_value, 1u, memory_order_relaxed, memory_order_relaxed);
+  uint2 res = uint2(prev_value, matched);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicCompareExchangeWeak_b2ab2c(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicCompareExchangeWeak_b2ab2c(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicExchange/0a5dca.wgsl.expected.msl b/test/intrinsics/gen/atomicExchange/0a5dca.wgsl.expected.msl
index 070becac55..1dafa27f9f 100644
--- a/test/intrinsics/gen/atomicExchange/0a5dca.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicExchange/0a5dca.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicExchange_0a5dca(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : u32 = atomicExchange(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicExchange_0a5dca(threadgroup atomic_uint* const tint_symbol_1) {
+  uint res = atomic_exchange_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicExchange_0a5dca(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicExchange_0a5dca(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicExchange/d59712.wgsl.expected.msl b/test/intrinsics/gen/atomicExchange/d59712.wgsl.expected.msl
index d3d01f5054..248b62acc9 100644
--- a/test/intrinsics/gen/atomicExchange/d59712.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicExchange/d59712.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicExchange_d59712() {
-  var res : u32 = atomicExchange(&(sb_rw.arg_0), 1u);
+void atomicExchange_d59712(device SB_RW& sb_rw) {
+  uint res = atomic_exchange_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicExchange_d59712();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicExchange_d59712(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicExchange_d59712();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicExchange_d59712(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicExchange/e114ba.wgsl.expected.msl b/test/intrinsics/gen/atomicExchange/e114ba.wgsl.expected.msl
index 9e1b24d361..5104612c12 100644
--- a/test/intrinsics/gen/atomicExchange/e114ba.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicExchange/e114ba.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicExchange_e114ba(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : i32 = atomicExchange(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicExchange_e114ba(threadgroup atomic_int* const tint_symbol_1) {
+  int res = atomic_exchange_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicExchange_e114ba(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicExchange_e114ba(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicExchange/f2e22f.wgsl.expected.msl b/test/intrinsics/gen/atomicExchange/f2e22f.wgsl.expected.msl
index be0f8972fd..9e0dd6b081 100644
--- a/test/intrinsics/gen/atomicExchange/f2e22f.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicExchange/f2e22f.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicExchange_f2e22f() {
-  var res : i32 = atomicExchange(&(sb_rw.arg_0), 1);
+void atomicExchange_f2e22f(device SB_RW& sb_rw) {
+  int res = atomic_exchange_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicExchange_f2e22f();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicExchange_f2e22f(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicExchange_f2e22f();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicExchange_f2e22f(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicLoad/0806ad.wgsl.expected.msl b/test/intrinsics/gen/atomicLoad/0806ad.wgsl.expected.msl
index 39eff932e4..8cc7837b71 100644
--- a/test/intrinsics/gen/atomicLoad/0806ad.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicLoad/0806ad.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicLoad_0806ad() {
-  var res : i32 = atomicLoad(&(sb_rw.arg_0));
+void atomicLoad_0806ad(device SB_RW& sb_rw) {
+  int res = atomic_load_explicit(&(sb_rw.arg_0), memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicLoad_0806ad();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicLoad_0806ad(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicLoad_0806ad();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicLoad_0806ad(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicLoad/361bf1.wgsl.expected.msl b/test/intrinsics/gen/atomicLoad/361bf1.wgsl.expected.msl
index 7d5661afab..b7c08b54a9 100644
--- a/test/intrinsics/gen/atomicLoad/361bf1.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicLoad/361bf1.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicLoad_361bf1(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : u32 = atomicLoad(&(*(tint_symbol)));
+using namespace metal;
+void atomicLoad_361bf1(threadgroup atomic_uint* const tint_symbol_1) {
+  uint res = atomic_load_explicit(&(*(tint_symbol_1)), memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicLoad_361bf1(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicLoad_361bf1(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicLoad/afcc03.wgsl.expected.msl b/test/intrinsics/gen/atomicLoad/afcc03.wgsl.expected.msl
index 6304ed0844..7b6310acc3 100644
--- a/test/intrinsics/gen/atomicLoad/afcc03.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicLoad/afcc03.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicLoad_afcc03(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : i32 = atomicLoad(&(*(tint_symbol)));
+using namespace metal;
+void atomicLoad_afcc03(threadgroup atomic_int* const tint_symbol_1) {
+  int res = atomic_load_explicit(&(*(tint_symbol_1)), memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicLoad_afcc03(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicLoad_afcc03(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicLoad/fe6cc3.wgsl.expected.msl b/test/intrinsics/gen/atomicLoad/fe6cc3.wgsl.expected.msl
index 2c1eeff2eb..aef98d3a3d 100644
--- a/test/intrinsics/gen/atomicLoad/fe6cc3.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicLoad/fe6cc3.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicLoad_fe6cc3() {
-  var res : u32 = atomicLoad(&(sb_rw.arg_0));
+void atomicLoad_fe6cc3(device SB_RW& sb_rw) {
+  uint res = atomic_load_explicit(&(sb_rw.arg_0), memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicLoad_fe6cc3();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicLoad_fe6cc3(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicLoad_fe6cc3();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicLoad_fe6cc3(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicMax/51b9be.wgsl.expected.msl b/test/intrinsics/gen/atomicMax/51b9be.wgsl.expected.msl
index 9a9fb1dfdf..8d45595d8d 100644
--- a/test/intrinsics/gen/atomicMax/51b9be.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicMax/51b9be.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicMax_51b9be() {
-  var res : u32 = atomicMax(&(sb_rw.arg_0), 1u);
+void atomicMax_51b9be(device SB_RW& sb_rw) {
+  uint res = atomic_fetch_max_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicMax_51b9be();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicMax_51b9be(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicMax_51b9be();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicMax_51b9be(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicMax/92aa72.wgsl.expected.msl b/test/intrinsics/gen/atomicMax/92aa72.wgsl.expected.msl
index a2997a1736..12babbdea1 100644
--- a/test/intrinsics/gen/atomicMax/92aa72.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicMax/92aa72.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicMax_92aa72() {
-  var res : i32 = atomicMax(&(sb_rw.arg_0), 1);
+void atomicMax_92aa72(device SB_RW& sb_rw) {
+  int res = atomic_fetch_max_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicMax_92aa72();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicMax_92aa72(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicMax_92aa72();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicMax_92aa72(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicMax/a89cc3.wgsl.expected.msl b/test/intrinsics/gen/atomicMax/a89cc3.wgsl.expected.msl
index bdb46ff9b2..26d0ee6648 100644
--- a/test/intrinsics/gen/atomicMax/a89cc3.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicMax/a89cc3.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicMax_a89cc3(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : i32 = atomicMax(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicMax_a89cc3(threadgroup atomic_int* const tint_symbol_1) {
+  int res = atomic_fetch_max_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicMax_a89cc3(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicMax_a89cc3(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicMax/beccfc.wgsl.expected.msl b/test/intrinsics/gen/atomicMax/beccfc.wgsl.expected.msl
index 5dc602ca17..dd139eca54 100644
--- a/test/intrinsics/gen/atomicMax/beccfc.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicMax/beccfc.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicMax_beccfc(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : u32 = atomicMax(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicMax_beccfc(threadgroup atomic_uint* const tint_symbol_1) {
+  uint res = atomic_fetch_max_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicMax_beccfc(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicMax_beccfc(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicMin/278235.wgsl.expected.msl b/test/intrinsics/gen/atomicMin/278235.wgsl.expected.msl
index 7eb65c5168..282c9571d3 100644
--- a/test/intrinsics/gen/atomicMin/278235.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicMin/278235.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicMin_278235(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : i32 = atomicMin(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicMin_278235(threadgroup atomic_int* const tint_symbol_1) {
+  int res = atomic_fetch_min_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicMin_278235(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicMin_278235(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicMin/69d383.wgsl.expected.msl b/test/intrinsics/gen/atomicMin/69d383.wgsl.expected.msl
index 5ed9a710f8..4e6ec9280a 100644
--- a/test/intrinsics/gen/atomicMin/69d383.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicMin/69d383.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicMin_69d383(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : u32 = atomicMin(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicMin_69d383(threadgroup atomic_uint* const tint_symbol_1) {
+  uint res = atomic_fetch_min_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicMin_69d383(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicMin_69d383(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicMin/8e38dc.wgsl.expected.msl b/test/intrinsics/gen/atomicMin/8e38dc.wgsl.expected.msl
index dc15ddd6d8..5de8f10205 100644
--- a/test/intrinsics/gen/atomicMin/8e38dc.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicMin/8e38dc.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicMin_8e38dc() {
-  var res : i32 = atomicMin(&(sb_rw.arg_0), 1);
+void atomicMin_8e38dc(device SB_RW& sb_rw) {
+  int res = atomic_fetch_min_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicMin_8e38dc();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicMin_8e38dc(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicMin_8e38dc();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicMin_8e38dc(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicMin/c67a74.wgsl.expected.msl b/test/intrinsics/gen/atomicMin/c67a74.wgsl.expected.msl
index 1b28fe83d8..04d9d54757 100644
--- a/test/intrinsics/gen/atomicMin/c67a74.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicMin/c67a74.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicMin_c67a74() {
-  var res : u32 = atomicMin(&(sb_rw.arg_0), 1u);
+void atomicMin_c67a74(device SB_RW& sb_rw) {
+  uint res = atomic_fetch_min_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicMin_c67a74();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicMin_c67a74(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicMin_c67a74();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicMin_c67a74(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicOr/5e3d61.wgsl.expected.msl b/test/intrinsics/gen/atomicOr/5e3d61.wgsl.expected.msl
index e38ef78c6d..3f85fc00a8 100644
--- a/test/intrinsics/gen/atomicOr/5e3d61.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicOr/5e3d61.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicOr_5e3d61(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : u32 = atomicOr(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicOr_5e3d61(threadgroup atomic_uint* const tint_symbol_1) {
+  uint res = atomic_fetch_or_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicOr_5e3d61(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicOr_5e3d61(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicOr/5e95d4.wgsl.expected.msl b/test/intrinsics/gen/atomicOr/5e95d4.wgsl.expected.msl
index e8ae23b0a5..a711f2d985 100644
--- a/test/intrinsics/gen/atomicOr/5e95d4.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicOr/5e95d4.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicOr_5e95d4() {
-  var res : u32 = atomicOr(&(sb_rw.arg_0), 1u);
+void atomicOr_5e95d4(device SB_RW& sb_rw) {
+  uint res = atomic_fetch_or_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicOr_5e95d4();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicOr_5e95d4(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicOr_5e95d4();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicOr_5e95d4(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicOr/8d96a0.wgsl.expected.msl b/test/intrinsics/gen/atomicOr/8d96a0.wgsl.expected.msl
index 22c839df6a..3602800462 100644
--- a/test/intrinsics/gen/atomicOr/8d96a0.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicOr/8d96a0.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicOr_8d96a0() {
-  var res : i32 = atomicOr(&(sb_rw.arg_0), 1);
+void atomicOr_8d96a0(device SB_RW& sb_rw) {
+  int res = atomic_fetch_or_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicOr_8d96a0();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicOr_8d96a0(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicOr_8d96a0();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicOr_8d96a0(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicOr/d09248.wgsl.expected.msl b/test/intrinsics/gen/atomicOr/d09248.wgsl.expected.msl
index e8c75dd052..620f392642 100644
--- a/test/intrinsics/gen/atomicOr/d09248.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicOr/d09248.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicOr_d09248(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : i32 = atomicOr(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicOr_d09248(threadgroup atomic_int* const tint_symbol_1) {
+  int res = atomic_fetch_or_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicOr_d09248(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicOr_d09248(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicStore/726882.wgsl.expected.msl b/test/intrinsics/gen/atomicStore/726882.wgsl.expected.msl
index c2468d6e5f..bebd8e58ce 100644
--- a/test/intrinsics/gen/atomicStore/726882.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicStore/726882.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicStore_726882(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  atomicStore(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicStore_726882(threadgroup atomic_uint* const tint_symbol_1) {
+  atomic_store_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicStore_726882(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicStore_726882(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicStore/8bea94.wgsl.expected.msl b/test/intrinsics/gen/atomicStore/8bea94.wgsl.expected.msl
index 651a1da6e7..9f9249a583 100644
--- a/test/intrinsics/gen/atomicStore/8bea94.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicStore/8bea94.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicStore_8bea94(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  atomicStore(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicStore_8bea94(threadgroup atomic_int* const tint_symbol_1) {
+  atomic_store_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicStore_8bea94(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicStore_8bea94(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicStore/cdc29e.wgsl.expected.msl b/test/intrinsics/gen/atomicStore/cdc29e.wgsl.expected.msl
index cefa61ef76..24231c1cd6 100644
--- a/test/intrinsics/gen/atomicStore/cdc29e.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicStore/cdc29e.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicStore_cdc29e() {
-  atomicStore(&(sb_rw.arg_0), 1u);
+void atomicStore_cdc29e(device SB_RW& sb_rw) {
+  atomic_store_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicStore_cdc29e();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicStore_cdc29e(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicStore_cdc29e();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicStore_cdc29e(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicStore/d1e9a6.wgsl.expected.msl b/test/intrinsics/gen/atomicStore/d1e9a6.wgsl.expected.msl
index 98cae7a3e5..4d3c46816e 100644
--- a/test/intrinsics/gen/atomicStore/d1e9a6.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicStore/d1e9a6.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicStore_d1e9a6() {
-  atomicStore(&(sb_rw.arg_0), 1);
+void atomicStore_d1e9a6(device SB_RW& sb_rw) {
+  atomic_store_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicStore_d1e9a6();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicStore_d1e9a6(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicStore_d1e9a6();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicStore_d1e9a6(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicXor/54510e.wgsl.expected.msl b/test/intrinsics/gen/atomicXor/54510e.wgsl.expected.msl
index 9bccb9b35c..c43fc6be63 100644
--- a/test/intrinsics/gen/atomicXor/54510e.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicXor/54510e.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<u32>;
+  /* 0x0000 */ atomic_uint arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicXor_54510e() {
-  var res : u32 = atomicXor(&(sb_rw.arg_0), 1u);
+void atomicXor_54510e(device SB_RW& sb_rw) {
+  uint res = atomic_fetch_xor_explicit(&(sb_rw.arg_0), 1u, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicXor_54510e();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicXor_54510e(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicXor_54510e();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicXor_54510e(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__u32
diff --git a/test/intrinsics/gen/atomicXor/75dc95.wgsl.expected.msl b/test/intrinsics/gen/atomicXor/75dc95.wgsl.expected.msl
index 32bbb93997..e1b8265805 100644
--- a/test/intrinsics/gen/atomicXor/75dc95.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicXor/75dc95.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicXor_75dc95(tint_symbol : ptr<workgroup, atomic<i32>>) {
-  var res : i32 = atomicXor(&(*(tint_symbol)), 1);
+using namespace metal;
+void atomicXor_75dc95(threadgroup atomic_int* const tint_symbol_1) {
+  int res = atomic_fetch_xor_explicit(&(*(tint_symbol_1)), 1, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<i32>;
-  atomicXor_75dc95(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), int(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicXor_75dc95(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope
diff --git a/test/intrinsics/gen/atomicXor/c1b78c.wgsl.expected.msl b/test/intrinsics/gen/atomicXor/c1b78c.wgsl.expected.msl
index bd565c88b5..d48ee317d4 100644
--- a/test/intrinsics/gen/atomicXor/c1b78c.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicXor/c1b78c.wgsl.expected.msl
@@ -1,25 +1,21 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-[[block]]
+using namespace metal;
 struct SB_RW {
-  arg_0 : atomic<i32>;
+  /* 0x0000 */ atomic_int arg_0;
 };
 
-[[group(0), binding(0)]] var<storage, read_write> sb_rw : SB_RW;
-
-fn atomicXor_c1b78c() {
-  var res : i32 = atomicXor(&(sb_rw.arg_0), 1);
+void atomicXor_c1b78c(device SB_RW& sb_rw) {
+  int res = atomic_fetch_xor_explicit(&(sb_rw.arg_0), 1, memory_order_relaxed);
 }
 
-[[stage(fragment)]]
-fn fragment_main() {
-  atomicXor_c1b78c();
+fragment void fragment_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicXor_c1b78c(sb_rw);
+  return;
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  atomicXor_c1b78c();
+kernel void compute_main(device SB_RW& sb_rw [[buffer(0)]]) {
+  atomicXor_c1b78c(sb_rw);
+  return;
 }
 
-Failed to generate: error: unknown type in EmitType: __atomic__i32
diff --git a/test/intrinsics/gen/atomicXor/c8e6be.wgsl.expected.msl b/test/intrinsics/gen/atomicXor/c8e6be.wgsl.expected.msl
index 4ee0172668..3f0283dd39 100644
--- a/test/intrinsics/gen/atomicXor/c8e6be.wgsl.expected.msl
+++ b/test/intrinsics/gen/atomicXor/c8e6be.wgsl.expected.msl
@@ -1,14 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
 
-
-fn atomicXor_c8e6be(tint_symbol : ptr<workgroup, atomic<u32>>) {
-  var res : u32 = atomicXor(&(*(tint_symbol)), 1u);
+using namespace metal;
+void atomicXor_c8e6be(threadgroup atomic_uint* const tint_symbol_1) {
+  uint res = atomic_fetch_xor_explicit(&(*(tint_symbol_1)), 1u, memory_order_relaxed);
 }
 
-[[stage(compute)]]
-fn compute_main() {
-  [[internal(disable_validation__function_var_storage_class)]] var<workgroup> tint_symbol_1 : atomic<u32>;
-  atomicXor_c8e6be(&(tint_symbol_1));
+kernel void compute_main(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup atomic_uint tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    atomic_store_explicit(&(tint_symbol_2), uint(), memory_order_relaxed);
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  atomicXor_c8e6be(&(tint_symbol_2));
+  return;
 }
 
-error: cannot declare an atomic var in a function scope