diff --git a/src/tint/lang/msl/builtin_fn.cc b/src/tint/lang/msl/builtin_fn.cc
index 1eedd33..b7547d4 100644
--- a/src/tint/lang/msl/builtin_fn.cc
+++ b/src/tint/lang/msl/builtin_fn.cc
@@ -72,6 +72,8 @@
             return "get_height";
         case BuiltinFn::kGetDepth:
             return "get_depth";
+        case BuiltinFn::kGetNumMipLevels:
+            return "get_num_mip_levels";
         case BuiltinFn::kRead:
             return "read";
         case BuiltinFn::kSample:
diff --git a/src/tint/lang/msl/builtin_fn.h b/src/tint/lang/msl/builtin_fn.h
index d1665b4..a92c166 100644
--- a/src/tint/lang/msl/builtin_fn.h
+++ b/src/tint/lang/msl/builtin_fn.h
@@ -62,6 +62,7 @@
     kGetWidth,
     kGetHeight,
     kGetDepth,
+    kGetNumMipLevels,
     kRead,
     kSample,
     kWrite,
diff --git a/src/tint/lang/msl/intrinsic/data.cc b/src/tint/lang/msl/intrinsic/data.cc
index 710f451..3ea527b 100644
--- a/src/tint/lang/msl/intrinsic/data.cc
+++ b/src/tint/lang/msl/intrinsic/data.cc
@@ -3250,6 +3250,116 @@
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
     /* num_parameters */ 1,
     /* num_explicit_templates */ 0,
+    /* num_templates   */ 1,
+    /* templates */ TemplateIndex(0),
+    /* parameters */ ParameterIndex(200),
+    /* return_matcher_indices */ MatcherIndicesIndex(59),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [100] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 1,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 1,
+    /* templates */ TemplateIndex(0),
+    /* parameters */ ParameterIndex(120),
+    /* return_matcher_indices */ MatcherIndicesIndex(59),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [101] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 1,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 1,
+    /* templates */ TemplateIndex(0),
+    /* parameters */ ParameterIndex(73),
+    /* return_matcher_indices */ MatcherIndicesIndex(59),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [102] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 1,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 1,
+    /* templates */ TemplateIndex(0),
+    /* parameters */ ParameterIndex(123),
+    /* return_matcher_indices */ MatcherIndicesIndex(59),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [103] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 1,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 1,
+    /* templates */ TemplateIndex(0),
+    /* parameters */ ParameterIndex(182),
+    /* return_matcher_indices */ MatcherIndicesIndex(59),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [104] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 1,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 1,
+    /* templates */ TemplateIndex(0),
+    /* parameters */ ParameterIndex(184),
+    /* return_matcher_indices */ MatcherIndicesIndex(59),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [105] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 1,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 0,
+    /* templates */ TemplateIndex(/* invalid */),
+    /* parameters */ ParameterIndex(48),
+    /* return_matcher_indices */ MatcherIndicesIndex(59),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [106] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 1,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 0,
+    /* templates */ TemplateIndex(/* invalid */),
+    /* parameters */ ParameterIndex(6),
+    /* return_matcher_indices */ MatcherIndicesIndex(59),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [107] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 1,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 0,
+    /* templates */ TemplateIndex(/* invalid */),
+    /* parameters */ ParameterIndex(97),
+    /* return_matcher_indices */ MatcherIndicesIndex(59),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [108] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 1,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 0,
+    /* templates */ TemplateIndex(/* invalid */),
+    /* parameters */ ParameterIndex(53),
+    /* return_matcher_indices */ MatcherIndicesIndex(59),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [109] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
+    /* num_parameters */ 1,
+    /* num_explicit_templates */ 0,
     /* num_templates   */ 2,
     /* templates */ TemplateIndex(5),
     /* parameters */ ParameterIndex(220),
@@ -3257,7 +3367,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [100] */
+    /* [110] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
     /* num_parameters */ 1,
     /* num_explicit_templates */ 0,
@@ -3268,7 +3378,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [101] */
+    /* [111] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
     /* num_parameters */ 1,
     /* num_explicit_templates */ 0,
@@ -3279,7 +3389,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [102] */
+    /* [112] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
     /* num_parameters */ 1,
     /* num_explicit_templates */ 0,
@@ -3290,7 +3400,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [103] */
+    /* [113] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
     /* num_parameters */ 2,
     /* num_explicit_templates */ 0,
@@ -3301,7 +3411,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [104] */
+    /* [114] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline, OverloadFlag::kMemberFunction),
     /* num_parameters */ 2,
     /* num_explicit_templates */ 0,
@@ -3312,7 +3422,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [105] */
+    /* [115] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 5,
     /* num_explicit_templates */ 0,
@@ -3323,7 +3433,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [106] */
+    /* [116] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 3,
     /* num_explicit_templates */ 0,
@@ -3334,7 +3444,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [107] */
+    /* [117] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 2,
     /* num_explicit_templates */ 0,
@@ -3345,7 +3455,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [108] */
+    /* [118] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 3,
     /* num_explicit_templates */ 0,
@@ -3356,7 +3466,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [109] */
+    /* [119] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 1,
     /* num_explicit_templates */ 0,
