diff --git a/src/tint/lang/msl/builtin_fn.cc b/src/tint/lang/msl/builtin_fn.cc
index bda3140..9bc5c80 100644
--- a/src/tint/lang/msl/builtin_fn.cc
+++ b/src/tint/lang/msl/builtin_fn.cc
@@ -94,6 +94,8 @@
             return "distance";
         case BuiltinFn::kDot:
             return "dot";
+        case BuiltinFn::kFmod:
+            return "fmod";
         case BuiltinFn::kLength:
             return "length";
         case BuiltinFn::kThreadgroupBarrier:
diff --git a/src/tint/lang/msl/builtin_fn.h b/src/tint/lang/msl/builtin_fn.h
index b0cac68..2ad0d28 100644
--- a/src/tint/lang/msl/builtin_fn.h
+++ b/src/tint/lang/msl/builtin_fn.h
@@ -73,6 +73,7 @@
     kWrite,
     kDistance,
     kDot,
+    kFmod,
     kLength,
     kThreadgroupBarrier,
     kNone,
diff --git a/src/tint/lang/msl/intrinsic/data.cc b/src/tint/lang/msl/intrinsic/data.cc
index 62e30c9..1272c5b 100644
--- a/src/tint/lang/msl/intrinsic/data.cc
+++ b/src/tint/lang/msl/intrinsic/data.cc
@@ -2610,6 +2610,16 @@
   },
   {
     /* [302] */
+    /* usage */ core::ParameterUsage::kNone,
+    /* matcher_indices */ MatcherIndicesIndex(3),
+  },
+  {
+    /* [303] */
+    /* usage */ core::ParameterUsage::kNone,
+    /* matcher_indices */ MatcherIndicesIndex(3),
+  },
+  {
+    /* [304] */
     /* usage */ core::ParameterUsage::kTexture,
     /* matcher_indices */ MatcherIndicesIndex(9),
   },
@@ -3557,7 +3567,7 @@
     /* num_explicit_templates */ 0,
     /* num_templates   */ 2,
     /* templates */ TemplateIndex(5),
-    /* parameters */ ParameterIndex(302),
+    /* parameters */ ParameterIndex(304),
     /* return_matcher_indices */ MatcherIndicesIndex(59),
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
@@ -4338,7 +4348,7 @@
     /* num_explicit_templates */ 0,
     /* num_templates   */ 2,
     /* templates */ TemplateIndex(5),
