[msl] Handle textureSample builtins

Replace them with texture.sample() member functions. Add support for
MemberFunctionCall to the printer.

Bug: 42251016
Change-Id: I24246253c9e3160ebe712d6e2e7f9d951e94a738
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/192305
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/src/tint/lang/msl/builtin_fn.cc b/src/tint/lang/msl/builtin_fn.cc
index c259530..46b324c 100644
--- a/src/tint/lang/msl/builtin_fn.cc
+++ b/src/tint/lang/msl/builtin_fn.cc
@@ -64,6 +64,8 @@
             return "atomic_load_explicit";
         case BuiltinFn::kAtomicStoreExplicit:
             return "atomic_store_explicit";
+        case BuiltinFn::kSample:
+            return "sample";
         case BuiltinFn::kThreadgroupBarrier:
             return "threadgroup_barrier";
     }
diff --git a/src/tint/lang/msl/builtin_fn.h b/src/tint/lang/msl/builtin_fn.h
index dcdf48e..17703bf 100644
--- a/src/tint/lang/msl/builtin_fn.h
+++ b/src/tint/lang/msl/builtin_fn.h
@@ -58,6 +58,7 @@
     kAtomicFetchXorExplicit,
     kAtomicLoadExplicit,
     kAtomicStoreExplicit,
+    kSample,
     kThreadgroupBarrier,
     kNone,
 };
diff --git a/src/tint/lang/msl/intrinsic/data.cc b/src/tint/lang/msl/intrinsic/data.cc
index c632629..c3afbb3 100644
--- a/src/tint/lang/msl/intrinsic/data.cc
+++ b/src/tint/lang/msl/intrinsic/data.cc
@@ -116,6 +116,80 @@
 };
 
 
+/// TypeMatcher for 'type f32'
+constexpr TypeMatcher kF32Matcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+    if (!MatchF32(state, ty)) {
+      return nullptr;
+    }
+    return BuildF32(state, ty);
+  },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
+    out << style::Type("f32");
+  }
+};
+
+
+/// TypeMatcher for 'type vec2'
+constexpr TypeMatcher kVec2Matcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+  const Type* T = nullptr;
+    if (!MatchVec2(state, ty, T)) {
+      return nullptr;
+    }
+    T = state.Type(T);
+    if (T == nullptr) {
+      return nullptr;
+    }
+    return BuildVec2(state, ty, T);
+  },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+  state->PrintType(T);
+    out << style::Type("vec2", "<", T, ">");
+  }
+};
+
+
+/// TypeMatcher for 'type vec3'
+constexpr TypeMatcher kVec3Matcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+  const Type* T = nullptr;
+    if (!MatchVec3(state, ty, T)) {
+      return nullptr;
+    }
+    T = state.Type(T);
+    if (T == nullptr) {
+      return nullptr;
+    }
+    return BuildVec3(state, ty, T);
+  },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+  state->PrintType(T);
+    out << style::Type("vec3", "<", T, ">");
+  }
+};
+
+
+/// TypeMatcher for 'type vec4'
+constexpr TypeMatcher kVec4Matcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+  const Type* T = nullptr;
+    if (!MatchVec4(state, ty, T)) {
+      return nullptr;
+    }
+    T = state.Type(T);
+    if (T == nullptr) {
+      return nullptr;
+    }
+    return BuildVec4(state, ty, T);
+  },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+  state->PrintType(T);
+    out << style::Type("vec4", "<", T, ">");
+  }
+};
+
+
 /// TypeMatcher for 'type atomic'
 constexpr TypeMatcher kAtomicMatcher {
 /* match */ [](MatchState& state, const Type* ty) -> const Type* {
@@ -168,6 +242,196 @@
 };
 
 
+/// TypeMatcher for 'type sampler'
+constexpr TypeMatcher kSamplerMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+    if (!MatchSampler(state, ty)) {
+      return nullptr;
+    }
+    return BuildSampler(state, ty);
+  },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
+    out << style::Type("sampler");
+  }
+};
+
+
+/// TypeMatcher for 'type texture_1d'
+constexpr TypeMatcher kTexture1DMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+  const Type* T = nullptr;
+    if (!MatchTexture1D(state, ty, T)) {
+      return nullptr;
+    }
+    T = state.Type(T);
+    if (T == nullptr) {
+      return nullptr;
+    }
+    return BuildTexture1D(state, ty, T);
+  },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+  state->PrintType(T);
+    out << style::Type("texture_1d", "<", T, ">");
+  }
+};
+
+
+/// TypeMatcher for 'type texture_2d'
+constexpr TypeMatcher kTexture2DMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+  const Type* T = nullptr;
+    if (!MatchTexture2D(state, ty, T)) {
+      return nullptr;
+    }
+    T = state.Type(T);
+    if (T == nullptr) {
+      return nullptr;
+    }
+    return BuildTexture2D(state, ty, T);
+  },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+  state->PrintType(T);
+    out << style::Type("texture_2d", "<", T, ">");
+  }
+};
+
+
+/// TypeMatcher for 'type texture_2d_array'
+constexpr TypeMatcher kTexture2DArrayMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+  const Type* T = nullptr;
+    if (!MatchTexture2DArray(state, ty, T)) {
+      return nullptr;
+    }
+    T = state.Type(T);
+    if (T == nullptr) {
+      return nullptr;
+    }
+    return BuildTexture2DArray(state, ty, T);
+  },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+  state->PrintType(T);
+    out << style::Type("texture_2d_array", "<", T, ">");
+  }
+};
+
+
+/// TypeMatcher for 'type texture_3d'
+constexpr TypeMatcher kTexture3DMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+  const Type* T = nullptr;
+    if (!MatchTexture3D(state, ty, T)) {
+      return nullptr;
+    }
+    T = state.Type(T);
+    if (T == nullptr) {
+      return nullptr;
+    }
+    return BuildTexture3D(state, ty, T);
+  },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+  state->PrintType(T);
+    out << style::Type("texture_3d", "<", T, ">");
+  }
+};
+
+
+/// TypeMatcher for 'type texture_cube'
+constexpr TypeMatcher kTextureCubeMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+  const Type* T = nullptr;
+    if (!MatchTextureCube(state, ty, T)) {
+      return nullptr;
+    }
+    T = state.Type(T);
+    if (T == nullptr) {
+      return nullptr;
+    }
+    return BuildTextureCube(state, ty, T);
+  },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+  state->PrintType(T);
+    out << style::Type("texture_cube", "<", T, ">");
+  }
+};
+
+
+/// TypeMatcher for 'type texture_cube_array'
+constexpr TypeMatcher kTextureCubeArrayMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+  const Type* T = nullptr;
+    if (!MatchTextureCubeArray(state, ty, T)) {
+      return nullptr;
+    }
+    T = state.Type(T);
+    if (T == nullptr) {
+      return nullptr;
+    }
+    return BuildTextureCubeArray(state, ty, T);
+  },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {StyledText T;
+  state->PrintType(T);
+    out << style::Type("texture_cube_array", "<", T, ">");
+  }
+};
+
+
+/// TypeMatcher for 'type texture_depth_2d'
+constexpr TypeMatcher kTextureDepth2DMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+    if (!MatchTextureDepth2D(state, ty)) {
+      return nullptr;
+    }
+    return BuildTextureDepth2D(state, ty);
+  },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
+    out << style::Type("texture_depth_2d");
+  }
+};
+
+
+/// TypeMatcher for 'type texture_depth_2d_array'
+constexpr TypeMatcher kTextureDepth2DArrayMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+    if (!MatchTextureDepth2DArray(state, ty)) {
+      return nullptr;
+    }
+    return BuildTextureDepth2DArray(state, ty);
+  },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
+    out << style::Type("texture_depth_2d_array");
+  }
+};
+
+
+/// TypeMatcher for 'type texture_depth_cube'
+constexpr TypeMatcher kTextureDepthCubeMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+    if (!MatchTextureDepthCube(state, ty)) {
+      return nullptr;
+    }
+    return BuildTextureDepthCube(state, ty);
+  },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
+    out << style::Type("texture_depth_cube");
+  }
+};
+
+
+/// TypeMatcher for 'type texture_depth_cube_array'
+constexpr TypeMatcher kTextureDepthCubeArrayMatcher {
+/* match */ [](MatchState& state, const Type* ty) -> const Type* {
+    if (!MatchTextureDepthCubeArray(state, ty)) {
+      return nullptr;
+    }
+    return BuildTextureDepthCubeArray(state, ty);
+  },
+/* print */ []([[maybe_unused]] MatchState* state, StyledText& out) {
+    out << style::Type("texture_depth_cube_array");
+  }
+};
+
+
 /// TypeMatcher for 'match iu32'
 constexpr TypeMatcher kIu32Matcher {
 /* match */ [](MatchState& state, const Type* ty) -> const Type* {
@@ -236,9 +500,24 @@
   /* [2] */ kBoolMatcher,
   /* [3] */ kI32Matcher,
   /* [4] */ kU32Matcher,
-  /* [5] */ kAtomicMatcher,
-  /* [6] */ kPtrMatcher,
-  /* [7] */ kIu32Matcher,
+  /* [5] */ kF32Matcher,
+  /* [6] */ kVec2Matcher,
+  /* [7] */ kVec3Matcher,
+  /* [8] */ kVec4Matcher,
+  /* [9] */ kAtomicMatcher,
+  /* [10] */ kPtrMatcher,
+  /* [11] */ kSamplerMatcher,
+  /* [12] */ kTexture1DMatcher,
+  /* [13] */ kTexture2DMatcher,
+  /* [14] */ kTexture2DArrayMatcher,
+  /* [15] */ kTexture3DMatcher,
+  /* [16] */ kTextureCubeMatcher,
+  /* [17] */ kTextureCubeArrayMatcher,
+  /* [18] */ kTextureDepth2DMatcher,
+  /* [19] */ kTextureDepth2DArrayMatcher,
+  /* [20] */ kTextureDepthCubeMatcher,
+  /* [21] */ kTextureDepthCubeArrayMatcher,
+  /* [22] */ kIu32Matcher,
 };
 
 /// The template numbers, and number matchers
@@ -251,17 +530,44 @@
 };
 
 constexpr MatcherIndex kMatcherIndices[] = {
-  /* [0] */ MatcherIndex(6),
+  /* [0] */ MatcherIndex(10),
   /* [1] */ MatcherIndex(1),
-  /* [2] */ MatcherIndex(5),
+  /* [2] */ MatcherIndex(9),
   /* [3] */ MatcherIndex(0),
   /* [4] */ MatcherIndex(2),
-  /* [5] */ MatcherIndex(6),
+  /* [5] */ MatcherIndex(10),
   /* [6] */ MatcherIndex(3),
   /* [7] */ MatcherIndex(0),
   /* [8] */ MatcherIndex(2),
-  /* [9] */ MatcherIndex(7),
-  /* [10] */ MatcherIndex(4),
+  /* [9] */ MatcherIndex(8),
+  /* [10] */ MatcherIndex(5),
+  /* [11] */ MatcherIndex(12),
+  /* [12] */ MatcherIndex(5),
+  /* [13] */ MatcherIndex(13),
+  /* [14] */ MatcherIndex(5),
+  /* [15] */ MatcherIndex(6),
+  /* [16] */ MatcherIndex(5),
+  /* [17] */ MatcherIndex(6),
+  /* [18] */ MatcherIndex(3),
+  /* [19] */ MatcherIndex(14),
+  /* [20] */ MatcherIndex(5),
+  /* [21] */ MatcherIndex(15),
+  /* [22] */ MatcherIndex(5),
+  /* [23] */ MatcherIndex(7),
+  /* [24] */ MatcherIndex(5),
+  /* [25] */ MatcherIndex(7),
+  /* [26] */ MatcherIndex(3),
+  /* [27] */ MatcherIndex(16),
+  /* [28] */ MatcherIndex(5),
+  /* [29] */ MatcherIndex(17),
+  /* [30] */ MatcherIndex(5),
+  /* [31] */ MatcherIndex(22),
+  /* [32] */ MatcherIndex(4),
+  /* [33] */ MatcherIndex(11),
+  /* [34] */ MatcherIndex(18),
+  /* [35] */ MatcherIndex(19),
+  /* [36] */ MatcherIndex(20),
+  /* [37] */ MatcherIndex(21),
 };
 
 static_assert(MatcherIndicesIndex::CanIndex(kMatcherIndices),
@@ -286,37 +592,232 @@
   {
     /* [3] */
     /* usage */ core::ParameterUsage::kNone,
-    /* matcher_indices */ MatcherIndicesIndex(10),
+    /* matcher_indices */ MatcherIndicesIndex(32),
   },
   {
     /* [4] */
     /* usage */ core::ParameterUsage::kNone,
-    /* matcher_indices */ MatcherIndicesIndex(10),
+    /* matcher_indices */ MatcherIndicesIndex(32),
   },
   {
     /* [5] */
+    /* usage */ core::ParameterUsage::kTexture,
+    /* matcher_indices */ MatcherIndicesIndex(19),
+  },
+  {
+    /* [6] */
+    /* usage */ core::ParameterUsage::kSampler,
+    /* matcher_indices */ MatcherIndicesIndex(33),
+  },
+  {
+    /* [7] */
+    /* usage */ core::ParameterUsage::kCoords,
+    /* matcher_indices */ MatcherIndicesIndex(15),
+  },
+  {
+    /* [8] */
+    /* usage */ core::ParameterUsage::kArrayIndex,
+    /* matcher_indices */ MatcherIndicesIndex(3),
+  },
+  {
+    /* [9] */
+    /* usage */ core::ParameterUsage::kOffset,
+    /* matcher_indices */ MatcherIndicesIndex(17),
+  },
+  {
+    /* [10] */
+    /* usage */ core::ParameterUsage::kTexture,
+    /* matcher_indices */ MatcherIndicesIndex(35),
+  },
+  {
+    /* [11] */
+    /* usage */ core::ParameterUsage::kSampler,
+    /* matcher_indices */ MatcherIndicesIndex(33),
+  },
+  {
+    /* [12] */
+    /* usage */ core::ParameterUsage::kCoords,
+    /* matcher_indices */ MatcherIndicesIndex(15),
+  },
+  {
+    /* [13] */
+    /* usage */ core::ParameterUsage::kArrayIndex,
+    /* matcher_indices */ MatcherIndicesIndex(3),
+  },
+  {
+    /* [14] */
+    /* usage */ core::ParameterUsage::kOffset,
+    /* matcher_indices */ MatcherIndicesIndex(17),
+  },
+  {
+    /* [15] */
+    /* usage */ core::ParameterUsage::kTexture,
+    /* matcher_indices */ MatcherIndicesIndex(13),
+  },
+  {
+    /* [16] */
+    /* usage */ core::ParameterUsage::kSampler,
+    /* matcher_indices */ MatcherIndicesIndex(33),
+  },
+  {
+    /* [17] */
+    /* usage */ core::ParameterUsage::kCoords,
+    /* matcher_indices */ MatcherIndicesIndex(15),
+  },
+  {
+    /* [18] */
+    /* usage */ core::ParameterUsage::kOffset,
+    /* matcher_indices */ MatcherIndicesIndex(17),
+  },
+  {
+    /* [19] */
+    /* usage */ core::ParameterUsage::kTexture,
+    /* matcher_indices */ MatcherIndicesIndex(21),
+  },
+  {
+    /* [20] */
+    /* usage */ core::ParameterUsage::kSampler,
+    /* matcher_indices */ MatcherIndicesIndex(33),
+  },
+  {
+    /* [21] */
+    /* usage */ core::ParameterUsage::kCoords,
+    /* matcher_indices */ MatcherIndicesIndex(23),
+  },
+  {
+    /* [22] */
+    /* usage */ core::ParameterUsage::kOffset,
+    /* matcher_indices */ MatcherIndicesIndex(25),
+  },
+  {
+    /* [23] */
+    /* usage */ core::ParameterUsage::kTexture,
+    /* matcher_indices */ MatcherIndicesIndex(29),
+  },
+  {
+    /* [24] */
+    /* usage */ core::ParameterUsage::kSampler,
+    /* matcher_indices */ MatcherIndicesIndex(33),
+  },
+  {
+    /* [25] */
+    /* usage */ core::ParameterUsage::kCoords,
+    /* matcher_indices */ MatcherIndicesIndex(23),
+  },
+  {
+    /* [26] */
+    /* usage */ core::ParameterUsage::kArrayIndex,
+    /* matcher_indices */ MatcherIndicesIndex(3),
+  },
+  {
+    /* [27] */
+    /* usage */ core::ParameterUsage::kTexture,
+    /* matcher_indices */ MatcherIndicesIndex(34),
+  },
+  {
+    /* [28] */
+    /* usage */ core::ParameterUsage::kSampler,
+    /* matcher_indices */ MatcherIndicesIndex(33),
+  },
+  {
+    /* [29] */
+    /* usage */ core::ParameterUsage::kCoords,
+    /* matcher_indices */ MatcherIndicesIndex(15),
+  },
+  {
+    /* [30] */
+    /* usage */ core::ParameterUsage::kOffset,
+    /* matcher_indices */ MatcherIndicesIndex(17),
+  },
+  {
+    /* [31] */
+    /* usage */ core::ParameterUsage::kTexture,
+    /* matcher_indices */ MatcherIndicesIndex(37),
+  },
+  {
+    /* [32] */
+    /* usage */ core::ParameterUsage::kSampler,
+    /* matcher_indices */ MatcherIndicesIndex(33),
+  },
+  {
+    /* [33] */
+    /* usage */ core::ParameterUsage::kCoords,
+    /* matcher_indices */ MatcherIndicesIndex(23),
+  },
+  {
+    /* [34] */
+    /* usage */ core::ParameterUsage::kArrayIndex,
+    /* matcher_indices */ MatcherIndicesIndex(3),
+  },
+  {
+    /* [35] */
     /* usage */ core::ParameterUsage::kNone,
     /* matcher_indices */ MatcherIndicesIndex(0),
   },
   {
-    /* [6] */
+    /* [36] */
     /* usage */ core::ParameterUsage::kNone,
     /* matcher_indices */ MatcherIndicesIndex(3),
   },
   {
-    /* [7] */
+    /* [37] */
     /* usage */ core::ParameterUsage::kNone,
+    /* matcher_indices */ MatcherIndicesIndex(32),
+  },
+  {
+    /* [38] */
+    /* usage */ core::ParameterUsage::kTexture,
+    /* matcher_indices */ MatcherIndicesIndex(11),
+  },
+  {
+    /* [39] */
+    /* usage */ core::ParameterUsage::kSampler,
+    /* matcher_indices */ MatcherIndicesIndex(33),
+  },
+  {
+    /* [40] */
+    /* usage */ core::ParameterUsage::kCoords,
     /* matcher_indices */ MatcherIndicesIndex(10),
   },
   {
-    /* [8] */
+    /* [41] */
+    /* usage */ core::ParameterUsage::kTexture,
+    /* matcher_indices */ MatcherIndicesIndex(27),
+  },
+  {
+    /* [42] */
+    /* usage */ core::ParameterUsage::kSampler,
+    /* matcher_indices */ MatcherIndicesIndex(33),
+  },
+  {
+    /* [43] */
+    /* usage */ core::ParameterUsage::kCoords,
+    /* matcher_indices */ MatcherIndicesIndex(23),
+  },
+  {
+    /* [44] */
+    /* usage */ core::ParameterUsage::kTexture,
+    /* matcher_indices */ MatcherIndicesIndex(36),
+  },
+  {
+    /* [45] */
+    /* usage */ core::ParameterUsage::kSampler,
+    /* matcher_indices */ MatcherIndicesIndex(33),
+  },
+  {
+    /* [46] */
+    /* usage */ core::ParameterUsage::kCoords,
+    /* matcher_indices */ MatcherIndicesIndex(23),
+  },
+  {
+    /* [47] */
     /* usage */ core::ParameterUsage::kNone,
     /* matcher_indices */ MatcherIndicesIndex(0),
   },
   {
-    /* [9] */
+    /* [48] */
     /* usage */ core::ParameterUsage::kNone,
-    /* matcher_indices */ MatcherIndicesIndex(10),
+    /* matcher_indices */ MatcherIndicesIndex(32),
   },
 };
 