@@ -3376,67 +3486,67 @@
     /* [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(105),
+    /* overloads */ OverloadIndex(115),
   },
   {
     /* [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(106),
+    /* overloads */ OverloadIndex(116),
   },
   {
     /* [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(106),
+    /* overloads */ OverloadIndex(116),
   },
   {
     /* [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(106),
+    /* overloads */ OverloadIndex(116),
   },
   {
     /* [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(106),
+    /* overloads */ OverloadIndex(116),
   },
   {
     /* [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(106),
+    /* overloads */ OverloadIndex(116),
   },
   {
     /* [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(106),
+    /* overloads */ OverloadIndex(116),
   },
   {
     /* [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(106),
+    /* overloads */ OverloadIndex(116),
   },
   {
     /* [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(106),
+    /* overloads */ OverloadIndex(116),
   },
   {
     /* [9] */
     /* fn atomic_load_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, u32) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(107),
+    /* overloads */ OverloadIndex(117),
   },
   {
     /* [10] */
     /* fn atomic_store_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(108),
+    /* overloads */ OverloadIndex(118),
   },
   {
     /* [11] */
@@ -3445,7 +3555,7 @@
     /* fn fence[F : texel_format, A : access](texture: texture_storage_2d_array<F, A>) */
     /* fn fence[F : texel_format, A : access](texture: texture_storage_3d<F, A>) */
     /* num overloads */ 4,
-    /* overloads */ OverloadIndex(99),
+    /* overloads */ OverloadIndex(109),
   },
   {
     /* [12] */
@@ -3492,10 +3602,25 @@
     /* fn get_depth[T : fiu32](texture: texture_3d<T>, u32) -> u32 */
     /* fn get_depth[F : texel_format, A : access](texture: texture_storage_3d<F, A>, u32) -> u32 */
     /* num overloads */ 2,
-    /* overloads */ OverloadIndex(103),
+    /* overloads */ OverloadIndex(113),
   },
   {
     /* [15] */
+    /* fn get_num_mip_levels[T : fiu32](texture: texture_1d<T>) -> u32 */
+    /* fn get_num_mip_levels[T : fiu32](texture: texture_2d<T>) -> u32 */
+    /* fn get_num_mip_levels[T : fiu32](texture: texture_2d_array<T>) -> u32 */
+    /* fn get_num_mip_levels[T : fiu32](texture: texture_3d<T>) -> u32 */
+    /* fn get_num_mip_levels[T : fiu32](texture: texture_cube<T>) -> u32 */
+    /* fn get_num_mip_levels[T : fiu32](texture: texture_cube_array<T>) -> u32 */
+    /* fn get_num_mip_levels(texture: texture_depth_2d) -> u32 */
+    /* fn get_num_mip_levels(texture: texture_depth_2d_array) -> u32 */
+    /* fn get_num_mip_levels(texture: texture_depth_cube) -> u32 */
+    /* fn get_num_mip_levels(texture: texture_depth_cube_array) -> u32 */
+    /* num overloads */ 10,
+    /* overloads */ OverloadIndex(99),
+  },
+  {
+    /* [16] */
     /* fn read[T : fiu32](texture: texture_1d<T>, coords: u32) -> vec4<T> */
     /* fn read[T : fiu32, L : iu32](texture: texture_2d<T>, coords: vec2<u32>, level: L) -> vec4<T> */
     /* fn read[T : fiu32, A : iu32, L : iu32](texture: texture_2d_array<T>, coords: vec2<u32>, array_index: A, level: L) -> vec4<T> */
@@ -3520,7 +3645,7 @@
     /* overloads */ OverloadIndex(37),
   },
   {
-    /* [16] */
+    /* [17] */
     /* 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> */
@@ -3562,7 +3687,7 @@
     /* overloads */ OverloadIndex(0),
   },
   {
-    /* [17] */
+    /* [18] */
     /* fn write(texture: texture_storage_1d<f32_texel_format, writable>, value: vec4<f32>, coords: u32) */
     /* fn write(texture: texture_storage_2d<f32_texel_format, writable>, value: vec4<f32>, coords: vec2<u32>) */
     /* fn write[A : iu32](texture: texture_storage_2d_array<f32_texel_format, writable>, value: vec4<f32>, coords: vec2<u32>, array_index: A) */
@@ -3579,10 +3704,10 @@
     /* overloads */ OverloadIndex(87),
   },
   {
-    /* [18] */
+    /* [19] */
     /* fn threadgroup_barrier(u32) */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(109),
+    /* overloads */ OverloadIndex(119),
   },
 };
 
diff --git a/src/tint/lang/msl/msl.def b/src/tint/lang/msl/msl.def
index eb111b7..d660f01 100644
--- a/src/tint/lang/msl/msl.def
+++ b/src/tint/lang/msl/msl.def
@@ -175,6 +175,17 @@
 @member_function fn get_depth[T: fiu32](texture: texture_3d<T>, u32) -> u32
 @member_function fn get_depth[F: texel_format, A: access](texture: texture_storage_3d<F, A>, u32) -> u32
 