-    /* parameters */ ParameterIndex(302),
+    /* parameters */ ParameterIndex(304),
     /* return_matcher_indices */ MatcherIndicesIndex(/* invalid */),
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
@@ -4421,6 +4431,28 @@
   },
   {
     /* [156] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+    /* num_parameters */ 2,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 1,
+    /* templates */ TemplateIndex(12),
+    /* parameters */ ParameterIndex(302),
+    /* return_matcher_indices */ MatcherIndicesIndex(3),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [157] */
+    /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
+    /* num_parameters */ 2,
+    /* num_explicit_templates */ 0,
+    /* num_templates   */ 2,
+    /* templates */ TemplateIndex(11),
+    /* parameters */ ParameterIndex(300),
+    /* return_matcher_indices */ MatcherIndicesIndex(93),
+    /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
+  },
+  {
+    /* [158] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 5,
     /* num_explicit_templates */ 0,
@@ -4431,7 +4463,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [157] */
+    /* [159] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 3,
     /* num_explicit_templates */ 0,
@@ -4442,7 +4474,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [158] */
+    /* [160] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 2,
     /* num_explicit_templates */ 0,
@@ -4453,7 +4485,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [159] */
+    /* [161] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 3,
     /* num_explicit_templates */ 0,
@@ -4464,7 +4496,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [160] */
+    /* [162] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 2,
     /* num_explicit_templates */ 0,
@@ -4475,7 +4507,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [161] */
+    /* [163] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsVertexPipeline, OverloadFlag::kSupportsFragmentPipeline, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 1,
     /* num_explicit_templates */ 0,
@@ -4486,7 +4518,7 @@
     /* const_eval_fn */ ConstEvalFunctionIndex(/* invalid */),
   },
   {
-    /* [162] */
+    /* [164] */
     /* flags */ OverloadFlags(OverloadFlag::kIsBuiltin, OverloadFlag::kSupportsComputePipeline),
     /* num_parameters */ 1,
     /* num_explicit_templates */ 0,
@@ -4506,67 +4538,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(156),
+    /* overloads */ OverloadIndex(158),
   },
   {
     /* [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(157),
+    /* overloads */ OverloadIndex(159),
   },
   {
     /* [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(157),
+    /* overloads */ OverloadIndex(159),
   },
   {
     /* [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(157),
+    /* overloads */ OverloadIndex(159),
   },
   {
     /* [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(157),
+    /* overloads */ OverloadIndex(159),
   },
   {
     /* [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(157),
+    /* overloads */ OverloadIndex(159),
   },
   {
     /* [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(157),
+    /* overloads */ OverloadIndex(159),
   },
   {
     /* [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(157),
+    /* overloads */ OverloadIndex(159),
   },
   {
     /* [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(157),
+    /* overloads */ OverloadIndex(159),
   },
   {
     /* [9] */
     /* fn atomic_load_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, u32) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(158),
+    /* overloads */ OverloadIndex(160),
   },
   {
     /* [10] */
     /* fn atomic_store_explicit[T : iu32, S : workgroup_or_storage](ptr<S, atomic<T>, read_write>, T, u32) */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(159),
+    /* overloads */ OverloadIndex(161),
   },
   {
     /* [11] */
@@ -4793,25 +4825,32 @@
     /* [24] */
     /* fn distance[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(160),
+    /* overloads */ OverloadIndex(162),
   },
   {
     /* [25] */
     /* fn dot[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> T */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(160),
+    /* overloads */ OverloadIndex(162),
   },
   {
     /* [26] */
-    /* fn length[N : num, T : f32_f16](vec<N, T>) -> T */
-    /* num overloads */ 1,
-    /* overloads */ OverloadIndex(161),
+    /* fn fmod[T : f32_f16](T, T) -> T */
+    /* fn fmod[N : num, T : f32_f16](vec<N, T>, vec<N, T>) -> vec<N, T> */
+    /* num overloads */ 2,
+    /* overloads */ OverloadIndex(156),
   },
   {
     /* [27] */
+    /* fn length[N : num, T : f32_f16](vec<N, T>) -> T */
+    /* num overloads */ 1,
+    /* overloads */ OverloadIndex(163),
+  },
+  {
+    /* [28] */
     /* fn threadgroup_barrier(u32) */
     /* num overloads */ 1,
-    /* overloads */ OverloadIndex(162),
+    /* overloads */ OverloadIndex(164),
   },
 };
 
diff --git a/src/tint/lang/msl/msl.def b/src/tint/lang/msl/msl.def
index 3894e23..4c887be 100644
--- a/src/tint/lang/msl/msl.def
+++ b/src/tint/lang/msl/msl.def
@@ -323,6 +323,8 @@
 
 fn distance[N: num, T: f32_f16](vec<N, T>, vec<N, T>) -> T
 fn dot[N: num, T: f32_f16](vec<N, T>, vec<N, T>) -> T