@@ -327,15 +828,21 @@
   {
     /* [0] */
     /* name */ "T",
-    /* matcher_indices */ MatcherIndicesIndex(9),
+    /* matcher_indices */ MatcherIndicesIndex(31),
     /* kind */ TemplateInfo::Kind::kType,
   },
   {
     /* [1] */
     /* name */ "S",
-    /* matcher_indices */ MatcherIndicesIndex(10),
+    /* matcher_indices */ MatcherIndicesIndex(32),
     /* kind */ TemplateInfo::Kind::kNumber,
   },
+  {
+    /* [2] */
+    /* name */ "A",
+    /* matcher_indices */ MatcherIndicesIndex(31),
+    /* kind */ TemplateInfo::Kind::kType,
+  },
 };
 
 static_assert(TemplateIndex::CanIndex(kTemplates),
@@ -344,6 +851,171 @@
 constexpr OverloadInfo kOverloads[] = {
   {
     /* [0] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 3,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 0,
+    /* templates */ TemplateIndex(/* invalid */),
+    /* parameters */ ParameterIndex(38),
+    /* return_matcher_indices */ MatcherIndicesIndex(9),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [1] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 3,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 0,
+    /* templates */ TemplateIndex(/* invalid */),
+    /* parameters */ ParameterIndex(15),
+    /* return_matcher_indices */ MatcherIndicesIndex(9),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [2] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 4,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 0,
+    /* templates */ TemplateIndex(/* invalid */),
+    /* parameters */ ParameterIndex(15),
+    /* return_matcher_indices */ MatcherIndicesIndex(9),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [3] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 4,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 1,
+    /* templates */ TemplateIndex(2),
+    /* parameters */ ParameterIndex(5),
+    /* return_matcher_indices */ MatcherIndicesIndex(9),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [4] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 5,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 1,
+    /* templates */ TemplateIndex(2),
+    /* parameters */ ParameterIndex(5),
+    /* return_matcher_indices */ MatcherIndicesIndex(9),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [5] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 3,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 0,
+    /* templates */ TemplateIndex(/* invalid */),
+    /* parameters */ ParameterIndex(19),
+    /* return_matcher_indices */ MatcherIndicesIndex(9),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [6] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 4,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 0,
+    /* templates */ TemplateIndex(/* invalid */),
+    /* parameters */ ParameterIndex(19),
+    /* return_matcher_indices */ MatcherIndicesIndex(9),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [7] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 3,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 0,
+    /* templates */ TemplateIndex(/* invalid */),
+    /* parameters */ ParameterIndex(41),
+    /* return_matcher_indices */ MatcherIndicesIndex(9),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [8] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 4,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 1,
+    /* templates */ TemplateIndex(2),
+    /* parameters */ ParameterIndex(23),
+    /* return_matcher_indices */ MatcherIndicesIndex(9),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [9] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 3,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 0,
+    /* templates */ TemplateIndex(/* invalid */),
+    /* parameters */ ParameterIndex(27),
+    /* return_matcher_indices */ MatcherIndicesIndex(10),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [10] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 4,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 0,
+    /* templates */ TemplateIndex(/* invalid */),
+    /* parameters */ ParameterIndex(27),
+    /* return_matcher_indices */ MatcherIndicesIndex(10),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [11] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 4,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 1,
+    /* templates */ TemplateIndex(2),
+    /* parameters */ ParameterIndex(10),
+    /* return_matcher_indices */ MatcherIndicesIndex(10),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [12] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 5,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 1,
+    /* templates */ TemplateIndex(2),
+    /* parameters */ ParameterIndex(10),
+    /* return_matcher_indices */ MatcherIndicesIndex(10),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [13] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 3,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 0,
+    /* templates */ TemplateIndex(/* invalid */),
+    /* parameters */ ParameterIndex(44),
+    /* return_matcher_indices */ MatcherIndicesIndex(10),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [14] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 4,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 1,
+    /* templates */ TemplateIndex(2),
+    /* parameters */ ParameterIndex(31),
+    /* return_matcher_indices */ MatcherIndicesIndex(10),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [15] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 5,
     /* num_explicit_templates */ 0,
@@ -354,40 +1026,40 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [1] */
+    /* [16] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 3,
     /* num_explicit_templates */ 0,
     /* num_templates   */ 2,
     /* templates */ TemplateIndex(0),
-    /* parameters */ ParameterIndex(5),
+    /* parameters */ ParameterIndex(35),
     /* return_matcher_indices */ MatcherIndicesIndex(3),
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [2] */
+    /* [17] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 2,
     /* num_explicit_templates */ 0,
     /* num_templates   */ 2,
     /* templates */ TemplateIndex(0),
-    /* parameters */ ParameterIndex(8),
+    /* parameters */ ParameterIndex(47),
     /* return_matcher_indices */ MatcherIndicesIndex(3),
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [3] */
+    /* [18] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 3,
     /* num_explicit_templates */ 0,
     /* num_templates   */ 2,
     /* templates */ TemplateIndex(0),
-    /* parameters */ ParameterIndex(5),
+    /* parameters */ ParameterIndex(35),
     /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [4] */
+    /* [19] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 1,
     /* num_explicit_templates */ 0,
@@ -407,73 +1079,93 @@
     /* [0] */
     /* fn atomic_compare_exchange_weak_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, ptr<function, T, read_write>, T, u32, u32) -> bool */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(0),
+    /* overloads */ OverloadIndex(15),
   },
   {
     /* [1] */
     /* fn atomic_exchange_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(1),
+    /* overloads */ OverloadIndex(16),
   },
   {
     /* [2] */
     /* fn atomic_fetch_add_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(1),
+    /* overloads */ OverloadIndex(16),
   },
   {
     /* [3] */
     /* fn atomic_fetch_and_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(1),
+    /* overloads */ OverloadIndex(16),
   },
   {
     /* [4] */
     /* fn atomic_fetch_max_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(1),
+    /* overloads */ OverloadIndex(16),
   },
   {
     /* [5] */
     /* fn atomic_fetch_min_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(1),
+    /* overloads */ OverloadIndex(16),
   },
   {
     /* [6] */
     /* fn atomic_fetch_or_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(1),
+    /* overloads */ OverloadIndex(16),
   },
   {
     /* [7] */
     /* fn atomic_fetch_sub_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(1),
+    /* overloads */ OverloadIndex(16),
   },
   {
     /* [8] */
     /* fn atomic_fetch_xor_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(1),
+    /* overloads */ OverloadIndex(16),
   },
   {
     /* [9] */
     /* fn atomic_load_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, u32) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(2),
+    /* overloads */ OverloadIndex(17),
   },
   {
     /* [10] */
     /* fn atomic_store_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(3),
+    /* overloads */ OverloadIndex(18),
   },
   {
     /* [11] */
+    /* fn sample(texture: texture_1d<f32>, sampler: sampler, coords: f32) -> vec4<f32> */
+    /* fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>) -> vec4<f32> */
+    /* fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32> */
+    /* fn sample[A : iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A) -> vec4<f32> */
+    /* fn sample[A : iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A, @const offset: vec2<i32>) -> vec4<f32> */
+    /* fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32> */
+    /* fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, @const offset: vec3<i32>) -> vec4<f32> */
+    /* fn sample(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32> */
+    /* fn sample[A : iu32](texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: A) -> vec4<f32> */
+    /* fn sample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>) -> f32 */
+    /* fn sample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> f32 */
+    /* fn sample[A : iu32](texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: A) -> f32 */
+    /* fn sample[A : iu32](texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: A, @const offset: vec2<i32>) -> f32 */
+    /* fn sample(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>) -> f32 */
+    /* fn sample[A : iu32](texture: texture_depth_cube_array, sampler: sampler, coords: vec3<f32>, array_index: A) -> f32 */
+    /* num overloads */ 15,
+    /* overloads */ OverloadIndex(0),
+  },
+  {
+    /* [12] */
     /* fn threadgroup_barrier(u32) */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(4),
+    /* overloads */ OverloadIndex(19),
   },
 };
 
diff --git a/src/tint/lang/msl/msl.def b/src/tint/lang/msl/msl.def
index 858ee55..235f697 100644
--- a/src/tint/lang/msl/msl.def
+++ b/src/tint/lang/msl/msl.def
@@ -56,8 +56,23 @@
 type bool
 type i32
 type u32
+type f32
+type vec2<T>
+type vec3<T>
+type vec4<T>
 type atomic<T>
 type ptr<S: address_space, T, A: access>
+type sampler
+type texture_1d<T>
+type texture_2d<T>
+type texture_2d_array<T>
+type texture_3d<T>
+type texture_cube<T>
+type texture_cube_array<T>
+type texture_depth_2d
+type texture_depth_2d_array
+type texture_depth_cube
+type texture_depth_cube_array
 
 ////////////////////////////////////////////////////////////////////////////////
 // Type matchers                                                              //
@@ -79,5 +94,22 @@
 @stage("fragment", "compute") fn atomic_fetch_xor_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) -> T
 @stage("fragment", "compute") fn atomic_load_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, u32) -> T
 @stage("fragment", "compute") fn atomic_store_explicit[T: iu32, S: workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32)