+@member_function fn get_num_mip_levels[T: fiu32](texture: texture_1d<T>) -> u32
+@member_function fn get_num_mip_levels[T: fiu32](texture: texture_2d<T>) -> u32
+@member_function fn get_num_mip_levels[T: fiu32](texture: texture_2d_array<T>) -> u32
+@member_function fn get_num_mip_levels[T: fiu32](texture: texture_3d<T>) -> u32
+@member_function fn get_num_mip_levels[T: fiu32](texture: texture_cube<T>) -> u32
+@member_function fn get_num_mip_levels[T: fiu32](texture: texture_cube_array<T>) -> u32
+@member_function fn get_num_mip_levels(texture: texture_depth_2d) -> u32
+@member_function fn get_num_mip_levels(texture: texture_depth_2d_array) -> u32
+@member_function fn get_num_mip_levels(texture: texture_depth_cube) -> u32
+@member_function fn get_num_mip_levels(texture: texture_depth_cube_array) -> u32
+
 @member_function fn read[T: fiu32](texture: texture_1d<T>, coords: u32) -> vec4<T>
 @member_function fn read[T: fiu32, L: iu32](texture: texture_2d<T>, coords: vec2<u32>, level: L) -> vec4<T>
 @member_function fn read[T: fiu32, A: iu32, L: iu32](texture: texture_2d_array<T>, coords: vec2<u32>, array_index: A, level: L) -> vec4<T>
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
index be8c955..b594663 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
@@ -91,6 +91,7 @@
                     case core::BuiltinFn::kAtomicXor:
                     case core::BuiltinFn::kTextureDimensions:
                     case core::BuiltinFn::kTextureLoad:
+                    case core::BuiltinFn::kTextureNumLevels:
                     case core::BuiltinFn::kTextureSample:
                     case core::BuiltinFn::kTextureSampleBias:
                     case core::BuiltinFn::kTextureSampleLevel:
@@ -152,6 +153,9 @@
                 case core::BuiltinFn::kTextureLoad:
                     TextureLoad(builtin);
                     break;
+                case core::BuiltinFn::kTextureNumLevels:
+                    TextureNumLevels(builtin);
+                    break;
                 case core::BuiltinFn::kTextureSample:
                     TextureSample(builtin);
                     break;
@@ -329,6 +333,17 @@
         builtin->Destroy();
     }
 
+    /// Replace a textureNumLevels call with the equivalent MSL intrinsic.
+    /// @param builtin the builtin call instruction
+    void TextureNumLevels(core::ir::CoreBuiltinCall* builtin) {
+        // The MSL intrinsic is a member function, so we split the first argument off as the object.
+        auto* tex = builtin->Args()[0];
+        auto* call = b.MemberCallWithResult<msl::ir::MemberBuiltinCall>(
+            builtin->DetachResult(), msl::BuiltinFn::kGetNumMipLevels, tex);
+        call->InsertBefore(builtin);
+        builtin->Destroy();
+    }
+
     /// Replace a textureSample call with the equivalent MSL intrinsic.
     /// @param builtin the builtin call instruction
     void TextureSample(core::ir::CoreBuiltinCall* builtin) {
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 5207738..f712723 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
@@ -1251,6 +1251,40 @@
     EXPECT_EQ(expect, str());
 }
 