+fn fmod[T: f32_f16](T, T) -> T
+fn fmod[N: num, T: f32_f16](vec<N, T>, vec<N, T>) -> vec<N, T>
 fn length[N: num, T: f32_f16](vec<N, T>) -> T
 @stage("compute") fn threadgroup_barrier(u32)
 
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
index 0d2e61b..c4ea261 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
@@ -80,8 +80,9 @@
 
     /// Process the module.
     void Process() {
-        // Find the builtins that need replacing.
-        Vector<core::ir::CoreBuiltinCall*, 4> worklist;
+        // Find the builtins and binary operators that need replacing.
+        Vector<core::ir::CoreBinary*, 4> fmod_worklist;
+        Vector<core::ir::CoreBuiltinCall*, 4> builtin_worklist;
         for (auto* inst : ir.Instructions()) {
             if (auto* builtin = inst->As<core::ir::CoreBuiltinCall>()) {
                 switch (builtin->Func()) {
@@ -118,16 +119,21 @@
                     case core::BuiltinFn::kWorkgroupBarrier:
                     case core::BuiltinFn::kTextureBarrier:
                     case core::BuiltinFn::kUnpack2X16Float:
-                        worklist.Push(builtin);
+                        builtin_worklist.Push(builtin);
                         break;
                     default:
                         break;
                 }
+            } else if (auto* binary = inst->As<core::ir::CoreBinary>()) {
+                if (binary->Op() == core::BinaryOp::kModulo &&
+                    binary->LHS()->Type()->is_float_scalar_or_vector()) {
+                    fmod_worklist.Push(binary);
+                }
             }
         }
 
         // Replace the builtins that we found.
-        for (auto* builtin : worklist) {
+        for (auto* builtin : builtin_worklist) {
             switch (builtin->Func()) {
                 // Atomics.
                 case core::BuiltinFn::kAtomicAdd:
@@ -242,6 +248,11 @@
                     break;
             }
         }
+
+        // Replace the fmod instructions that we found.
+        for (auto* fmod : fmod_worklist) {
+            FMod(fmod);
+        }
     }
 
     /// Replace an atomic builtin call with an equivalent MSL intrinsic.
@@ -778,6 +789,15 @@
         });
         builtin->Destroy();
     }
+
+    /// Replace a floating point modulo binary instruction with the equivalent MSL intrinsic.
+    /// @param binary the float point modulo binary instruction
+    void FMod(core::ir::CoreBinary* binary) {
+        auto* call = b.CallWithResult<msl::ir::BuiltinCall>(
+            binary->DetachResult(), msl::BuiltinFn::kFmod, binary->Operands());
+        call->InsertBefore(binary);
+        binary->Destroy();
+    }
 };
 
 }  // namespace
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 467bb5e..6280415 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill_test.cc
@@ -2872,5 +2872,73 @@
     EXPECT_EQ(expect, str());
 }
 
+TEST_F(MslWriter_BuiltinPolyfillTest, FMod_Scalar) {
+    auto* lhs = b.FunctionParam<f32>("lhs");
+    auto* rhs = b.FunctionParam<f32>("rhs");
+    auto* func = b.Function("foo", ty.f32());
+    func->SetParams({lhs, rhs});
+    b.Append(func->Block(), [&] {
+        auto* result = b.Modulo<f32>(lhs, rhs);
+        b.Return(func, result);
+    });
+
+    auto* src = R"(
+%foo = func(%lhs:f32, %rhs:f32):f32 {
+  $B1: {
+    %4:f32 = mod %lhs, %rhs
+    ret %4
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%foo = func(%lhs:f32, %rhs:f32):f32 {
+  $B1: {
+    %4:f32 = msl.fmod %lhs, %rhs
+    ret %4
+  }
+}
+)";
+
+    Run(BuiltinPolyfill);
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BuiltinPolyfillTest, FMod_Vector) {
+    auto* lhs = b.FunctionParam<vec4<f16>>("lhs");
+    auto* rhs = b.FunctionParam<vec4<f16>>("rhs");
+    auto* func = b.Function("foo", ty.vec4<f16>());
+    func->SetParams({lhs, rhs});
+    b.Append(func->Block(), [&] {
+        auto* result = b.Modulo<vec4<f16>>(lhs, rhs);
+        b.Return(func, result);
+    });
+
+    auto* src = R"(
+%foo = func(%lhs:vec4<f16>, %rhs:vec4<f16>):vec4<f16> {
+  $B1: {
+    %4:vec4<f16> = mod %lhs, %rhs
+    ret %4
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+%foo = func(%lhs:vec4<f16>, %rhs:vec4<f16>):vec4<f16> {
+  $B1: {
+    %4:vec4<f16> = msl.fmod %lhs, %rhs
+    ret %4
+  }
+}
+)";
+
+    Run(BuiltinPolyfill);
+
+    EXPECT_EQ(expect, str());
+}
+
 }  // namespace
 }  // namespace tint::msl::writer::raise