+
+@member_function @stage("fragment") fn sample(texture: texture_1d<f32>, sampler: sampler, coords: f32) -> vec4<f32>
+@member_function @stage("fragment") fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>) -> vec4<f32>
+@member_function @stage("fragment") fn sample(texture: texture_2d<f32>, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> vec4<f32>
+@member_function @stage("fragment") fn sample[A: iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A) -> vec4<f32>
+@member_function @stage("fragment") fn sample[A: iu32](texture: texture_2d_array<f32>, sampler: sampler, coords: vec2<f32>, array_index: A, @const offset: vec2<i32>) -> vec4<f32>
+@member_function @stage("fragment") fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32>
+@member_function @stage("fragment") fn sample(texture: texture_3d<f32>, sampler: sampler, coords: vec3<f32>, @const offset: vec3<i32>) -> vec4<f32>
+@member_function @stage("fragment") fn sample(texture: texture_cube<f32>, sampler: sampler, coords: vec3<f32>) -> vec4<f32>
+@member_function @stage("fragment") fn sample[A: iu32](texture: texture_cube_array<f32>, sampler: sampler, coords: vec3<f32>, array_index: A) -> vec4<f32>
+@member_function @stage("fragment") fn sample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>) -> f32
+@member_function @stage("fragment") fn sample(texture: texture_depth_2d, sampler: sampler, coords: vec2<f32>, @const offset: vec2<i32>) -> f32
+@member_function @stage("fragment") fn sample[A: iu32](texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: A) -> f32
+@member_function @stage("fragment") fn sample[A: iu32](texture: texture_depth_2d_array, sampler: sampler, coords: vec2<f32>, array_index: A, @const offset: vec2<i32>) -> f32
+@member_function @stage("fragment") fn sample(texture: texture_depth_cube, sampler: sampler, coords: vec3<f32>) -> f32
+@member_function @stage("fragment") fn sample[A: iu32](texture: texture_depth_cube_array, sampler: sampler, coords: vec3<f32>, array_index: A) -> f32
+
 @stage("compute") fn threadgroup_barrier(u32)
 
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index 5d74063..2617d6d 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -89,6 +89,7 @@
 #include "src/tint/lang/msl/barrier_type.h"
 #include "src/tint/lang/msl/builtin_fn.h"
 #include "src/tint/lang/msl/ir/builtin_call.h"
+#include "src/tint/lang/msl/ir/member_builtin_call.h"
 #include "src/tint/lang/msl/ir/memory_order.h"
 #include "src/tint/lang/msl/writer/common/printer_support.h"
 #include "src/tint/utils/containers/map.h"