+TEST_F(MslWriter_BuiltinPolyfillTest, TextureNumLevels) {
+    auto* t = b.FunctionParam(
+        "t", ty.Get<core::type::SampledTexture>(core::type::TextureDimension::k2d, ty.f32()));
+    auto* func = b.Function("foo", ty.u32());
+    func->SetParams({t});
+    b.Append(func->Block(), [&] {
+        auto* result = b.Call<u32>(core::BuiltinFn::kTextureNumLevels, t);
+        b.Return(func, result);
+    });
+
+    auto* src = R"(
+%foo = func(%t:texture_2d<f32>):u32 {
+  $B1: {
+    %3:u32 = textureNumLevels %t
+    ret %3
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%foo = func(%t:texture_2d<f32>):u32 {
+  $B1: {
+    %3:u32 = %t.get_num_mip_levels
+    ret %3
+  }
+}
+)";
+
+    Run(BuiltinPolyfill);
+
+    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()));
diff --git a/test/tint/builtins/gen/literal/textureNumLevels/181090.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/181090.wgsl.expected.ir.msl
index 1d6870d..146e4f7 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/181090.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/181090.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d_array<uint, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d_array<u32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_181090(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_181090 = func():void {
-  $B2: {
-    %4:texture_2d_array<u32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_181090(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_181090
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture2d_array<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_181090(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_181090
-    ret
-  }
+kernel void compute_main(texture2d_array<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_181090(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_181090
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture2d_array<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/1a3fa9.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/1a3fa9.wgsl.expected.ir.msl
index add0e08..cbec735 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/1a3fa9.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/1a3fa9.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<int, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d<i32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_1a3fa9(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_1a3fa9 = func():void {
-  $B2: {
-    %4:texture_2d<i32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_1a3fa9(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_1a3fa9
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture2d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_1a3fa9(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_1a3fa9
-    ret
-  }
+kernel void compute_main(texture2d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_1a3fa9(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_1a3fa9
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture2d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/1a7fc3.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/1a7fc3.wgsl.expected.ir.msl
index 0e2d138..6f3ceac 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/1a7fc3.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/1a7fc3.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture1d<uint, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_1d<u32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_1a7fc3(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_1a7fc3 = func():void {
-  $B2: {
-    %4:texture_1d<u32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_1a7fc3(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_1a7fc3
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture1d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_1a7fc3(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_1a7fc3
-    ret
-  }
+kernel void compute_main(texture1d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_1a7fc3(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_1a7fc3
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture1d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/2267d8.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/2267d8.wgsl.expected.ir.msl
index b5543c7..fff04f4 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/2267d8.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/2267d8.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube<int, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_cube<i32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_2267d8(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_2267d8 = func():void {
-  $B2: {
-    %4:texture_cube<i32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_2267d8(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_2267d8
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texturecube<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_2267d8(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_2267d8
-    ret
-  }
+kernel void compute_main(texturecube<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_2267d8(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_2267d8
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texturecube<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/24b2c6.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/24b2c6.wgsl.expected.ir.msl
index df3f5f9..3298605 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/24b2c6.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/24b2c6.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_24b2c6(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_24b2c6 = func():void {
-  $B2: {
-    %4:texture_2d<f32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_24b2c6(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_24b2c6
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture2d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_24b2c6(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_24b2c6
-    ret
-  }
+kernel void compute_main(texture2d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_24b2c6(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_24b2c6
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture2d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/2bea6c.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/2bea6c.wgsl.expected.ir.msl
index c2a98e4..658d0b5 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/2bea6c.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/2bea6c.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depthcube_array<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_2bea6c(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_2bea6c = func():void {
-  $B2: {
-    %4:texture_depth_cube_array = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_2bea6c(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_2bea6c
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(depthcube_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_2bea6c(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_2bea6c
-    ret
-  }
+kernel void compute_main(depthcube_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_2bea6c(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_2bea6c
-    ret
-  }
+vertex vertex_main_outputs vertex_main(depthcube_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/2df1ab.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/2df1ab.wgsl.expected.ir.msl
index 9562080..ffef839 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/2df1ab.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/2df1ab.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_2df1ab(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_2df1ab = func():void {
-  $B2: {
-    %4:texture_cube<f32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_2df1ab(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_2df1ab
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texturecube<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_2df1ab(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_2df1ab
-    ret
-  }
+kernel void compute_main(texturecube<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_2df1ab(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_2df1ab
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texturecube<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/46dbd8.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/46dbd8.wgsl.expected.ir.msl
index 7cf81ff..b688a4d 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/46dbd8.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/46dbd8.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d_array<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_46dbd8(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_46dbd8 = func():void {
-  $B2: {
-    %4:texture_2d_array<f32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_46dbd8(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_46dbd8
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_46dbd8(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_46dbd8
-    ret
-  }
+kernel void compute_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_46dbd8(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_46dbd8
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/60d9b8.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/60d9b8.wgsl.expected.ir.msl
index 51be4e4..78fd8bd 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/60d9b8.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/60d9b8.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture3d<uint, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_3d<u32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_60d9b8(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_60d9b8 = func():void {
-  $B2: {
-    %4:texture_3d<u32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_60d9b8(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_60d9b8
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture3d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_60d9b8(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_60d9b8
-    ret
-  }
+kernel void compute_main(texture3d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_60d9b8(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_60d9b8
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture3d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/903920.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/903920.wgsl.expected.ir.msl
index cdb68aa..f8901ee 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/903920.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/903920.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube_array<int, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_cube_array<i32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_903920(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_903920 = func():void {
-  $B2: {
-    %4:texture_cube_array<i32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_903920(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_903920
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texturecube_array<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_903920(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_903920
-    ret
-  }
+kernel void compute_main(texturecube_array<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_903920(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_903920
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texturecube_array<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/9a1a65.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/9a1a65.wgsl.expected.ir.msl
index 56868d8..4896390 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/9a1a65.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/9a1a65.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture3d<int, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_3d<i32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_9a1a65(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_9a1a65 = func():void {
-  $B2: {
-    %4:texture_3d<i32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_9a1a65(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_9a1a65
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture3d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_9a1a65(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_9a1a65
-    ret
-  }
+kernel void compute_main(texture3d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_9a1a65(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_9a1a65
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture3d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/adc783.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/adc783.wgsl.expected.ir.msl
index 358b38a..fdc6276 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/adc783.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/adc783.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d_array<int, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d_array<i32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_adc783(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_adc783 = func():void {
-  $B2: {
-    %4:texture_2d_array<i32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_adc783(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_adc783
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture2d_array<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_adc783(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_adc783
-    ret
-  }
+kernel void compute_main(texture2d_array<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_adc783(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_adc783
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture2d_array<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/ae911c.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/ae911c.wgsl.expected.ir.msl
index 352fc7f..f766da3 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/ae911c.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/ae911c.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depth2d_array<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_ae911c(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_ae911c = func():void {
-  $B2: {
-    %4:texture_depth_2d_array = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_ae911c(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_ae911c
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_ae911c(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_ae911c
-    ret
-  }
+kernel void compute_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_ae911c(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_ae911c
-    ret
-  }
+vertex vertex_main_outputs vertex_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/c386c8.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/c386c8.wgsl.expected.ir.msl
index a381977..add247c 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/c386c8.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/c386c8.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube<uint, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_cube<u32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_c386c8(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_c386c8 = func():void {
-  $B2: {
-    %4:texture_cube<u32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_c386c8(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_c386c8
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texturecube<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_c386c8(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_c386c8
-    ret
-  }
+kernel void compute_main(texturecube<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_c386c8(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_c386c8
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texturecube<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/c399f9.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/c399f9.wgsl.expected.ir.msl
index ab8b06c..1292577 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/c399f9.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/c399f9.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture1d<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_c399f9(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_c399f9 = func():void {
-  $B2: {
-    %4:texture_1d<f32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_c399f9(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_c399f9
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture1d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_c399f9(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_c399f9
-    ret
-  }
+kernel void compute_main(texture1d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_c399f9(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_c399f9
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture1d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/c8c25c.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/c8c25c.wgsl.expected.ir.msl
index ac6bf19..bdc2dc2 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/c8c25c.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/c8c25c.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depthcube<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_c8c25c(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_c8c25c = func():void {
-  $B2: {
-    %4:texture_depth_cube = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_c8c25c(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_c8c25c
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(depthcube<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_c8c25c(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_c8c25c
-    ret
-  }
+kernel void compute_main(depthcube<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_c8c25c(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_c8c25c
-    ret
-  }
+vertex vertex_main_outputs vertex_main(depthcube<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/d63126.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/d63126.wgsl.expected.ir.msl
index a3f9b35..60c08d1 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/d63126.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/d63126.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depth2d<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_d63126(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_d63126 = func():void {
-  $B2: {
-    %4:texture_depth_2d = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_d63126(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_d63126
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(depth2d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_d63126(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_d63126
-    ret
-  }
+kernel void compute_main(depth2d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_d63126(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_d63126
-    ret
-  }
+vertex vertex_main_outputs vertex_main(depth2d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/d8f73b.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/d8f73b.wgsl.expected.ir.msl
index 73607e8..8cad080 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/d8f73b.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/d8f73b.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube_array<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_d8f73b(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_d8f73b = func():void {
-  $B2: {
-    %4:texture_cube_array<f32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_d8f73b(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_d8f73b
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texturecube_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_d8f73b(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_d8f73b
-    ret
-  }
+kernel void compute_main(texturecube_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_d8f73b(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_d8f73b
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texturecube_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/ef7944.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/ef7944.wgsl.expected.ir.msl
index 52292ff..57138ba 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/ef7944.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/ef7944.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture3d<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_ef7944(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_ef7944 = func():void {
-  $B2: {
-    %4:texture_3d<f32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_ef7944(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_ef7944
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture3d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_ef7944(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_ef7944
-    ret
-  }
+kernel void compute_main(texture3d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_ef7944(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_ef7944
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture3d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/efd6df.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/efd6df.wgsl.expected.ir.msl
index 0d6eac8..575ad08 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/efd6df.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/efd6df.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<uint, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d<u32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_efd6df(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_efd6df = func():void {
-  $B2: {
-    %4:texture_2d<u32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_efd6df(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_efd6df
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture2d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_efd6df(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_efd6df
-    ret
-  }
+kernel void compute_main(texture2d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_efd6df(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_efd6df
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture2d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/f742c0.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/f742c0.wgsl.expected.ir.msl
index d26c684..96dc82a 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/f742c0.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/f742c0.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture1d<int, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_1d<i32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_f742c0(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_f742c0 = func():void {
-  $B2: {
-    %4:texture_1d<i32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_f742c0(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_f742c0
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture1d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_f742c0(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_f742c0
-    ret
-  }
+kernel void compute_main(texture1d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_f742c0(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_f742c0
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture1d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/fe2171.wgsl.expected.ir.msl b/test/tint/builtins/gen/literal/textureNumLevels/fe2171.wgsl.expected.ir.msl
index 7ee8eef..e8ff625 100644
--- a/test/tint/builtins/gen/literal/textureNumLevels/fe2171.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/literal/textureNumLevels/fe2171.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube_array<uint, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_cube_array<u32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_fe2171(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_fe2171 = func():void {
-  $B2: {
-    %4:texture_cube_array<u32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_fe2171(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_fe2171
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texturecube_array<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_fe2171(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_fe2171
-    ret
-  }
+kernel void compute_main(texturecube_array<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_fe2171(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_fe2171
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texturecube_array<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/181090.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/181090.wgsl.expected.ir.msl
index 1d6870d..146e4f7 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/181090.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/181090.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d_array<uint, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d_array<u32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_181090(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_181090 = func():void {
-  $B2: {
-    %4:texture_2d_array<u32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_181090(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_181090
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture2d_array<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_181090(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_181090
-    ret
-  }
+kernel void compute_main(texture2d_array<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_181090(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_181090
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture2d_array<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/1a3fa9.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/1a3fa9.wgsl.expected.ir.msl
index add0e08..cbec735 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/1a3fa9.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/1a3fa9.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<int, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d<i32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_1a3fa9(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_1a3fa9 = func():void {
-  $B2: {
-    %4:texture_2d<i32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_1a3fa9(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_1a3fa9
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture2d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_1a3fa9(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_1a3fa9
-    ret
-  }
+kernel void compute_main(texture2d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_1a3fa9(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_1a3fa9
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture2d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/1a7fc3.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/1a7fc3.wgsl.expected.ir.msl
index 0e2d138..6f3ceac 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/1a7fc3.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/1a7fc3.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture1d<uint, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_1d<u32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_1a7fc3(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_1a7fc3 = func():void {
-  $B2: {
-    %4:texture_1d<u32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_1a7fc3(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_1a7fc3
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture1d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_1a7fc3(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_1a7fc3
-    ret
-  }
+kernel void compute_main(texture1d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_1a7fc3(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_1a7fc3
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture1d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/2267d8.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/2267d8.wgsl.expected.ir.msl
index b5543c7..fff04f4 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/2267d8.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/2267d8.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube<int, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_cube<i32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_2267d8(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_2267d8 = func():void {
-  $B2: {
-    %4:texture_cube<i32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_2267d8(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_2267d8
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texturecube<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_2267d8(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_2267d8
-    ret
-  }
+kernel void compute_main(texturecube<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_2267d8(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_2267d8
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texturecube<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/24b2c6.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/24b2c6.wgsl.expected.ir.msl
index df3f5f9..3298605 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/24b2c6.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/24b2c6.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_24b2c6(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_24b2c6 = func():void {
-  $B2: {
-    %4:texture_2d<f32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_24b2c6(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_24b2c6
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture2d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_24b2c6(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_24b2c6
-    ret
-  }
+kernel void compute_main(texture2d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_24b2c6(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_24b2c6
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture2d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/2bea6c.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/2bea6c.wgsl.expected.ir.msl
index c2a98e4..658d0b5 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/2bea6c.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/2bea6c.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depthcube_array<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_2bea6c(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_2bea6c = func():void {
-  $B2: {
-    %4:texture_depth_cube_array = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_2bea6c(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_2bea6c
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(depthcube_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_2bea6c(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_2bea6c
-    ret
-  }
+kernel void compute_main(depthcube_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_2bea6c(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_2bea6c
-    ret
-  }
+vertex vertex_main_outputs vertex_main(depthcube_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/2df1ab.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/2df1ab.wgsl.expected.ir.msl
index 9562080..ffef839 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/2df1ab.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/2df1ab.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_2df1ab(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_2df1ab = func():void {
-  $B2: {
-    %4:texture_cube<f32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_2df1ab(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_2df1ab
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texturecube<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_2df1ab(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_2df1ab
-    ret
-  }
+kernel void compute_main(texturecube<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_2df1ab(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_2df1ab
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texturecube<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/46dbd8.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/46dbd8.wgsl.expected.ir.msl
index 7cf81ff..b688a4d 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/46dbd8.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/46dbd8.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d_array<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_46dbd8(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_46dbd8 = func():void {
-  $B2: {
-    %4:texture_2d_array<f32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_46dbd8(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_46dbd8
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_46dbd8(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_46dbd8
-    ret
-  }
+kernel void compute_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_46dbd8(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_46dbd8
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture2d_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/60d9b8.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/60d9b8.wgsl.expected.ir.msl
index 51be4e4..78fd8bd 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/60d9b8.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/60d9b8.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture3d<uint, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_3d<u32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_60d9b8(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_60d9b8 = func():void {
-  $B2: {
-    %4:texture_3d<u32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_60d9b8(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_60d9b8
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture3d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_60d9b8(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_60d9b8
-    ret
-  }
+kernel void compute_main(texture3d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_60d9b8(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_60d9b8
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture3d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/903920.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/903920.wgsl.expected.ir.msl
index cdb68aa..f8901ee 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/903920.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/903920.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube_array<int, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_cube_array<i32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_903920(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_903920 = func():void {
-  $B2: {
-    %4:texture_cube_array<i32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_903920(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_903920
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texturecube_array<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_903920(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_903920
-    ret
-  }
+kernel void compute_main(texturecube_array<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_903920(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_903920
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texturecube_array<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/9a1a65.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/9a1a65.wgsl.expected.ir.msl
index 56868d8..4896390 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/9a1a65.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/9a1a65.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture3d<int, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_3d<i32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_9a1a65(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_9a1a65 = func():void {
-  $B2: {
-    %4:texture_3d<i32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_9a1a65(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_9a1a65
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture3d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_9a1a65(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_9a1a65
-    ret
-  }
+kernel void compute_main(texture3d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_9a1a65(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_9a1a65
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture3d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/adc783.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/adc783.wgsl.expected.ir.msl
index 358b38a..fdc6276 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/adc783.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/adc783.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d_array<int, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d_array<i32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_adc783(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_adc783 = func():void {
-  $B2: {
-    %4:texture_2d_array<i32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_adc783(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_adc783
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture2d_array<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_adc783(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_adc783
-    ret
-  }
+kernel void compute_main(texture2d_array<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_adc783(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_adc783
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture2d_array<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/ae911c.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/ae911c.wgsl.expected.ir.msl
index 352fc7f..f766da3 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/ae911c.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/ae911c.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depth2d_array<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_ae911c(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_ae911c = func():void {
-  $B2: {
-    %4:texture_depth_2d_array = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_ae911c(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_ae911c
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_ae911c(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_ae911c
-    ret
-  }
+kernel void compute_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_ae911c(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_ae911c
-    ret
-  }
+vertex vertex_main_outputs vertex_main(depth2d_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/c386c8.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/c386c8.wgsl.expected.ir.msl
index a381977..add247c 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/c386c8.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/c386c8.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube<uint, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_cube<u32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_c386c8(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_c386c8 = func():void {
-  $B2: {
-    %4:texture_cube<u32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_c386c8(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_c386c8
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texturecube<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_c386c8(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_c386c8
-    ret
-  }
+kernel void compute_main(texturecube<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_c386c8(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_c386c8
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texturecube<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/c399f9.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/c399f9.wgsl.expected.ir.msl
index ab8b06c..1292577 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/c399f9.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/c399f9.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture1d<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_c399f9(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_c399f9 = func():void {
-  $B2: {
-    %4:texture_1d<f32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_c399f9(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_c399f9
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture1d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_c399f9(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_c399f9
-    ret
-  }
+kernel void compute_main(texture1d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_c399f9(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_c399f9
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture1d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/c8c25c.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/c8c25c.wgsl.expected.ir.msl
index ac6bf19..bdc2dc2 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/c8c25c.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/c8c25c.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depthcube<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_c8c25c(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_c8c25c = func():void {
-  $B2: {
-    %4:texture_depth_cube = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_c8c25c(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_c8c25c
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(depthcube<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_c8c25c(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_c8c25c
-    ret
-  }
+kernel void compute_main(depthcube<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_c8c25c(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_c8c25c
-    ret
-  }
+vertex vertex_main_outputs vertex_main(depthcube<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/d63126.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/d63126.wgsl.expected.ir.msl
index a3f9b35..60c08d1 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/d63126.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/d63126.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  depth2d<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_d63126(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_d63126 = func():void {
-  $B2: {
-    %4:texture_depth_2d = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_d63126(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_d63126
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(depth2d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_d63126(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_d63126
-    ret
-  }
+kernel void compute_main(depth2d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_d63126(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_d63126
-    ret
-  }
+vertex vertex_main_outputs vertex_main(depth2d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/d8f73b.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/d8f73b.wgsl.expected.ir.msl
index 73607e8..8cad080 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/d8f73b.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/d8f73b.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube_array<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_d8f73b(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_d8f73b = func():void {
-  $B2: {
-    %4:texture_cube_array<f32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_d8f73b(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_d8f73b
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texturecube_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_d8f73b(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_d8f73b
-    ret
-  }
+kernel void compute_main(texturecube_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_d8f73b(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_d8f73b
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texturecube_array<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/ef7944.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/ef7944.wgsl.expected.ir.msl
index 52292ff..57138ba 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/ef7944.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/ef7944.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture3d<float, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../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)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_ef7944(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_ef7944 = func():void {
-  $B2: {
-    %4:texture_3d<f32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_ef7944(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_ef7944
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture3d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_ef7944(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_ef7944
-    ret
-  }
+kernel void compute_main(texture3d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_ef7944(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_ef7944
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture3d<float, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/efd6df.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/efd6df.wgsl.expected.ir.msl
index 0d6eac8..575ad08 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/efd6df.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/efd6df.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture2d<uint, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_2d<u32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_efd6df(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_efd6df = func():void {
-  $B2: {
-    %4:texture_2d<u32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_efd6df(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_efd6df
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture2d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_efd6df(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_efd6df
-    ret
-  }
+kernel void compute_main(texture2d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_efd6df(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_efd6df
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture2d<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/f742c0.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/f742c0.wgsl.expected.ir.msl
index d26c684..96dc82a 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/f742c0.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/f742c0.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texture1d<int, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_1d<i32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_f742c0(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_f742c0 = func():void {
-  $B2: {
-    %4:texture_1d<i32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_f742c0(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_f742c0
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texture1d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_f742c0(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_f742c0
-    ret
-  }
+kernel void compute_main(texture1d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_f742c0(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_f742c0
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texture1d<int, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/textureNumLevels/fe2171.wgsl.expected.ir.msl b/test/tint/builtins/gen/var/textureNumLevels/fe2171.wgsl.expected.ir.msl
index 7ee8eef..e8ff625 100644
--- a/test/tint/builtins/gen/var/textureNumLevels/fe2171.wgsl.expected.ir.msl
+++ b/test/tint/builtins/gen/var/textureNumLevels/fe2171.wgsl.expected.ir.msl
@@ -1,43 +1,30 @@
-SKIP: FAILED
+#include <metal_stdlib>
+using namespace metal;
+struct tint_module_vars_struct {
+  texturecube_array<uint, access::sample> arg_0;
+  device uint* prevent_dce;
+};
+struct vertex_main_outputs {
+  float4 tint_symbol [[position]];
+};
 
-../../src/tint/lang/msl/writer/printer/printer.cc:500 internal compiler error: $B1: {  # root
-  %arg_0:ptr<handle, texture_cube_array<u32>, read> = var @binding_point(1, 0)
-  %prevent_dce:ptr<storage, u32, read_write> = var @binding_point(2, 0)
+void textureNumLevels_fe2171(tint_module_vars_struct tint_module_vars) {
+  uint res = tint_module_vars.arg_0.get_num_mip_levels();
+  (*tint_module_vars.prevent_dce) = res;
 }
-
-%textureNumLevels_fe2171 = func():void {
-  $B2: {
-    %4:texture_cube_array<u32> = load %arg_0
-    %5:u32 = textureNumLevels %4
-    %res:ptr<function, u32, read_write> = var, %5
-    %7:u32 = load %res
-    store %prevent_dce, %7
-    ret
-  }
+float4 vertex_main_inner(tint_module_vars_struct tint_module_vars) {
+  textureNumLevels_fe2171(tint_module_vars);
+  return float4(0.0f);
 }
-%vertex_main = @vertex func():vec4<f32> [@position] {
-  $B3: {
-    %9:void = call %textureNumLevels_fe2171
-    ret vec4<f32>(0.0f)
-  }
+fragment void fragment_main(texturecube_array<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_fe2171(tint_module_vars);
 }
-%fragment_main = @fragment func():void {
-  $B4: {
-    %11:void = call %textureNumLevels_fe2171
-    ret
-  }
+kernel void compute_main(texturecube_array<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  textureNumLevels_fe2171(tint_module_vars);
 }
-%compute_main = @compute @workgroup_size(1, 1, 1) func():void {
-  $B5: {
-    %13:void = call %textureNumLevels_fe2171
-    ret
-  }
+vertex vertex_main_outputs vertex_main(texturecube_array<uint, access::sample> arg_0 [[texture(0)]], device uint* prevent_dce [[buffer(0)]]) {
+  tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.arg_0=arg_0, .prevent_dce=prevent_dce};
+  return vertex_main_outputs{.tint_symbol=vertex_main_inner(tint_module_vars)};
 }
-
-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/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_0.spvasm.expected.ir.msl b/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_0.spvasm.expected.ir.msl
deleted file mode 100644
index c341f8a..0000000
--- a/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_0.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: textureNumLevels
-********************************************************************
-*  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/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_1.spvasm.expected.ir.msl b/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_1.spvasm.expected.ir.msl
deleted file mode 100644
index c341f8a..0000000
--- a/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_1.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: textureNumLevels
-********************************************************************
-*  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/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_2.spvasm.expected.ir.msl b/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_2.spvasm.expected.ir.msl
deleted file mode 100644
index c341f8a..0000000
--- a/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_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: textureNumLevels
-********************************************************************
-*  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/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_3.spvasm.expected.ir.msl b/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_3.spvasm.expected.ir.msl
deleted file mode 100644
index c341f8a..0000000
--- a/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_3.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: textureNumLevels
-********************************************************************
-*  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/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_4.spvasm.expected.ir.msl b/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_4.spvasm.expected.ir.msl
deleted file mode 100644
index c341f8a..0000000
--- a/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_4.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: textureNumLevels
-********************************************************************
-*  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/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_5.spvasm.expected.ir.msl b/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_5.spvasm.expected.ir.msl
deleted file mode 100644
index c341f8a..0000000
--- a/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_5.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: textureNumLevels
-********************************************************************
-*  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/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_6.spvasm.expected.ir.msl b/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_6.spvasm.expected.ir.msl
deleted file mode 100644
index c341f8a..0000000
--- a/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_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: textureNumLevels
-********************************************************************
-*  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/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_7.spvasm.expected.ir.msl b/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_7.spvasm.expected.ir.msl
deleted file mode 100644
index c341f8a..0000000
--- a/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_7.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: textureNumLevels
-********************************************************************
-*  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/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_8.spvasm.expected.ir.msl b/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_8.spvasm.expected.ir.msl
deleted file mode 100644
index c341f8a..0000000
--- a/test/tint/unittest/reader/spirv/ImageQueryLevels_SignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_8.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: textureNumLevels
-********************************************************************
-*  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/ImageQueryLevels_UnsignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_0.spvasm.expected.ir.msl b/test/tint/unittest/reader/spirv/ImageQueryLevels_UnsignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_0.spvasm.expected.ir.msl
deleted file mode 100644
index c341f8a..0000000
--- a/test/tint/unittest/reader/spirv/ImageQueryLevels_UnsignedResult_SpvParserHandleTest_SampledImageAccessTest_Variable_0.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: textureNumLevels
-********************************************************************
-*  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_RawImage_Variable_5.spvasm.expected.ir.msl b/test/tint/unittest/reader/spirv/Samples_SpvParserHandleTest_RegisterHandleUsage_RawImage_Variable_5.spvasm.expected.ir.msl
deleted file mode 100644
index c341f8a..0000000
--- a/test/tint/unittest/reader/spirv/Samples_SpvParserHandleTest_RegisterHandleUsage_RawImage_Variable_5.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: textureNumLevels
-********************************************************************
-*  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.  *
-********************************************************************