diff --git a/test/tint/bug/chromium/1434271.wgsl.expected.ir.msl b/test/tint/bug/chromium/1434271.wgsl.expected.ir.msl
index 028ef20..e15baf4 100644
--- a/test/tint/bug/chromium/1434271.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/1434271.wgsl.expected.ir.msl
@@ -141,7 +141,7 @@
     float const c = (*tint_module_vars.buf_in).weights[((src_offset + 1u) + (*tint_module_vars.ubo).width)];
     float const d = (*tint_module_vars.buf_in).weights[((src_offset + 1u) + (*tint_module_vars.ubo).width)];
     float const sum = dot(float4(a, b, c, d), float4(1.0f));
-    (*tint_module_vars.buf_out).weights[dst_offset] = (sum % 4.0f);
+    (*tint_module_vars.buf_out).weights[dst_offset] = fmod(sum, 4.0f);
     float4 const v_4 = float4(a, (a * b), ((a / b) + c), sum);
     float4 const probabilities = (v_4 + max(sum, 0.0f));
     tint_module_vars.tex_out.write(probabilities, uint2(int2(coord.xy)));
@@ -171,9 +171,6 @@
 program_source:89:8: warning: unused variable 'res' [-Wunused-variable]
   half res = asinh(arg_0);
        ^
-program_source:142:60: error: invalid operands to binary expression ('const float' and 'float')
-    (*tint_module_vars.buf_out).weights[dst_offset] = (sum % 4.0f);
-                                                       ~~~ ^ ~~~~
 program_source:154:104: error: call to deleted constructor of 'texture1d<float, access::sample>'
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.render_params=render_params};
                                                                                                        ^
diff --git a/test/tint/bug/tint/948.wgsl.expected.ir.msl b/test/tint/bug/tint/948.wgsl.expected.ir.msl
index 65f1035..2d0fd9c 100644
--- a/test/tint/bug/tint/948.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/948.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
@@ -144,7 +142,7 @@
       if ((x_174 > 0.0f)) {
         float const x_181 = (*tint_module_vars.x_20).time;
         float const x_184 = animationData[2u];
-        (*tint_module_vars.mt) = ((x_181 * x_184) % 1.0f);
+        (*tint_module_vars.mt) = fmod((x_181 * x_184), 1.0f);
         f = 0.0f;
         {
           while(true) {
@@ -257,16 +255,3 @@
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.x_20=x_20, .frameMapTexture=frameMapTexture, .frameMapSampler=frameMapSampler, .tUV=(&tUV), .tileMapsTexture0=tileMapsTexture0, .tileMapsSampler=tileMapsSampler, .tileMapsTexture1=tileMapsTexture1, .animationMapTexture=animationMapTexture, .animationMapSampler=animationMapSampler, .mt=(&mt), .spriteSheetTexture=spriteSheetTexture, .spriteSheetSampler=spriteSheetSampler, .glFragColor=(&glFragColor), .tileID_1=(&tileID_1), .levelUnits=(&levelUnits), .stageUnits_1=(&stageUnits_1), .vPosition=(&vPosition), .vUV=(&vUV)};
   return tint_symbol_outputs{.main_out_glFragColor_1=tint_symbol_inner(inputs.tUV_param, inputs.tileID_1_param, inputs.levelUnits_param, inputs.stageUnits_1_param, inputs.vPosition_param, inputs.vUV_param, tint_module_vars).glFragColor_1};
 }
-program_source:145:51: error: invalid operands to binary expression ('float' and 'float')
-        (*tint_module_vars.mt) = ((x_181 * x_184) % 1.0f);
-                                  ~~~~~~~~~~~~~~~ ^ ~~~~
-program_source:161:25: warning: unused variable 'x_208' [-Wunused-variable]
-            float const x_208 = frameID_1;
-                        ^
-program_source:162:25: warning: unused variable 'x_211' [-Wunused-variable]
-            float const x_211 = (*tint_module_vars.x_20).spriteCount;
-                        ^
-program_source:163:25: warning: unused variable 'x_214' [-Wunused-variable]
-            float const x_214 = f;
-                        ^
-
diff --git a/test/tint/expressions/binary/mod/scalar-scalar/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/scalar-scalar/f16.wgsl.expected.ir.msl
index 812cbe6..fa2c3ee 100644
--- a/test/tint/expressions/binary/mod/scalar-scalar/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/scalar-scalar/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   half const a = 1.0h;
   half const b = 2.0h;
-  half const r = (a % b);
+  half const r = fmod(a, b);
 }