@@ -408,17 +409,20 @@
             [&](const core::ir::Constant* c) { EmitConstant(out, c); },  //
             [&](const core::ir::InstructionResult* r) {
                 Switch(
-                    r->Instruction(),                                                          //
-                    [&](const core::ir::CoreBinary* b) { EmitBinary(out, b); },                //
-                    [&](const core::ir::CoreUnary* u) { EmitUnary(out, u); },                  //
-                    [&](const core::ir::Convert* b) { EmitConvert(out, b); },                  //
-                    [&](const core::ir::Let* l) { out << NameOf(l->Result(0)); },              //
-                    [&](const core::ir::Load* l) { EmitLoad(out, l); },                        //
-                    [&](const core::ir::Construct* c) { EmitConstruct(out, c); },              //
-                    [&](const core::ir::Var* var) { out << NameOf(var->Result(0)); },          //
-                    [&](const core::ir::Bitcast* b) { EmitBitcast(out, b); },                  //
-                    [&](const core::ir::Access* a) { EmitAccess(out, a); },                    //
-                    [&](const msl::ir::BuiltinCall* c) { EmitMslBuiltinCall(out, c); },        //
+                    r->Instruction(),                                                    //
+                    [&](const core::ir::CoreBinary* b) { EmitBinary(out, b); },          //
+                    [&](const core::ir::CoreUnary* u) { EmitUnary(out, u); },            //
+                    [&](const core::ir::Convert* b) { EmitConvert(out, b); },            //
+                    [&](const core::ir::Let* l) { out << NameOf(l->Result(0)); },        //
+                    [&](const core::ir::Load* l) { EmitLoad(out, l); },                  //
+                    [&](const core::ir::Construct* c) { EmitConstruct(out, c); },        //
+                    [&](const core::ir::Var* var) { out << NameOf(var->Result(0)); },    //
+                    [&](const core::ir::Bitcast* b) { EmitBitcast(out, b); },            //
+                    [&](const core::ir::Access* a) { EmitAccess(out, a); },              //
+                    [&](const msl::ir::BuiltinCall* c) { EmitMslBuiltinCall(out, c); },  //
+                    [&](const msl::ir::MemberBuiltinCall* c) {
+                        EmitMslMemberBuiltinCall(out, c);
+                    },                                                                         //
                     [&](const core::ir::CoreBuiltinCall* c) { EmitCoreBuiltinCall(out, c); },  //
                     [&](const core::ir::UserCall* c) { EmitUserCall(out, c); },                //
                     [&](const core::ir::LoadVectorElement* e) {
@@ -854,6 +858,20 @@
         out << ")";
     }
 
+    void EmitMslMemberBuiltinCall(StringStream& out, const msl::ir::MemberBuiltinCall* c) {
+        EmitValue(out, c->Object());
+        out << "." << c->Func() << "(";
+        bool needs_comma = false;
+        for (const auto* arg : c->Args()) {
+            if (needs_comma) {
+                out << ", ";
+            }
+            EmitAndTakeAddressIfNeeded(out, arg);
+            needs_comma = true;
+        }
+        out << ")";
+    }
+
     void EmitCoreBuiltinCall(StringStream& out, const core::ir::CoreBuiltinCall* c) {
         EmitCoreBuiltinName(out, c->Func());
         out << "(";
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
index f5c36f4..b5bbaf1 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
@@ -39,6 +39,7 @@
 #include "src/tint/lang/msl/barrier_type.h"
 #include "src/tint/lang/msl/builtin_fn.h"
 #include "src/tint/lang/msl/ir/builtin_call.h"
+#include "src/tint/lang/msl/ir/member_builtin_call.h"
 #include "src/tint/lang/msl/ir/memory_order.h"
 #include "src/tint/utils/containers/hashmap.h"
 
@@ -79,6 +80,7 @@
                     case core::BuiltinFn::kAtomicStore:
                     case core::BuiltinFn::kAtomicSub:
                     case core::BuiltinFn::kAtomicXor:
+                    case core::BuiltinFn::kTextureSample:
                     case core::BuiltinFn::kStorageBarrier:
                     case core::BuiltinFn::kWorkgroupBarrier:
                     case core::BuiltinFn::kTextureBarrier:
@@ -93,6 +95,7 @@
         // Replace the builtins that we found.
         for (auto* builtin : worklist) {
             switch (builtin->Func()) {
+                // Atomics.
                 case core::BuiltinFn::kAtomicAdd:
                     AtomicCall(builtin, msl::BuiltinFn::kAtomicFetchAddExplicit);
                     break;
@@ -126,6 +129,13 @@
                 case core::BuiltinFn::kAtomicXor:
                     AtomicCall(builtin, msl::BuiltinFn::kAtomicFetchXorExplicit);
                     break;
+
+                // Texture builtins.
+                case core::BuiltinFn::kTextureSample:
+                    TextureSample(builtin);
+                    break;
+
+                // Barriers.
                 case core::BuiltinFn::kStorageBarrier:
                     ThreadgroupBarrier(builtin, BarrierType::kDevice);
                     break;
@@ -135,6 +145,7 @@
                 case core::BuiltinFn::kTextureBarrier:
                     ThreadgroupBarrier(builtin, BarrierType::kTexture);
                     break;
+
                 default:
                     break;
             }
@@ -194,6 +205,17 @@
         builtin->Destroy();
     }
 
+    /// Replace a textureSample call with the equivalent MSL intrinsic.
+    /// @param builtin the builtin call instruction
+    void TextureSample(core::ir::CoreBuiltinCall* builtin) {
+        // The MSL intrinsic is a member function, so we split the first argument off as the object.
+        auto args = Vector<core::ir::Value*, 4>(builtin->Args().Offset(1));
+        auto* call = b.MemberCallWithResult<msl::ir::MemberBuiltinCall>(
+            builtin->DetachResult(), msl::BuiltinFn::kSample, builtin->Args()[0], std::move(args));
+        call->InsertBefore(builtin);
+        builtin->Destroy();
+    }
+
     /// Replace a barrier builtin with the `threadgroupBarrier()` intrinsic.
     /// @param builtin the builtin call instruction
     /// @param type the barrier type
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
index 527c7c7..904e1fa 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
@@ -35,6 +35,7 @@
 #include "src/tint/lang/core/number.h"
 #include "src/tint/lang/core/type/atomic.h"
 #include "src/tint/lang/core/type/builtin_structs.h"
+#include "src/tint/lang/core/type/sampled_texture.h"
 
 using namespace tint::core::fluent_types;     // NOLINT
 using namespace tint::core::number_suffixes;  // NOLINT
@@ -751,6 +752,42 @@
     EXPECT_EQ(expect, str());
 }
 
+TEST_F(MslWriter_BuiltinPolyfillTest, TextureSample) {
+    auto* t = b.FunctionParam(
+        "t", ty.Get<core::type::SampledTexture>(core::type::TextureDimension::k2d, ty.f32()));
+    auto* s = b.FunctionParam("s", ty.sampler());
+    auto* coords = b.FunctionParam("coords", ty.vec2<f32>());
+    auto* func = b.Function("foo", ty.vec4<f32>());
+    func->SetParams({t, s, coords});
+    b.Append(func->Block(), [&] {
+        auto* result = b.Call<vec4<f32>>(core::BuiltinFn::kTextureSample, t, s, coords);
+        b.Return(func, result);
+    });
+
+    auto* src = R"(
+%foo = func(%t:texture_2d<f32>, %s:sampler, %coords:vec2<f32>):vec4<f32> {
+  $B1: {
+    %5:vec4<f32> = textureSample %t, %s, %coords
+    ret %5
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%foo = func(%t:texture_2d<f32>, %s:sampler, %coords:vec2<f32>):vec4<f32> {
+  $B1: {
+    %5:vec4<f32> = %t.sample %s, %coords
+    ret %5
+  }
+}
+)";
+
+    Run(BuiltinPolyfill);
+
+    EXPECT_EQ(expect, str());
+}
+
 TEST_F(MslWriter_BuiltinPolyfillTest, WorkgroupBarrier) {
     auto* func = b.Function("foo", ty.void_());
     func->SetStage(core::ir::Function::PipelineStage::kCompute);
diff --git a/test/tint/bug/tint/1183.wgsl.expected.ir.msl b/test/tint/bug/tint/1183.wgsl.expected.ir.msl
index 3543919..c6e4754 100644
--- a/test/tint/bug/tint/1183.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1183.wgsl.expected.ir.msl
@@ -1,9 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> t;
+  sampler s;
+};
+struct f_outputs {
+  float4 tint_symbol [[color(0)]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
+float4 f_inner(tint_module_vars_struct tint_module_vars) {
+  return tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f), int2(4, 6));
+}
+fragment f_outputs f(texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+  return f_outputs{.tint_symbol=f_inner(tint_module_vars)};
+}
diff --git a/test/tint/bug/tint/978.wgsl.expected.ir.msl b/test/tint/bug/tint/978.wgsl.expected.ir.msl
index 414ed77..e9fbfff 100644
--- a/test/tint/bug/tint/978.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/978.wgsl.expected.ir.msl
@@ -1,40 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct FragmentOutput {
+  float4 color;
+};
+struct FragmentInput {
+  float2 vUv;
+};
+struct tint_module_vars_struct {
+  depth2d<float, access::sample> depthMap;
+  sampler texSampler;
+};
+struct tint_symbol_outputs {
+  float4 FragmentOutput_color [[color(0)]];
+};
+struct tint_symbol_inputs {
+  float2 FragmentInput_vUv [[user(locn2)]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:1257 internal compiler error: FragmentOutput = struct @align(16) {
-  color:vec4<f32> @offset(0), @location(0)
+FragmentOutput tint_symbol_inner(FragmentInput fIn, tint_module_vars_struct tint_module_vars) {
+  float const sample = tint_module_vars.depthMap.sample(tint_module_vars.texSampler, fIn.vUv);
+  float3 const color = float3(sample, sample, sample);
+  FragmentOutput fOut = {};
+  fOut.color = float4(color, 1.0f);
+  return fOut;
 }
-
-FragmentInput = struct @align(8) {
-  vUv:vec2<f32> @offset(0), @location(2)
+fragment tint_symbol_outputs tint_symbol(tint_symbol_inputs inputs [[stage_in]], depth2d<float, access::sample> depthMap [[texture(0)]], sampler texSampler [[sampler(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.depthMap=depthMap, .texSampler=texSampler};
+  return tint_symbol_outputs{.FragmentOutput_color=tint_symbol_inner(FragmentInput{.vUv=inputs.FragmentInput_vUv}, tint_module_vars).color};
 }
-
-$B1: {  # root
-  %depthMap:ptr<handle, texture_depth_2d, read> = var @binding_point(1, 5)
-  %texSampler:ptr<handle, sampler, read> = var @binding_point(1, 3)
-}
-
-%tint_symbol = @fragment func(%fIn:FragmentInput):FragmentOutput {
-  $B2: {
-    %5:texture_depth_2d = load %depthMap
-    %6:sampler = load %texSampler
-    %7:vec2<f32> = access %fIn, 0u
-    %8:f32 = textureSample %5, %6, %7
-    %sample:f32 = let %8
-    %10:vec3<f32> = construct %sample, %sample, %sample
-    %color:vec3<f32> = let %10
-    %fOut:ptr<function, FragmentOutput, read_write> = var
-    %13:ptr<function, vec4<f32>, read_write> = access %fOut, 0u
-    %14:vec4<f32> = construct %color, 1.0f
-    store %13, %14
-    %15:FragmentOutput = load %fOut
-    ret %15
-  }
-}
-
-invalid entry point IO struct uses
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/0dff6c.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/0dff6c.wgsl.expected.ir.msl
index 55819ad..adce6d3 100644
--- a/test/tint/builtins/gen/literal/textureSample/0dff6c.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/0dff6c.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depth2d<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_2d, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_0dff6c(tint_module_vars_struct tint_module_vars) {
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), int2(1));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_0dff6c = func():void {
-  $B2: {
-    %5:texture_depth_2d = load %arg_0
-    %6:sampler = load %arg_1
-    %7:f32 = textureSample %5, %6, vec2<f32>(1.0f), vec2<i32>(1i)
-    %res:ptr<function, f32, read_write> = var, %7
-    %9:f32 = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(depth2d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_0dff6c(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_0dff6c
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/17e988.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/17e988.wgsl.expected.ir.msl
index 1da17ad..754c9df 100644
--- a/test/tint/builtins/gen/literal/textureSample/17e988.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/17e988.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d_array<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_17e988(tint_module_vars_struct tint_module_vars) {
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), 1, int2(1));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_17e988 = func():void {
-  $B2: {
-    %5:texture_2d_array<f32> = load %arg_0
-    %6:sampler = load %arg_1
-    %7:vec4<f32> = textureSample %5, %6, vec2<f32>(1.0f), 1i, vec2<i32>(1i)
-    %res:ptr<function, vec4<f32>, read_write> = var, %7
-    %9:vec4<f32> = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_17e988(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_17e988
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/193203.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/193203.wgsl.expected.ir.msl
index ae6b1e6..7441785 100644
--- a/test/tint/builtins/gen/literal/textureSample/193203.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/193203.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d_array<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_193203(tint_module_vars_struct tint_module_vars) {
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), 1u, int2(1));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_193203 = func():void {
-  $B2: {
-    %5:texture_2d_array<f32> = load %arg_0
-    %6:sampler = load %arg_1
-    %7:vec4<f32> = textureSample %5, %6, vec2<f32>(1.0f), 1u, vec2<i32>(1i)
-    %res:ptr<function, vec4<f32>, read_write> = var, %7
-    %9:vec4<f32> = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_193203(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_193203
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/1a4e1b.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/1a4e1b.wgsl.expected.ir.msl
index a0e7fe3..4d76dd5 100644
--- a/test/tint/builtins/gen/literal/textureSample/1a4e1b.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/1a4e1b.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depth2d_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_2d_array, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_1a4e1b(tint_module_vars_struct tint_module_vars) {
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), 1u);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_1a4e1b = func():void {
-  $B2: {
-    %5:texture_depth_2d_array = load %arg_0
-    %6:sampler = load %arg_1
-    %7:f32 = textureSample %5, %6, vec2<f32>(1.0f), 1u
-    %res:ptr<function, f32, read_write> = var, %7
-    %9:f32 = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_1a4e1b(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_1a4e1b
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/2149ec.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/2149ec.wgsl.expected.ir.msl
index 8b191d3..fa08bbe 100644
--- a/test/tint/builtins/gen/literal/textureSample/2149ec.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/2149ec.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture3d<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_3d<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_2149ec(tint_module_vars_struct tint_module_vars) {
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float3(1.0f), int3(1));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_2149ec = func():void {
-  $B2: {
-    %5:texture_3d<f32> = load %arg_0
-    %6:sampler = load %arg_1
-    %7:vec4<f32> = textureSample %5, %6, vec3<f32>(1.0f), vec3<i32>(1i)
-    %res:ptr<function, vec4<f32>, read_write> = var, %7
-    %9:vec4<f32> = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(texture3d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_2149ec(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_2149ec
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/38bbb9.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/38bbb9.wgsl.expected.ir.msl
index a4a4204..f5743b7 100644
--- a/test/tint/builtins/gen/literal/textureSample/38bbb9.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/38bbb9.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depth2d<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_2d, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_38bbb9(tint_module_vars_struct tint_module_vars) {
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_38bbb9 = func():void {
-  $B2: {
-    %5:texture_depth_2d = load %arg_0
-    %6:sampler = load %arg_1
-    %7:f32 = textureSample %5, %6, vec2<f32>(1.0f)
-    %res:ptr<function, f32, read_write> = var, %7
-    %9:f32 = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(depth2d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_38bbb9(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_38bbb9
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/3b50bd.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/3b50bd.wgsl.expected.ir.msl
index 3f682f9..bd92d77 100644
--- a/test/tint/builtins/gen/literal/textureSample/3b50bd.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/3b50bd.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture3d<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_3d<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_3b50bd(tint_module_vars_struct tint_module_vars) {
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float3(1.0f));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_3b50bd = func():void {
-  $B2: {
-    %5:texture_3d<f32> = load %arg_0
-    %6:sampler = load %arg_1
-    %7:vec4<f32> = textureSample %5, %6, vec3<f32>(1.0f)
-    %res:ptr<function, vec4<f32>, read_write> = var, %7
-    %9:vec4<f32> = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(texture3d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_3b50bd(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_3b50bd
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/4703d0.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/4703d0.wgsl.expected.ir.msl
index 979bf90..83b2790 100644
--- a/test/tint/builtins/gen/literal/textureSample/4703d0.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/4703d0.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depth2d_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_2d_array, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_4703d0(tint_module_vars_struct tint_module_vars) {
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), 1u, int2(1));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_4703d0 = func():void {
-  $B2: {
-    %5:texture_depth_2d_array = load %arg_0
-    %6:sampler = load %arg_1
-    %7:f32 = textureSample %5, %6, vec2<f32>(1.0f), 1u, vec2<i32>(1i)
-    %res:ptr<function, f32, read_write> = var, %7
-    %9:f32 = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_4703d0(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_4703d0
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/4dd1bf.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/4dd1bf.wgsl.expected.ir.msl
index 9ebb5c2..b7252ad 100644
--- a/test/tint/builtins/gen/literal/textureSample/4dd1bf.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/4dd1bf.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_cube_array<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_4dd1bf(tint_module_vars_struct tint_module_vars) {
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float3(1.0f), 1);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_4dd1bf = func():void {
-  $B2: {
-    %5:texture_cube_array<f32> = load %arg_0
-    %6:sampler = load %arg_1
-    %7:vec4<f32> = textureSample %5, %6, vec3<f32>(1.0f), 1i
-    %res:ptr<function, vec4<f32>, read_write> = var, %7
-    %9:vec4<f32> = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(texturecube_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_4dd1bf(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_4dd1bf
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/51b514.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/51b514.wgsl.expected.ir.msl
index e019e2b..850b230 100644
--- a/test/tint/builtins/gen/literal/textureSample/51b514.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/51b514.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_51b514(tint_module_vars_struct tint_module_vars) {
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_51b514 = func():void {
-  $B2: {
-    %5:texture_2d<f32> = load %arg_0
-    %6:sampler = load %arg_1
-    %7:vec4<f32> = textureSample %5, %6, vec2<f32>(1.0f)
-    %res:ptr<function, vec4<f32>, read_write> = var, %7
-    %9:vec4<f32> = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(texture2d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_51b514(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_51b514
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/60bf45.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/60bf45.wgsl.expected.ir.msl
index 03697ed..90a9283 100644
--- a/test/tint/builtins/gen/literal/textureSample/60bf45.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/60bf45.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depth2d_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_2d_array, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_60bf45(tint_module_vars_struct tint_module_vars) {
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), 1, int2(1));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_60bf45 = func():void {
-  $B2: {
-    %5:texture_depth_2d_array = load %arg_0
-    %6:sampler = load %arg_1
-    %7:f32 = textureSample %5, %6, vec2<f32>(1.0f), 1i, vec2<i32>(1i)
-    %res:ptr<function, f32, read_write> = var, %7
-    %9:f32 = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_60bf45(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_60bf45
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/6717ca.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/6717ca.wgsl.expected.ir.msl
index c5d638e..34f57aa 100644
--- a/test/tint/builtins/gen/literal/textureSample/6717ca.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/6717ca.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d_array<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_6717ca(tint_module_vars_struct tint_module_vars) {
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), 1);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_6717ca = func():void {
-  $B2: {
-    %5:texture_2d_array<f32> = load %arg_0
-    %6:sampler = load %arg_1
-    %7:vec4<f32> = textureSample %5, %6, vec2<f32>(1.0f), 1i
-    %res:ptr<function, vec4<f32>, read_write> = var, %7
-    %9:vec4<f32> = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_6717ca(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_6717ca
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/6e64fb.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/6e64fb.wgsl.expected.ir.msl
index a3c0665..971c6fb 100644
--- a/test/tint/builtins/gen/literal/textureSample/6e64fb.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/6e64fb.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture1d<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_1d<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_6e64fb(tint_module_vars_struct tint_module_vars) {
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, 1.0f);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_6e64fb = func():void {
-  $B2: {
-    %5:texture_1d<f32> = load %arg_0
-    %6:sampler = load %arg_1
-    %7:vec4<f32> = textureSample %5, %6, 1.0f
-    %res:ptr<function, vec4<f32>, read_write> = var, %7
-    %9:vec4<f32> = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(texture1d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_6e64fb(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_6e64fb
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/7e9ffd.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/7e9ffd.wgsl.expected.ir.msl
index 203f61c..ccd1290 100644
--- a/test/tint/builtins/gen/literal/textureSample/7e9ffd.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/7e9ffd.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depth2d_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_2d_array, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_7e9ffd(tint_module_vars_struct tint_module_vars) {
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), 1);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_7e9ffd = func():void {
-  $B2: {
-    %5:texture_depth_2d_array = load %arg_0
-    %6:sampler = load %arg_1
-    %7:f32 = textureSample %5, %6, vec2<f32>(1.0f), 1i
-    %res:ptr<function, f32, read_write> = var, %7
-    %9:f32 = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_7e9ffd(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_7e9ffd
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/7fd8cb.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/7fd8cb.wgsl.expected.ir.msl
index f6d8f84..b24c4de 100644
--- a/test/tint/builtins/gen/literal/textureSample/7fd8cb.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/7fd8cb.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depthcube_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_cube_array, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_7fd8cb(tint_module_vars_struct tint_module_vars) {
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float3(1.0f), 1u);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_7fd8cb = func():void {
-  $B2: {
-    %5:texture_depth_cube_array = load %arg_0
-    %6:sampler = load %arg_1
-    %7:f32 = textureSample %5, %6, vec3<f32>(1.0f), 1u
-    %res:ptr<function, f32, read_write> = var, %7
-    %9:f32 = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(depthcube_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_7fd8cb(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_7fd8cb
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/85c4ba.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/85c4ba.wgsl.expected.ir.msl
index df38079..38b8474 100644
--- a/test/tint/builtins/gen/literal/textureSample/85c4ba.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/85c4ba.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_85c4ba(tint_module_vars_struct tint_module_vars) {
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), int2(1));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_85c4ba = func():void {
-  $B2: {
-    %5:texture_2d<f32> = load %arg_0
-    %6:sampler = load %arg_1
-    %7:vec4<f32> = textureSample %5, %6, vec2<f32>(1.0f), vec2<i32>(1i)
-    %res:ptr<function, vec4<f32>, read_write> = var, %7
-    %9:vec4<f32> = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(texture2d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_85c4ba(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_85c4ba
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/bc7477.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/bc7477.wgsl.expected.ir.msl
index 63be6d6..dfeb0c3 100644
--- a/test/tint/builtins/gen/literal/textureSample/bc7477.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/bc7477.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_cube_array<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_bc7477(tint_module_vars_struct tint_module_vars) {
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float3(1.0f), 1u);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_bc7477 = func():void {
-  $B2: {
-    %5:texture_cube_array<f32> = load %arg_0
-    %6:sampler = load %arg_1
-    %7:vec4<f32> = textureSample %5, %6, vec3<f32>(1.0f), 1u
-    %res:ptr<function, vec4<f32>, read_write> = var, %7
-    %9:vec4<f32> = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(texturecube_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_bc7477(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_bc7477
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/c2f4e8.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/c2f4e8.wgsl.expected.ir.msl
index cc9e0ac..5d3648c 100644
--- a/test/tint/builtins/gen/literal/textureSample/c2f4e8.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/c2f4e8.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depthcube_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_cube_array, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_c2f4e8(tint_module_vars_struct tint_module_vars) {
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float3(1.0f), 1);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_c2f4e8 = func():void {
-  $B2: {
-    %5:texture_depth_cube_array = load %arg_0
-    %6:sampler = load %arg_1
-    %7:f32 = textureSample %5, %6, vec3<f32>(1.0f), 1i
-    %res:ptr<function, f32, read_write> = var, %7
-    %9:f32 = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(depthcube_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_c2f4e8(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_c2f4e8
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/d6b281.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/d6b281.wgsl.expected.ir.msl
index d663331..c729426 100644
--- a/test/tint/builtins/gen/literal/textureSample/d6b281.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/d6b281.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d_array<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_d6b281(tint_module_vars_struct tint_module_vars) {
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float2(1.0f), 1u);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_d6b281 = func():void {
-  $B2: {
-    %5:texture_2d_array<f32> = load %arg_0
-    %6:sampler = load %arg_1
-    %7:vec4<f32> = textureSample %5, %6, vec2<f32>(1.0f), 1u
-    %res:ptr<function, vec4<f32>, read_write> = var, %7
-    %9:vec4<f32> = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_d6b281(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_d6b281
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/e53267.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/e53267.wgsl.expected.ir.msl
index 3babe54..6be1416 100644
--- a/test/tint/builtins/gen/literal/textureSample/e53267.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/e53267.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_cube<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_e53267(tint_module_vars_struct tint_module_vars) {
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float3(1.0f));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_e53267 = func():void {
-  $B2: {
-    %5:texture_cube<f32> = load %arg_0
-    %6:sampler = load %arg_1
-    %7:vec4<f32> = textureSample %5, %6, vec3<f32>(1.0f)
-    %res:ptr<function, vec4<f32>, read_write> = var, %7
-    %9:vec4<f32> = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(texturecube<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_e53267(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_e53267
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/literal/textureSample/ea7030.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureSample/ea7030.wgsl.expected.ir.msl
index 62669b6..db29db5 100644
--- a/test/tint/builtins/gen/literal/textureSample/ea7030.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureSample/ea7030.wgsl.expected.ir.msl
@@ -1,33 +1,16 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depthcube<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_cube, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_ea7030(tint_module_vars_struct tint_module_vars) {
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, float3(1.0f));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_ea7030 = func():void {
-  $B2: {
-    %5:texture_depth_cube = load %arg_0
-    %6:sampler = load %arg_1
-    %7:f32 = textureSample %5, %6, vec3<f32>(1.0f)
-    %res:ptr<function, f32, read_write> = var, %7
-    %9:f32 = load %res
-    store %prevent_dce, %9
-    ret
-  }
+fragment void fragment_main(depthcube<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_ea7030(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %11:void = call %textureSample_ea7030
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/0dff6c.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/0dff6c.wgsl.expected.ir.msl
index 3c40846..c730055 100644
--- a/test/tint/builtins/gen/var/textureSample/0dff6c.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/0dff6c.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depth2d<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_2d, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_0dff6c(tint_module_vars_struct tint_module_vars) {
+  float2 arg_2 = float2(1.0f);
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, int2(1));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_0dff6c = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
-    %6:texture_depth_2d = load %arg_0
-    %7:sampler = load %arg_1
-    %8:vec2<f32> = load %arg_2
-    %9:f32 = textureSample %6, %7, %8, vec2<i32>(1i)
-    %res:ptr<function, f32, read_write> = var, %9
-    %11:f32 = load %res
-    store %prevent_dce, %11
-    ret
-  }
+fragment void fragment_main(depth2d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_0dff6c(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %13:void = call %textureSample_0dff6c
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/17e988.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/17e988.wgsl.expected.ir.msl
index 5f62935..759875c9 100644
--- a/test/tint/builtins/gen/var/textureSample/17e988.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/17e988.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d_array<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_17e988(tint_module_vars_struct tint_module_vars) {
+  float2 arg_2 = float2(1.0f);
+  int arg_3 = 1;
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3, int2(1));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_17e988 = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
-    %arg_3:ptr<function, i32, read_write> = var, 1i
-    %7:texture_2d_array<f32> = load %arg_0
-    %8:sampler = load %arg_1
-    %9:vec2<f32> = load %arg_2
-    %10:i32 = load %arg_3
-    %11:vec4<f32> = textureSample %7, %8, %9, %10, vec2<i32>(1i)
-    %res:ptr<function, vec4<f32>, read_write> = var, %11
-    %13:vec4<f32> = load %res
-    store %prevent_dce, %13
-    ret
-  }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_17e988(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %15:void = call %textureSample_17e988
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/193203.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/193203.wgsl.expected.ir.msl
index b0393af..08d4cf4 100644
--- a/test/tint/builtins/gen/var/textureSample/193203.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/193203.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d_array<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_193203(tint_module_vars_struct tint_module_vars) {
+  float2 arg_2 = float2(1.0f);
+  uint arg_3 = 1u;
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3, int2(1));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_193203 = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
-    %arg_3:ptr<function, u32, read_write> = var, 1u
-    %7:texture_2d_array<f32> = load %arg_0
-    %8:sampler = load %arg_1
-    %9:vec2<f32> = load %arg_2
-    %10:u32 = load %arg_3
-    %11:vec4<f32> = textureSample %7, %8, %9, %10, vec2<i32>(1i)
-    %res:ptr<function, vec4<f32>, read_write> = var, %11
-    %13:vec4<f32> = load %res
-    store %prevent_dce, %13
-    ret
-  }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_193203(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %15:void = call %textureSample_193203
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/1a4e1b.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/1a4e1b.wgsl.expected.ir.msl
index ddca6b6..0a84e65 100644
--- a/test/tint/builtins/gen/var/textureSample/1a4e1b.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/1a4e1b.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depth2d_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_2d_array, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_1a4e1b(tint_module_vars_struct tint_module_vars) {
+  float2 arg_2 = float2(1.0f);
+  uint arg_3 = 1u;
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_1a4e1b = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
-    %arg_3:ptr<function, u32, read_write> = var, 1u
-    %7:texture_depth_2d_array = load %arg_0
-    %8:sampler = load %arg_1
-    %9:vec2<f32> = load %arg_2
-    %10:u32 = load %arg_3
-    %11:f32 = textureSample %7, %8, %9, %10
-    %res:ptr<function, f32, read_write> = var, %11
-    %13:f32 = load %res
-    store %prevent_dce, %13
-    ret
-  }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_1a4e1b(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %15:void = call %textureSample_1a4e1b
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/2149ec.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/2149ec.wgsl.expected.ir.msl
index 4ad8af4..f25e8dd 100644
--- a/test/tint/builtins/gen/var/textureSample/2149ec.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/2149ec.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture3d<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_3d<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_2149ec(tint_module_vars_struct tint_module_vars) {
+  float3 arg_2 = float3(1.0f);
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, int3(1));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_2149ec = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
-    %6:texture_3d<f32> = load %arg_0
-    %7:sampler = load %arg_1
-    %8:vec3<f32> = load %arg_2
-    %9:vec4<f32> = textureSample %6, %7, %8, vec3<i32>(1i)
-    %res:ptr<function, vec4<f32>, read_write> = var, %9
-    %11:vec4<f32> = load %res
-    store %prevent_dce, %11
-    ret
-  }
+fragment void fragment_main(texture3d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_2149ec(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %13:void = call %textureSample_2149ec
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/38bbb9.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/38bbb9.wgsl.expected.ir.msl
index a113cd4..dc1f22e 100644
--- a/test/tint/builtins/gen/var/textureSample/38bbb9.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/38bbb9.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depth2d<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_2d, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_38bbb9(tint_module_vars_struct tint_module_vars) {
+  float2 arg_2 = float2(1.0f);
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_38bbb9 = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
-    %6:texture_depth_2d = load %arg_0
-    %7:sampler = load %arg_1
-    %8:vec2<f32> = load %arg_2
-    %9:f32 = textureSample %6, %7, %8
-    %res:ptr<function, f32, read_write> = var, %9
-    %11:f32 = load %res
-    store %prevent_dce, %11
-    ret
-  }
+fragment void fragment_main(depth2d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_38bbb9(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %13:void = call %textureSample_38bbb9
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/3b50bd.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/3b50bd.wgsl.expected.ir.msl
index 6e29909..7d4dccf 100644
--- a/test/tint/builtins/gen/var/textureSample/3b50bd.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/3b50bd.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture3d<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_3d<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_3b50bd(tint_module_vars_struct tint_module_vars) {
+  float3 arg_2 = float3(1.0f);
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_3b50bd = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
-    %6:texture_3d<f32> = load %arg_0
-    %7:sampler = load %arg_1
-    %8:vec3<f32> = load %arg_2
-    %9:vec4<f32> = textureSample %6, %7, %8
-    %res:ptr<function, vec4<f32>, read_write> = var, %9
-    %11:vec4<f32> = load %res
-    store %prevent_dce, %11
-    ret
-  }
+fragment void fragment_main(texture3d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_3b50bd(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %13:void = call %textureSample_3b50bd
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/4703d0.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/4703d0.wgsl.expected.ir.msl
index 17b17fa..8713856 100644
--- a/test/tint/builtins/gen/var/textureSample/4703d0.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/4703d0.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depth2d_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_2d_array, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_4703d0(tint_module_vars_struct tint_module_vars) {
+  float2 arg_2 = float2(1.0f);
+  uint arg_3 = 1u;
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3, int2(1));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_4703d0 = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
-    %arg_3:ptr<function, u32, read_write> = var, 1u
-    %7:texture_depth_2d_array = load %arg_0
-    %8:sampler = load %arg_1
-    %9:vec2<f32> = load %arg_2
-    %10:u32 = load %arg_3
-    %11:f32 = textureSample %7, %8, %9, %10, vec2<i32>(1i)
-    %res:ptr<function, f32, read_write> = var, %11
-    %13:f32 = load %res
-    store %prevent_dce, %13
-    ret
-  }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_4703d0(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %15:void = call %textureSample_4703d0
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/4dd1bf.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/4dd1bf.wgsl.expected.ir.msl
index b021b59..78bf83a 100644
--- a/test/tint/builtins/gen/var/textureSample/4dd1bf.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/4dd1bf.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_cube_array<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_4dd1bf(tint_module_vars_struct tint_module_vars) {
+  float3 arg_2 = float3(1.0f);
+  int arg_3 = 1;
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_4dd1bf = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
-    %arg_3:ptr<function, i32, read_write> = var, 1i
-    %7:texture_cube_array<f32> = load %arg_0
-    %8:sampler = load %arg_1
-    %9:vec3<f32> = load %arg_2
-    %10:i32 = load %arg_3
-    %11:vec4<f32> = textureSample %7, %8, %9, %10
-    %res:ptr<function, vec4<f32>, read_write> = var, %11
-    %13:vec4<f32> = load %res
-    store %prevent_dce, %13
-    ret
-  }
+fragment void fragment_main(texturecube_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_4dd1bf(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %15:void = call %textureSample_4dd1bf
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/51b514.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/51b514.wgsl.expected.ir.msl
index aa66ed2..f3bd051 100644
--- a/test/tint/builtins/gen/var/textureSample/51b514.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/51b514.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_51b514(tint_module_vars_struct tint_module_vars) {
+  float2 arg_2 = float2(1.0f);
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_51b514 = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
-    %6:texture_2d<f32> = load %arg_0
-    %7:sampler = load %arg_1
-    %8:vec2<f32> = load %arg_2
-    %9:vec4<f32> = textureSample %6, %7, %8
-    %res:ptr<function, vec4<f32>, read_write> = var, %9
-    %11:vec4<f32> = load %res
-    store %prevent_dce, %11
-    ret
-  }
+fragment void fragment_main(texture2d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_51b514(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %13:void = call %textureSample_51b514
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/60bf45.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/60bf45.wgsl.expected.ir.msl
index 31b95cf..55c40ab 100644
--- a/test/tint/builtins/gen/var/textureSample/60bf45.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/60bf45.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depth2d_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_2d_array, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_60bf45(tint_module_vars_struct tint_module_vars) {
+  float2 arg_2 = float2(1.0f);
+  int arg_3 = 1;
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3, int2(1));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_60bf45 = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
-    %arg_3:ptr<function, i32, read_write> = var, 1i
-    %7:texture_depth_2d_array = load %arg_0
-    %8:sampler = load %arg_1
-    %9:vec2<f32> = load %arg_2
-    %10:i32 = load %arg_3
-    %11:f32 = textureSample %7, %8, %9, %10, vec2<i32>(1i)
-    %res:ptr<function, f32, read_write> = var, %11
-    %13:f32 = load %res
-    store %prevent_dce, %13
-    ret
-  }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_60bf45(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %15:void = call %textureSample_60bf45
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/6717ca.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/6717ca.wgsl.expected.ir.msl
index 3adc780..199640d 100644
--- a/test/tint/builtins/gen/var/textureSample/6717ca.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/6717ca.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d_array<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_6717ca(tint_module_vars_struct tint_module_vars) {
+  float2 arg_2 = float2(1.0f);
+  int arg_3 = 1;
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_6717ca = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
-    %arg_3:ptr<function, i32, read_write> = var, 1i
-    %7:texture_2d_array<f32> = load %arg_0
-    %8:sampler = load %arg_1
-    %9:vec2<f32> = load %arg_2
-    %10:i32 = load %arg_3
-    %11:vec4<f32> = textureSample %7, %8, %9, %10
-    %res:ptr<function, vec4<f32>, read_write> = var, %11
-    %13:vec4<f32> = load %res
-    store %prevent_dce, %13
-    ret
-  }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_6717ca(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %15:void = call %textureSample_6717ca
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/6e64fb.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/6e64fb.wgsl.expected.ir.msl
index 2d7a629..2de338d 100644
--- a/test/tint/builtins/gen/var/textureSample/6e64fb.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/6e64fb.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture1d<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_1d<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_6e64fb(tint_module_vars_struct tint_module_vars) {
+  float arg_2 = 1.0f;
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_6e64fb = func():void {
-  $B2: {
-    %arg_2:ptr<function, f32, read_write> = var, 1.0f
-    %6:texture_1d<f32> = load %arg_0
-    %7:sampler = load %arg_1
-    %8:f32 = load %arg_2
-    %9:vec4<f32> = textureSample %6, %7, %8
-    %res:ptr<function, vec4<f32>, read_write> = var, %9
-    %11:vec4<f32> = load %res
-    store %prevent_dce, %11
-    ret
-  }
+fragment void fragment_main(texture1d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_6e64fb(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %13:void = call %textureSample_6e64fb
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/7e9ffd.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/7e9ffd.wgsl.expected.ir.msl
index 2ff2807..d26d806 100644
--- a/test/tint/builtins/gen/var/textureSample/7e9ffd.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/7e9ffd.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depth2d_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_2d_array, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_7e9ffd(tint_module_vars_struct tint_module_vars) {
+  float2 arg_2 = float2(1.0f);
+  int arg_3 = 1;
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_7e9ffd = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
-    %arg_3:ptr<function, i32, read_write> = var, 1i
-    %7:texture_depth_2d_array = load %arg_0
-    %8:sampler = load %arg_1
-    %9:vec2<f32> = load %arg_2
-    %10:i32 = load %arg_3
-    %11:f32 = textureSample %7, %8, %9, %10
-    %res:ptr<function, f32, read_write> = var, %11
-    %13:f32 = load %res
-    store %prevent_dce, %13
-    ret
-  }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_7e9ffd(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %15:void = call %textureSample_7e9ffd
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/7fd8cb.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/7fd8cb.wgsl.expected.ir.msl
index cba8a48..d2fde33 100644
--- a/test/tint/builtins/gen/var/textureSample/7fd8cb.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/7fd8cb.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depthcube_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_cube_array, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_7fd8cb(tint_module_vars_struct tint_module_vars) {
+  float3 arg_2 = float3(1.0f);
+  uint arg_3 = 1u;
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_7fd8cb = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
-    %arg_3:ptr<function, u32, read_write> = var, 1u
-    %7:texture_depth_cube_array = load %arg_0
-    %8:sampler = load %arg_1
-    %9:vec3<f32> = load %arg_2
-    %10:u32 = load %arg_3
-    %11:f32 = textureSample %7, %8, %9, %10
-    %res:ptr<function, f32, read_write> = var, %11
-    %13:f32 = load %res
-    store %prevent_dce, %13
-    ret
-  }
+fragment void fragment_main(depthcube_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_7fd8cb(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %15:void = call %textureSample_7fd8cb
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/85c4ba.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/85c4ba.wgsl.expected.ir.msl
index 4515efb..8eed4cb 100644
--- a/test/tint/builtins/gen/var/textureSample/85c4ba.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/85c4ba.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_85c4ba(tint_module_vars_struct tint_module_vars) {
+  float2 arg_2 = float2(1.0f);
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, int2(1));
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_85c4ba = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
-    %6:texture_2d<f32> = load %arg_0
-    %7:sampler = load %arg_1
-    %8:vec2<f32> = load %arg_2
-    %9:vec4<f32> = textureSample %6, %7, %8, vec2<i32>(1i)
-    %res:ptr<function, vec4<f32>, read_write> = var, %9
-    %11:vec4<f32> = load %res
-    store %prevent_dce, %11
-    ret
-  }
+fragment void fragment_main(texture2d<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_85c4ba(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %13:void = call %textureSample_85c4ba
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/bc7477.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/bc7477.wgsl.expected.ir.msl
index 6f235ed..1344e67 100644
--- a/test/tint/builtins/gen/var/textureSample/bc7477.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/bc7477.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_cube_array<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_bc7477(tint_module_vars_struct tint_module_vars) {
+  float3 arg_2 = float3(1.0f);
+  uint arg_3 = 1u;
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_bc7477 = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
-    %arg_3:ptr<function, u32, read_write> = var, 1u
-    %7:texture_cube_array<f32> = load %arg_0
-    %8:sampler = load %arg_1
-    %9:vec3<f32> = load %arg_2
-    %10:u32 = load %arg_3
-    %11:vec4<f32> = textureSample %7, %8, %9, %10
-    %res:ptr<function, vec4<f32>, read_write> = var, %11
-    %13:vec4<f32> = load %res
-    store %prevent_dce, %13
-    ret
-  }
+fragment void fragment_main(texturecube_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_bc7477(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %15:void = call %textureSample_bc7477
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/c2f4e8.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/c2f4e8.wgsl.expected.ir.msl
index 5e96887..ac62210 100644
--- a/test/tint/builtins/gen/var/textureSample/c2f4e8.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/c2f4e8.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depthcube_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_cube_array, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_c2f4e8(tint_module_vars_struct tint_module_vars) {
+  float3 arg_2 = float3(1.0f);
+  int arg_3 = 1;
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_c2f4e8 = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
-    %arg_3:ptr<function, i32, read_write> = var, 1i
-    %7:texture_depth_cube_array = load %arg_0
-    %8:sampler = load %arg_1
-    %9:vec3<f32> = load %arg_2
-    %10:i32 = load %arg_3
-    %11:f32 = textureSample %7, %8, %9, %10
-    %res:ptr<function, f32, read_write> = var, %11
-    %13:f32 = load %res
-    store %prevent_dce, %13
-    ret
-  }
+fragment void fragment_main(depthcube_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_c2f4e8(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %15:void = call %textureSample_c2f4e8
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/d6b281.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/d6b281.wgsl.expected.ir.msl
index 465b0e7..1f78789 100644
--- a/test/tint/builtins/gen/var/textureSample/d6b281.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/d6b281.wgsl.expected.ir.msl
@@ -1,37 +1,18 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d_array<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d_array<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_d6b281(tint_module_vars_struct tint_module_vars) {
+  float2 arg_2 = float2(1.0f);
+  uint arg_3 = 1u;
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2, arg_3);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_d6b281 = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec2<f32>, read_write> = var, vec2<f32>(1.0f)
-    %arg_3:ptr<function, u32, read_write> = var, 1u
-    %7:texture_2d_array<f32> = load %arg_0
-    %8:sampler = load %arg_1
-    %9:vec2<f32> = load %arg_2
-    %10:u32 = load %arg_3
-    %11:vec4<f32> = textureSample %7, %8, %9, %10
-    %res:ptr<function, vec4<f32>, read_write> = var, %11
-    %13:vec4<f32> = load %res
-    store %prevent_dce, %13
-    ret
-  }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_d6b281(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %15:void = call %textureSample_d6b281
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/e53267.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/e53267.wgsl.expected.ir.msl
index 8c1cd90..a3053c5 100644
--- a/test/tint/builtins/gen/var/textureSample/e53267.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/e53267.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube<float, access::sample> arg_0;
+  sampler arg_1;
+  device float4* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_cube<f32>, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, vec4<f32>, read_write> = var @binding_point(2, 0)
+void textureSample_e53267(tint_module_vars_struct tint_module_vars) {
+  float3 arg_2 = float3(1.0f);
+  float4 res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_e53267 = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
-    %6:texture_cube<f32> = load %arg_0
-    %7:sampler = load %arg_1
-    %8:vec3<f32> = load %arg_2
-    %9:vec4<f32> = textureSample %6, %7, %8
-    %res:ptr<function, vec4<f32>, read_write> = var, %9
-    %11:vec4<f32> = load %res
-    store %prevent_dce, %11
-    ret
-  }
+fragment void fragment_main(texturecube<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float4* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_e53267(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %13:void = call %textureSample_e53267
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/builtins/gen/var/textureSample/ea7030.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureSample/ea7030.wgsl.expected.ir.msl
index 0962884..964f1a1 100644
--- a/test/tint/builtins/gen/var/textureSample/ea7030.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureSample/ea7030.wgsl.expected.ir.msl
@@ -1,35 +1,17 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depthcube<float, access::sample> arg_0;
+  sampler arg_1;
+  device float* prevent_dce;
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_depth_cube, read> = var @binding_point(1, 0)
-  %arg_1:ptr<handle, sampler, read> = var @binding_point(1, 1)
-  %prevent_dce:ptr<storage, f32, read_write> = var @binding_point(2, 0)
+void textureSample_ea7030(tint_module_vars_struct tint_module_vars) {
+  float3 arg_2 = float3(1.0f);
+  float res = tint_module_vars.arg_0.sample(tint_module_vars.arg_1, arg_2);
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureSample_ea7030 = func():void {
-  $B2: {
-    %arg_2:ptr<function, vec3<f32>, read_write> = var, vec3<f32>(1.0f)
-    %6:texture_depth_cube = load %arg_0
-    %7:sampler = load %arg_1
-    %8:vec3<f32> = load %arg_2
-    %9:f32 = textureSample %6, %7, %8
-    %res:ptr<function, f32, read_write> = var, %9
-    %11:f32 = load %res
-    store %prevent_dce, %11
-    ret
-  }
+fragment void fragment_main(depthcube<float, access::sample> arg_0 [[texture(0)]], sampler arg_1 [[sampler(0)]], device float* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .arg_1=arg_1, .prevent_dce=prevent_dce};
+  textureSample_ea7030(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B3: {
-    %13:void = call %textureSample_ea7030
-    ret
-  }
-}
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/diagnostic_filtering/case_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/case_body_attribute.wgsl.expected.ir.msl
index 14f9fde..61c2299 100644
--- a/test/tint/diagnostic_filtering/case_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/case_body_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
 <dawn>/test/tint/diagnostic_filtering/case_body_attribute.wgsl:8:11 warning: 'textureSample' must only be called from uniform control flow
       _ = textureSample(t, s, vec2(0, 0));
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,33 @@
   switch (i32(x)) {
               ^
 
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> t;
+  sampler s;
+};
+struct tint_symbol_inputs {
+  float x [[user(locn0)]];
+};
+
+int tint_f32_to_i32(float value) {
+  return select(2147483647, select((-2147483647 - 1), int(value), (value >= -2147483648.0f)), (value <= 2147483520.0f));
+}
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+  switch(tint_f32_to_i32(x)) {
+    case 0:
+    {
+      tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+      break;
+    }
+    default:
+    {
+      break;
+    }
+  }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+  tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/compound_statement_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/compound_statement_attribute.wgsl.expected.ir.msl
index 36bb093..a83e3b5 100644
--- a/test/tint/diagnostic_filtering/compound_statement_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/compound_statement_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
 <dawn>/test/tint/diagnostic_filtering/compound_statement_attribute.wgsl:8:11 warning: 'textureSample' must only be called from uniform control flow
       _ = textureSample(t, s, vec2(0, 0));
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,22 @@
     if (x > 0) {
         ^
 
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> t;
+  sampler s;
+};
+struct tint_symbol_inputs {
+  float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+  if ((x > 0.0f)) {
+    tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+  }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+  tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/default_case_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/default_case_body_attribute.wgsl.expected.ir.msl
index 3790f50..b9f6609 100644
--- a/test/tint/diagnostic_filtering/default_case_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/default_case_body_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
 <dawn>/test/tint/diagnostic_filtering/default_case_body_attribute.wgsl:8:11 warning: 'textureSample' must only be called from uniform control flow
       _ = textureSample(t, s, vec2(0, 0));
           ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,29 @@
   switch (i32(x)) {
               ^
 
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> t;
+  sampler s;
+};
+struct tint_symbol_inputs {
+  float x [[user(locn0)]];
+};
+
+int tint_f32_to_i32(float value) {
+  return select(2147483647, select((-2147483647 - 1), int(value), (value >= -2147483648.0f)), (value <= 2147483520.0f));
+}
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+  switch(tint_f32_to_i32(x)) {
+    default:
+    {
+      tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+      break;
+    }
+  }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+  tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/directive.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/directive.wgsl.expected.ir.msl
index 5a1ef83..9e92edd 100644
--- a/test/tint/diagnostic_filtering/directive.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/directive.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
 <dawn>/test/tint/diagnostic_filtering/directive.wgsl:9:9 warning: 'textureSample' must only be called from uniform control flow
     _ = textureSample(t, s, vec2(0, 0));
         ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,22 @@
   if (x > 0) {
       ^
 
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> t;
+  sampler s;
+};
+struct tint_symbol_inputs {
+  float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+  if ((x > 0.0f)) {
+    tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+  }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+  tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/else_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/else_body_attribute.wgsl.expected.ir.msl
index a49460b..432c17b 100644
--- a/test/tint/diagnostic_filtering/else_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/else_body_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
 <dawn>/test/tint/diagnostic_filtering/else_body_attribute.wgsl:8:9 warning: 'textureSample' must only be called from uniform control flow
     _ = textureSample(t, s, vec2(0, 0));
         ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,23 @@
   if (x > 0) {
       ^
 
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> t;
+  sampler s;
+};
+struct tint_symbol_inputs {
+  float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+  if ((x > 0.0f)) {
+  } else {
+    tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+  }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+  tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/else_if_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/else_if_body_attribute.wgsl.expected.ir.msl
index 74ecb5a..0890905 100644
--- a/test/tint/diagnostic_filtering/else_if_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/else_if_body_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
 <dawn>/test/tint/diagnostic_filtering/else_if_body_attribute.wgsl:8:9 warning: 'textureSample' must only be called from uniform control flow
     _ = textureSample(t, s, vec2(0, 0));
         ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,25 @@
   if (x > 0) {
       ^
 
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> t;
+  sampler s;
+};
+struct tint_symbol_inputs {
+  float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+  if ((x > 0.0f)) {
+  } else {
+    if ((x < 0.0f)) {
+      tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+    }
+  }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+  tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl
index b0a68a3..d13f01f 100644
--- a/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
 <dawn>/test/tint/diagnostic_filtering/for_loop_body_attribute.wgsl:8:9 warning: 'textureSample' must only be called from uniform control flow
     v = textureSample(t, s, vec2(0, 0));
         ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,32 @@
     v = textureSample(t, s, vec2(0, 0));
         ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
 
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> t;
+  sampler s;
+};
+struct tint_symbol_inputs {
+  float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+  float4 v = float4(0.0f);
+  {
+    while(true) {
+      if ((x > v[0u])) {
+      } else {
+        break;
+      }
+      v = tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+      {
+      }
+      continue;
+    }
+  }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+  tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.ir.msl
index 47e965a..b654cba 100644
--- a/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/function_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
 <dawn>/test/tint/diagnostic_filtering/function_attribute.wgsl:7:9 warning: 'textureSample' must only be called from uniform control flow
     _ = textureSample(t, s, vec2(0, 0));
         ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,22 @@
   if (x > 0) {
       ^
 
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> t;
+  sampler s;
+};
+struct tint_symbol_inputs {
+  float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+  if ((x > 0.0f)) {
+    tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+  }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+  tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/function_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/function_body_attribute.wgsl.expected.ir.msl
index 2c424b1..86b399d 100644
--- a/test/tint/diagnostic_filtering/function_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/function_body_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
 <dawn>/test/tint/diagnostic_filtering/function_body_attribute.wgsl:7:9 warning: 'textureSample' must only be called from uniform control flow
     _ = textureSample(t, s, vec2(0, 0));
         ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,22 @@
   if (x > 0) {
       ^
 
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> t;
+  sampler s;
+};
+struct tint_symbol_inputs {
+  float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+  if ((x > 0.0f)) {
+    tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+  }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+  tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/if_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/if_body_attribute.wgsl.expected.ir.msl
index 4e582b6..dde7eb0 100644
--- a/test/tint/diagnostic_filtering/if_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/if_body_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
 <dawn>/test/tint/diagnostic_filtering/if_body_attribute.wgsl:7:9 warning: 'textureSample' must only be called from uniform control flow
     _ = textureSample(t, s, vec2(0, 0));
         ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,22 @@
   if (x > 0) @diagnostic(warning, derivative_uniformity) {
       ^
 
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> t;
+  sampler s;
+};
+struct tint_symbol_inputs {
+  float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+  if ((x > 0.0f)) {
+    tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+  }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+  tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl b/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl
index 930218e..689a0bf 100644
--- a/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl
+++ b/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
 <dawn>/test/tint/diagnostic_filtering/while_loop_body_attribute.wgsl:8:9 warning: 'textureSample' must only be called from uniform control flow
     v = textureSample(t, s, vec2(0, 0));
         ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
@@ -12,10 +10,32 @@
     v = textureSample(t, s, vec2(0, 0));
         ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
 
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> t;
+  sampler s;
+};
+struct tint_symbol_inputs {
+  float x [[user(locn0)]];
+};
+
+void tint_symbol_inner(float x, tint_module_vars_struct tint_module_vars) {
+  float4 v = float4(0.0f);
+  {
+    while(true) {
+      if ((x > v[0u])) {
+      } else {
+        break;
+      }
+      v = tint_module_vars.t.sample(tint_module_vars.s, float2(0.0f));
+      {
+      }
+      continue;
+    }
+  }
+}
+fragment void tint_symbol(tint_symbol_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s};
+  tint_symbol_inner(inputs.x, tint_module_vars);
+}
diff --git a/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl b/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl
index 4a13f56..935a3a0 100644
--- a/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl
@@ -1,88 +1,52 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> t;
+  sampler s;
+  device atomic_int* a;
+  thread bool* continue_execution;
+};
+struct foo_outputs {
+  int tint_symbol [[color(0)]];
+};
+struct foo_inputs {
+  float in [[user(locn0)]];
+  float2 coord [[user(locn1)]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %t:ptr<handle, texture_2d<f32>, read> = var @binding_point(0, 0)
-  %s:ptr<handle, sampler, read> = var @binding_point(0, 1)
-  %a:ptr<storage, atomic<i32>, read_write> = var @binding_point(0, 2)
-  %continue_execution:ptr<private, bool, read_write> = var, true
+int tint_f32_to_i32(float value) {
+  return select(2147483647, select((-2147483647 - 1), int(value), (value >= -2147483648.0f)), (value <= 2147483520.0f));
 }
-
-%foo = @fragment func(%in:f32 [@location(0)], %coord:vec2<f32> [@location(1)]):i32 [@location(0)] {
-  $B2: {
-    %8:bool = eq %in, 0.0f
-    if %8 [t: $B3] {  # if_1
-      $B3: {  # true
-        store %continue_execution, false
-        exit_if  # if_1
-      }
-    }
-    %9:texture_2d<f32> = load %t
-    %10:sampler = load %s
-    %11:vec4<f32> = textureSample %9, %10, %coord
-    %12:f32 = access %11, 0u
-    %13:i32 = call %tint_f32_to_i32, %12
-    %result:ptr<function, i32, read_write> = var, %13
-    loop [i: $B4, b: $B5, c: $B6] {  # loop_1
-      $B4: {  # initializer
-        %i:ptr<function, i32, read_write> = var, 0i
-        next_iteration  # -> $B5
-      }
-      $B5: {  # body
-        %17:i32 = load %i
-        %18:bool = lt %17, 10i
-        if %18 [t: $B7, f: $B8] {  # if_2
-          $B7: {  # true
-            exit_if  # if_2
-          }
-          $B8: {  # false
-            exit_loop  # loop_1
-          }
-        }
-        %19:i32 = load %i
-        %20:i32 = load %result
-        %21:i32 = add %20, %19
-        store %result, %21
-        continue  # -> $B6
-      }
-      $B6: {  # continuing
-        %22:bool = load %continue_execution
-        %23:i32 = if %22 [t: $B9] {  # if_3
-          $B9: {  # true
-            %24:i32 = atomicAdd %a, 1i
-            exit_if %24  # if_3
-          }
-          # implicit false block: exit_if undef
-        }
-        store %i, %23
-        next_iteration  # -> $B5
-      }
-    }
-    %25:i32 = load %result
-    %26:bool = load %continue_execution
-    %27:bool = eq %26, false
-    if %27 [t: $B10] {  # if_4
-      $B10: {  # true
-        terminate_invocation
-      }
-    }
-    ret %25
+int foo_inner(float in, float2 coord, tint_module_vars_struct tint_module_vars) {
+  if ((in == 0.0f)) {
+    (*tint_module_vars.continue_execution) = false;
   }
-}
-%tint_f32_to_i32 = func(%value:f32):i32 {
-  $B11: {
-    %29:i32 = convert %value
-    %30:bool = gte %value, -2147483648.0f
-    %31:i32 = select -2147483648i, %29, %30
-    %32:bool = lte %value, 2147483520.0f
-    %33:i32 = select 2147483647i, %31, %32
-    ret %33
+  int result = tint_f32_to_i32(tint_module_vars.t.sample(tint_module_vars.s, coord)[0u]);
+  {
+    int i = 0;
+    while(true) {
+      if ((i < 10)) {
+      } else {
+        break;
+      }
+      result = (result + i);
+      {
+        int v = 0;
+        if ((*tint_module_vars.continue_execution)) {
+          v = atomic_fetch_add_explicit(tint_module_vars.a, 1, memory_order_relaxed);
+        }
+        i = v;
+      }
+      continue;
+    }
   }
+  if (!((*tint_module_vars.continue_execution))) {
+    discard_fragment();
+  }
+  return result;
 }
-
-unhandled variable address space
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
+fragment foo_outputs foo(foo_inputs inputs [[stage_in]], texture2d<float, access::sample> t [[texture(0)]], sampler s [[sampler(0)]], device atomic_int* a [[buffer(0)]]) {
+  thread bool continue_execution = true;
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.t=t, .s=s, .a=a, .continue_execution=(&continue_execution)};
+  return foo_outputs{.tint_symbol=foo_inner(inputs.in, inputs.coord, tint_module_vars)};
+}
diff --git a/test/tint/unittest/reader/spirv/Samples_SpvParserHandleTest_RegisterHandleUsage_SampledImage_Variable_2.spvasm.expected.ir.msl b/test/tint/unittest/reader/spirv/Samples_SpvParserHandleTest_RegisterHandleUsage_SampledImage_Variable_2.spvasm.expected.ir.msl
deleted file mode 100644
index 3543919..0000000
--- a/test/tint/unittest/reader/spirv/Samples_SpvParserHandleTest_RegisterHandleUsage_SampledImage_Variable_2.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,9 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/unittest/reader/spirv/Samples_SpvParserHandleTest_RegisterHandleUsage_SampledImage_Variable_6.spvasm.expected.ir.msl b/test/tint/unittest/reader/spirv/Samples_SpvParserHandleTest_RegisterHandleUsage_SampledImage_Variable_6.spvasm.expected.ir.msl
deleted file mode 100644
index 3543919..0000000
--- a/test/tint/unittest/reader/spirv/Samples_SpvParserHandleTest_RegisterHandleUsage_SampledImage_Variable_6.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,9 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************
diff --git a/test/tint/unittest/reader/spirv/SpvParserHandleTest_NeverGenerateConstDeclForHandle_UseVariableDirectly.spvasm.expected.ir.msl b/test/tint/unittest/reader/spirv/SpvParserHandleTest_NeverGenerateConstDeclForHandle_UseVariableDirectly.spvasm.expected.ir.msl
deleted file mode 100644
index 3543919..0000000
--- a/test/tint/unittest/reader/spirv/SpvParserHandleTest_NeverGenerateConstDeclForHandle_UseVariableDirectly.spvasm.expected.ir.msl
+++ /dev/null
@@ -1,9 +0,0 @@
-SKIP: FAILED
-
-../../src/tint/lang/msl/writer/printer/printer.cc:924 internal compiler error: TINT_UNREACHABLE unhandled: textureSample
-********************************************************************
-*  The tint shader compiler has encountered an unexpected error.   *
-*                                                                  *
-*  Please help us fix this issue by submitting a bug report at     *
-*  crbug.com/tint with the source program that triggered the bug.  *
-********************************************************************