-program_source:7:21: error: invalid operands to binary expression ('const half' and 'const half')
-  half const r = (a % b);
-                  ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod/scalar-scalar/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/scalar-scalar/f32.wgsl.expected.ir.msl
index c55eb34..7360105 100644
--- a/test/tint/expressions/binary/mod/scalar-scalar/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/scalar-scalar/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   float const a = 1.0f;
   float const b = 2.0f;
-  float const r = (a % b);
+  float const r = fmod(a, b);
 }
-program_source:7:22: error: invalid operands to binary expression ('const float' and 'const float')
-  float const r = (a % b);
-                   ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/scalar-vec3/f16.wgsl.expected.ir.msl
index 4e7e3dd..9d36d1e 100644
--- a/test/tint/expressions/binary/mod/scalar-vec3/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/scalar-vec3/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   half const a = 4.0h;
   half3 const b = half3(1.0h, 2.0h, 3.0h);
-  half3 const r = (a % b);
+  half3 const r = fmod(a, b);
 }
-program_source:7:22: error: invalid operands to binary expression ('const half' and 'const half3' (vector of 3 'half' values))
-  half3 const r = (a % b);
-                   ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.ir.msl
index c20de57..00748aa 100644
--- a/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/scalar-vec3/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   float const a = 4.0f;
   float3 const b = float3(1.0f, 2.0f, 3.0f);
-  float3 const r = (a % b);
+  float3 const r = fmod(a, b);
 }
-program_source:7:23: error: invalid operands to binary expression ('const float' and 'const float3' (vector of 3 'float' values))
-  float3 const r = (a % b);
-                    ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/vec3-scalar/f16.wgsl.expected.ir.msl
index 339b67d..eb99c54 100644
--- a/test/tint/expressions/binary/mod/vec3-scalar/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/vec3-scalar/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   half3 const a = half3(1.0h, 2.0h, 3.0h);
   half const b = 4.0h;
-  half3 const r = (a % b);
+  half3 const r = fmod(a, b);
 }
-program_source:7:22: error: invalid operands to binary expression ('const half3' (vector of 3 'half' values) and 'const half')
-  half3 const r = (a % b);
-                   ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.ir.msl
index fcd961f..3e9ccce 100644
--- a/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/vec3-scalar/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   float3 const a = float3(1.0f, 2.0f, 3.0f);
   float const b = 4.0f;
-  float3 const r = (a % b);
+  float3 const r = fmod(a, b);
 }
-program_source:7:23: error: invalid operands to binary expression ('const float3' (vector of 3 'float' values) and 'const float')
-  float3 const r = (a % b);
-                    ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod/vec3-vec3/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/vec3-vec3/f16.wgsl.expected.ir.msl
index 5eb95c9..c331b93 100644
--- a/test/tint/expressions/binary/mod/vec3-vec3/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/vec3-vec3/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   half3 const a = half3(1.0h, 2.0h, 3.0h);
   half3 const b = half3(4.0h, 5.0h, 6.0h);
-  half3 const r = (a % b);
+  half3 const r = fmod(a, b);
 }
-program_source:7:22: error: invalid operands to binary expression ('const half3' (vector of 3 'half' values) and 'const half3')
-  half3 const r = (a % b);
-                   ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod/vec3-vec3/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/vec3-vec3/f32.wgsl.expected.ir.msl
index 7ad70c2..c1fcab3 100644
--- a/test/tint/expressions/binary/mod/vec3-vec3/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/vec3-vec3/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   float3 const a = float3(1.0f, 2.0f, 3.0f);
   float3 const b = float3(4.0f, 5.0f, 6.0f);
-  float3 const r = (a % b);
+  float3 const r = fmod(a, b);
 }
-program_source:7:23: error: invalid operands to binary expression ('const float3' (vector of 3 'float' values) and 'const float3')
-  float3 const r = (a % b);
-                    ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f16.wgsl.expected.ir.msl
index 3883afa..d169baa 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   half const a = 1.0h;
   half const b = 0.0h;
-  half const r = (a % b);
+  half const r = fmod(a, b);
 }
-program_source:7:21: error: invalid operands to binary expression ('const half' and 'const half')
-  half const r = (a % b);
-                  ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.ir.msl
index f4e0d01..d46406f 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   float const a = 1.0f;
   float const b = 0.0f;
-  float const r = (a % b);
+  float const r = fmod(a, b);
 }
-program_source:7:22: error: invalid operands to binary expression ('const float' and 'const float')
-  float const r = (a % b);
-                   ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f16.wgsl.expected.ir.msl
index 500fa80..8826b2a 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   half3 const a = half3(1.0h, 2.0h, 3.0h);
   half3 const b = half3(0.0h, 5.0h, 0.0h);
-  half3 const r = (a % b);
+  half3 const r = fmod(a, b);
 }
-program_source:7:22: error: invalid operands to binary expression ('const half3' (vector of 3 'half' values) and 'const half3')
-  half3 const r = (a % b);
-                   ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.ir.msl
index b771d78..6496f21 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   float3 const a = float3(1.0f, 2.0f, 3.0f);
   float3 const b = float3(0.0f, 5.0f, 0.0f);
-  float3 const r = (a % b);
+  float3 const r = fmod(a, b);
 }
-program_source:7:23: error: invalid operands to binary expression ('const float3' (vector of 3 'float' values) and 'const float3')
-  float3 const r = (a % b);
-                    ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f16.wgsl.expected.ir.msl
index 347f215..cfb7230 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   half a = 1.0h;
   half b = 0.0h;
-  half const r = (a % (b + b));
+  half const r = fmod(a, (b + b));
 }
-program_source:7:21: error: invalid operands to binary expression ('half' and 'half')
-  half const r = (a % (b + b));
-                  ~ ^ ~~~~~~~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.ir.msl
index 5b44ae7..2a021b0 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   float a = 1.0f;
   float b = 0.0f;
-  float const r = (a % (b + b));
+  float const r = fmod(a, (b + b));
 }
-program_source:7:22: error: invalid operands to binary expression ('float' and 'float')
-  float const r = (a % (b + b));
-                   ~ ^ ~~~~~~~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f16.wgsl.expected.ir.msl
index cfcd9fa..5cb5c8d 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   half3 a = half3(1.0h, 2.0h, 3.0h);
   half3 b = half3(0.0h, 5.0h, 0.0h);
-  half3 const r = (a % (b + b));
+  half3 const r = fmod(a, (b + b));
 }
-program_source:7:22: error: invalid operands to binary expression ('half3' (vector of 3 'half' values) and 'half3')
-  half3 const r = (a % (b + b));
-                   ~ ^ ~~~~~~~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.ir.msl
index 3238577..b6f792c 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   float3 a = float3(1.0f, 2.0f, 3.0f);
   float3 b = float3(0.0f, 5.0f, 0.0f);
-  float3 const r = (a % (b + b));
+  float3 const r = fmod(a, (b + b));
 }
-program_source:7:23: error: invalid operands to binary expression ('float3' (vector of 3 'float' values) and 'float3')
-  float3 const r = (a % (b + b));
-                    ~ ^ ~~~~~~~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f16.wgsl.expected.ir.msl
index b925c09..eddd331 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   half a = 1.0h;
   half b = 0.0h;
-  half const r = (a % b);
+  half const r = fmod(a, b);
 }
-program_source:7:21: error: invalid operands to binary expression ('half' and 'half')
-  half const r = (a % b);
-                  ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.ir.msl
index 730fb96..bf6392d 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   float a = 1.0f;
   float b = 0.0f;
-  float const r = (a % b);
+  float const r = fmod(a, b);
 }
-program_source:7:22: error: invalid operands to binary expression ('float' and 'float')
-  float const r = (a % b);
-                   ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f16.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f16.wgsl.expected.ir.msl
index a81f54b..b04f111 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f16.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f16.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   half3 a = half3(1.0h, 2.0h, 3.0h);
   half3 b = half3(0.0h, 5.0h, 0.0h);
-  half3 const r = (a % b);
+  half3 const r = fmod(a, b);
 }
-program_source:7:22: error: invalid operands to binary expression ('half3' (vector of 3 'half' values) and 'half3')
-  half3 const r = (a % b);
-                   ~ ^ ~
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.ir.msl
index 0dbb18d..29578fd 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/f32.wgsl.expected.ir.msl
@@ -1,14 +1,8 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
 kernel void f() {
   float3 a = float3(1.0f, 2.0f, 3.0f);
   float3 b = float3(0.0f, 5.0f, 0.0f);
-  float3 const r = (a % b);
+  float3 const r = fmod(a, b);
 }
-program_source:7:23: error: invalid operands to binary expression ('float3' (vector of 3 'float' values) and 'float3')
-  float3 const r = (a % b);
-                    ~ ^ ~
-
diff --git a/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.ir.msl b/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.ir.msl
index 4ab3cb3..8494843 100644
--- a/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.ir.msl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
@@ -23,16 +21,9 @@
   (*tint_module_vars.a) = tint_div_i32((*tint_module_vars.a), maybe_zero);
   (*tint_module_vars.a) = tint_mod_i32((*tint_module_vars.a), maybe_zero);
   (*tint_module_vars.b) = ((*tint_module_vars.b) / 0.0f);
-  (*tint_module_vars.b) = ((*tint_module_vars.b) % 0.0f);
+  (*tint_module_vars.b) = fmod((*tint_module_vars.b), 0.0f);
   float const v_1 = float(maybe_zero);
   (*tint_module_vars.b) = ((*tint_module_vars.b) / v_1);
   float const v_2 = float(maybe_zero);
-  (*tint_module_vars.b) = ((*tint_module_vars.b) % v_2);
+  (*tint_module_vars.b) = fmod((*tint_module_vars.b), v_2);
 }
-program_source:24:50: error: invalid operands to binary expression ('float' and 'float')
-  (*tint_module_vars.b) = ((*tint_module_vars.b) % 0.0f);
-                           ~~~~~~~~~~~~~~~~~~~~~ ^ ~~~~
-program_source:28:50: error: invalid operands to binary expression ('float' and 'const float')
-  (*tint_module_vars.b) = ((*tint_module_vars.b) % v_2);
-                           ~~~~~~~~~~~~~~~~~~~~~ ^ ~~~
-