diff --git a/src/tint/lang/msl/writer/printer/binary_test.cc b/src/tint/lang/msl/writer/printer/binary_test.cc
index 495b156..fe0a464 100644
--- a/src/tint/lang/msl/writer/printer/binary_test.cc
+++ b/src/tint/lang/msl/writer/printer/binary_test.cc
@@ -91,14 +91,14 @@
 
     ASSERT_TRUE(Generate()) << err_ << output_;
     EXPECT_EQ(output_, MetalHeader() + R"(
+uint tint_div_u32(uint lhs, uint rhs) {
+  return (lhs / select(rhs, 1u, (rhs == 0u)));
+}
 void foo() {
   uint const left = 1u;
   uint const right = 2u;
   uint const val = tint_div_u32(left, right);
 }
-uint tint_div_u32(uint lhs, uint rhs) {
-  return (lhs / select(rhs, 1u, (rhs == 0u)));
-}
 )");
 }
 
@@ -114,15 +114,15 @@
 
     ASSERT_TRUE(Generate()) << err_ << output_;
     EXPECT_EQ(output_, MetalHeader() + R"(
+uint tint_mod_u32(uint lhs, uint rhs) {
+  uint const v = select(rhs, 1u, (rhs == 0u));
+  return (lhs - ((lhs / v) * v));
+}
 void foo() {
   uint const left = 1u;
   uint const right = 2u;
   uint const val = tint_mod_u32(left, right);
 }
-uint tint_mod_u32(uint lhs, uint rhs) {
-  uint const v = select(rhs, 1u, (rhs == 0u));
-  return (lhs - ((lhs / v) * v));
-}
 )");
 }
 
diff --git a/src/tint/lang/msl/writer/printer/printer.cc b/src/tint/lang/msl/writer/printer/printer.cc
index c39b030..40ffcdf 100644
--- a/src/tint/lang/msl/writer/printer/printer.cc
+++ b/src/tint/lang/msl/writer/printer/printer.cc
@@ -122,7 +122,7 @@
         EmitBlockInstructions(ir_.root_block);
 
         // Emit functions.
-        for (auto& func : ir_.functions) {
+        for (auto* func : ir_.DependencyOrderedFunctions()) {
             EmitFunction(func);
         }
 
diff --git a/test/tint/bug/chromium/1386647.wgsl.expected.ir.msl b/test/tint/bug/chromium/1386647.wgsl.expected.ir.msl
index 288ce4b..69ab288 100644
--- a/test/tint/bug/chromium/1386647.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/1386647.wgsl.expected.ir.msl
@@ -1,16 +1,10 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
-kernel void f(uint3 v [[thread_position_in_grid]]) {
-  uint const l = (v[0u] << (tint_mod_u32(v[1u], 1u) & 31u));
-}
 uint tint_mod_u32(uint lhs, uint rhs) {
   uint const v_1 = select(rhs, 1u, (rhs == 0u));
   return (lhs - ((lhs / v_1) * v_1));
 }
-program_source:5:29: error: use of undeclared identifier 'tint_mod_u32'
+kernel void f(uint3 v [[thread_position_in_grid]]) {
   uint const l = (v[0u] << (tint_mod_u32(v[1u], 1u) & 31u));
-                            ^
-
+}
diff --git a/test/tint/bug/tint/1083.wgsl.expected.ir.msl b/test/tint/bug/tint/1083.wgsl.expected.ir.msl
index 75a514b..90d70be 100644
--- a/test/tint/bug/tint/1083.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1083.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int tint_div_i32(int lhs, int rhs) {
+  return (lhs / select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1)))));
+}
 kernel void f() {
   int const a = 1;
   int const b = 0;
   int const c = tint_div_i32(a, b);
 }
-int tint_div_i32(int lhs, int rhs) {
-  return (lhs / select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1)))));
-}
-program_source:7:17: error: use of undeclared identifier 'tint_div_i32'
-  int const c = tint_div_i32(a, b);
-                ^
-
diff --git a/test/tint/expressions/binary/div/scalar-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div/scalar-scalar/i32.wgsl.expected.ir.msl
index cd92c60..d4968f9 100644
--- a/test/tint/expressions/binary/div/scalar-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div/scalar-scalar/i32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int tint_div_i32(int lhs, int rhs) {
+  return (lhs / select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1)))));
+}
 kernel void f() {
   int const a = 1;
   int const b = 2;
   int const r = tint_div_i32(a, b);
 }
-int tint_div_i32(int lhs, int rhs) {
-  return (lhs / select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1)))));
-}
-program_source:7:17: error: use of undeclared identifier 'tint_div_i32'
-  int const r = tint_div_i32(a, b);
-                ^
-
diff --git a/test/tint/expressions/binary/div/scalar-scalar/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div/scalar-scalar/u32.wgsl.expected.ir.msl
index 048f61e..2d1b8b5 100644
--- a/test/tint/expressions/binary/div/scalar-scalar/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div/scalar-scalar/u32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint tint_div_u32(uint lhs, uint rhs) {
+  return (lhs / select(rhs, 1u, (rhs == 0u)));
+}
 kernel void f() {
   uint const a = 1u;
   uint const b = 2u;
   uint const r = tint_div_u32(a, b);
 }
-uint tint_div_u32(uint lhs, uint rhs) {
-  return (lhs / select(rhs, 1u, (rhs == 0u)));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_div_u32'
-  uint const r = tint_div_u32(a, b);
-                 ^
-
diff --git a/test/tint/expressions/binary/div/scalar-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div/scalar-vec3/i32.wgsl.expected.ir.msl
index 2fa961a..7065a45 100644
--- a/test/tint/expressions/binary/div/scalar-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div/scalar-vec3/i32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_div_v3i32(int3 lhs, int3 rhs) {
+  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
+}
 kernel void f() {
   int const a = 4;
   int3 const b = int3(1, 2, 3);
   int3 const r = tint_div_v3i32(int3(a), b);
 }
-int3 tint_div_v3i32(int3 lhs, int3 rhs) {
-  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_div_v3i32'
-  int3 const r = tint_div_v3i32(int3(a), b);
-                 ^
-
diff --git a/test/tint/expressions/binary/div/scalar-vec3/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div/scalar-vec3/u32.wgsl.expected.ir.msl
index e836cce..f4fe943 100644
--- a/test/tint/expressions/binary/div/scalar-vec3/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div/scalar-vec3/u32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
+  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
+}
 kernel void f() {
   uint const a = 4u;
   uint3 const b = uint3(1u, 2u, 3u);
   uint3 const r = tint_div_v3u32(uint3(a), b);
 }
-uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
-  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
-}
-program_source:7:19: error: use of undeclared identifier 'tint_div_v3u32'
-  uint3 const r = tint_div_v3u32(uint3(a), b);
-                  ^
-
diff --git a/test/tint/expressions/binary/div/vec3-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div/vec3-scalar/i32.wgsl.expected.ir.msl
index 574b1c1..92081a4 100644
--- a/test/tint/expressions/binary/div/vec3-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div/vec3-scalar/i32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_div_v3i32(int3 lhs, int3 rhs) {
+  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
+}
 kernel void f() {
   int3 const a = int3(1, 2, 3);
   int const b = 4;
   int3 const r = tint_div_v3i32(a, int3(b));
 }
-int3 tint_div_v3i32(int3 lhs, int3 rhs) {
-  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_div_v3i32'
-  int3 const r = tint_div_v3i32(a, int3(b));
-                 ^
-
diff --git a/test/tint/expressions/binary/div/vec3-scalar/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div/vec3-scalar/u32.wgsl.expected.ir.msl
index db9d145..1c9c0a2 100644
--- a/test/tint/expressions/binary/div/vec3-scalar/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div/vec3-scalar/u32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
+  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
+}
 kernel void f() {
   uint3 const a = uint3(1u, 2u, 3u);
   uint const b = 4u;
   uint3 const r = tint_div_v3u32(a, uint3(b));
 }
-uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
-  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
-}
-program_source:7:19: error: use of undeclared identifier 'tint_div_v3u32'
-  uint3 const r = tint_div_v3u32(a, uint3(b));
-                  ^
-
diff --git a/test/tint/expressions/binary/div/vec3-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div/vec3-vec3/i32.wgsl.expected.ir.msl
index 2755ecd..5b3e9c3 100644
--- a/test/tint/expressions/binary/div/vec3-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div/vec3-vec3/i32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_div_v3i32(int3 lhs, int3 rhs) {
+  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
+}
 kernel void f() {
   int3 const a = int3(1, 2, 3);
   int3 const b = int3(4, 5, 6);
   int3 const r = tint_div_v3i32(a, b);
 }
-int3 tint_div_v3i32(int3 lhs, int3 rhs) {
-  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_div_v3i32'
-  int3 const r = tint_div_v3i32(a, b);
-                 ^
-
diff --git a/test/tint/expressions/binary/div/vec3-vec3/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div/vec3-vec3/u32.wgsl.expected.ir.msl
index 1baed17..ece9968 100644
--- a/test/tint/expressions/binary/div/vec3-vec3/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div/vec3-vec3/u32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
+  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
+}
 kernel void f() {
   uint3 const a = uint3(1u, 2u, 3u);
   uint3 const b = uint3(4u, 5u, 6u);
   uint3 const r = tint_div_v3u32(a, b);
 }
-uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
-  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
-}
-program_source:7:19: error: use of undeclared identifier 'tint_div_v3u32'
-  uint3 const r = tint_div_v3u32(a, b);
-                  ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.ir.msl
index 2266c21..ed490a4 100644
--- a/test/tint/expressions/binary/div_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int tint_div_i32(int lhs, int rhs) {
+  return (lhs / select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1)))));
+}
 kernel void f() {
   int const a = 1;
   int const b = 0;
   int const r = tint_div_i32(a, b);
 }
-int tint_div_i32(int lhs, int rhs) {
-  return (lhs / select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1)))));
-}
-program_source:7:17: error: use of undeclared identifier 'tint_div_i32'
-  int const r = tint_div_i32(a, b);
-                ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.ir.msl
index 738eeab..51a5e13 100644
--- a/test/tint/expressions/binary/div_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint tint_div_u32(uint lhs, uint rhs) {
+  return (lhs / select(rhs, 1u, (rhs == 0u)));
+}
 kernel void f() {
   uint const a = 1u;
   uint const b = 0u;
   uint const r = tint_div_u32(a, b);
 }
-uint tint_div_u32(uint lhs, uint rhs) {
-  return (lhs / select(rhs, 1u, (rhs == 0u)));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_div_u32'
-  uint const r = tint_div_u32(a, b);
-                 ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.ir.msl
index 162982e..4bd3cbe 100644
--- a/test/tint/expressions/binary/div_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_div_v3i32(int3 lhs, int3 rhs) {
+  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
+}
 kernel void f() {
   int const a = 4;
   int3 const b = int3(0, 2, 0);
   int3 const r = tint_div_v3i32(int3(a), b);
 }
-int3 tint_div_v3i32(int3 lhs, int3 rhs) {
-  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_div_v3i32'
-  int3 const r = tint_div_v3i32(int3(a), b);
-                 ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.ir.msl
index 0231a4e..b09d7b0 100644
--- a/test/tint/expressions/binary/div_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
+  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
+}
 kernel void f() {
   uint const a = 4u;
   uint3 const b = uint3(0u, 2u, 0u);
   uint3 const r = tint_div_v3u32(uint3(a), b);
 }
-uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
-  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
-}
-program_source:7:19: error: use of undeclared identifier 'tint_div_v3u32'
-  uint3 const r = tint_div_v3u32(uint3(a), b);
-                  ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.ir.msl
index b96f324..d8abc3a 100644
--- a/test/tint/expressions/binary/div_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_div_v3i32(int3 lhs, int3 rhs) {
+  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
+}
 kernel void f() {
   int3 const a = int3(1, 2, 3);
   int const b = 0;
   int3 const r = tint_div_v3i32(a, int3(b));
 }
-int3 tint_div_v3i32(int3 lhs, int3 rhs) {
-  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_div_v3i32'
-  int3 const r = tint_div_v3i32(a, int3(b));
-                 ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.ir.msl
index 8904fcc..98e6493 100644
--- a/test/tint/expressions/binary/div_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
+  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
+}
 kernel void f() {
   uint3 const a = uint3(1u, 2u, 3u);
   uint const b = 0u;
   uint3 const r = tint_div_v3u32(a, uint3(b));
 }
-uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
-  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
-}
-program_source:7:19: error: use of undeclared identifier 'tint_div_v3u32'
-  uint3 const r = tint_div_v3u32(a, uint3(b));
-                  ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.ir.msl
index 3a14a9d..890eced 100644
--- a/test/tint/expressions/binary/div_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_div_v3i32(int3 lhs, int3 rhs) {
+  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
+}
 kernel void f() {
   int3 const a = int3(1, 2, 3);
   int3 const b = int3(0, 5, 0);
   int3 const r = tint_div_v3i32(a, b);
 }
-int3 tint_div_v3i32(int3 lhs, int3 rhs) {
-  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_div_v3i32'
-  int3 const r = tint_div_v3i32(a, b);
-                 ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.ir.msl
index 3c7663c..5b17c76 100644
--- a/test/tint/expressions/binary/div_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
+  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
+}
 kernel void f() {
   uint3 const a = uint3(1u, 2u, 3u);
   uint3 const b = uint3(0u, 5u, 0u);
   uint3 const r = tint_div_v3u32(a, b);
 }
-uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
-  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
-}
-program_source:7:19: error: use of undeclared identifier 'tint_div_v3u32'
-  uint3 const r = tint_div_v3u32(a, b);
-                  ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.ir.msl
index 0c3ce11..1f2f7b2 100644
--- a/test/tint/expressions/binary/div_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int tint_div_i32(int lhs, int rhs) {
+  return (lhs / select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1)))));
+}
 kernel void f() {
   int a = 1;
   int b = 0;
   int const r = tint_div_i32(a, (b + b));
 }
-int tint_div_i32(int lhs, int rhs) {
-  return (lhs / select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1)))));
-}
-program_source:7:17: error: use of undeclared identifier 'tint_div_i32'
-  int const r = tint_div_i32(a, (b + b));
-                ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.ir.msl
index d0fe939..41dac77 100644
--- a/test/tint/expressions/binary/div_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint tint_div_u32(uint lhs, uint rhs) {
+  return (lhs / select(rhs, 1u, (rhs == 0u)));
+}
 kernel void f() {
   uint a = 1u;
   uint b = 0u;
   uint const r = tint_div_u32(a, (b + b));
 }
-uint tint_div_u32(uint lhs, uint rhs) {
-  return (lhs / select(rhs, 1u, (rhs == 0u)));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_div_u32'
-  uint const r = tint_div_u32(a, (b + b));
-                 ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.ir.msl
index 885624b..8884b3d 100644
--- a/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_div_v3i32(int3 lhs, int3 rhs) {
+  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
+}
 kernel void f() {
   int a = 4;
   int3 b = int3(0, 2, 0);
   int3 const v = (b + b);
   int3 const r = tint_div_v3i32(int3(a), v);
 }
-int3 tint_div_v3i32(int3 lhs, int3 rhs) {
-  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
-}
-program_source:8:18: error: use of undeclared identifier 'tint_div_v3i32'
-  int3 const r = tint_div_v3i32(int3(a), v);
-                 ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.ir.msl
index 8608a64..f29db42 100644
--- a/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
+  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
+}
 kernel void f() {
   uint a = 4u;
   uint3 b = uint3(0u, 2u, 0u);
   uint3 const v = (b + b);
   uint3 const r = tint_div_v3u32(uint3(a), v);
 }
-uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
-  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
-}
-program_source:8:19: error: use of undeclared identifier 'tint_div_v3u32'
-  uint3 const r = tint_div_v3u32(uint3(a), v);
-                  ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.ir.msl
index 45a199e..54107b1 100644
--- a/test/tint/expressions/binary/div_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_div_v3i32(int3 lhs, int3 rhs) {
+  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
+}
 kernel void f() {
   int3 a = int3(1, 2, 3);
   int b = 0;
   int3 const v = a;
   int3 const r = tint_div_v3i32(v, int3((b + b)));
 }
-int3 tint_div_v3i32(int3 lhs, int3 rhs) {
-  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
-}
-program_source:8:18: error: use of undeclared identifier 'tint_div_v3i32'
-  int3 const r = tint_div_v3i32(v, int3((b + b)));
-                 ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.ir.msl
index 230534b..42a2e8c 100644
--- a/test/tint/expressions/binary/div_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
+  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
+}
 kernel void f() {
   uint3 a = uint3(1u, 2u, 3u);
   uint b = 0u;
   uint3 const v = a;
   uint3 const r = tint_div_v3u32(v, uint3((b + b)));
 }
-uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
-  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
-}
-program_source:8:19: error: use of undeclared identifier 'tint_div_v3u32'
-  uint3 const r = tint_div_v3u32(v, uint3((b + b)));
-                  ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.ir.msl
index 84fd79e..5de7811 100644
--- a/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_div_v3i32(int3 lhs, int3 rhs) {
+  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
+}
 kernel void f() {
   int3 a = int3(1, 2, 3);
   int3 b = int3(0, 5, 0);
   int3 const r = tint_div_v3i32(a, (b + b));
 }
-int3 tint_div_v3i32(int3 lhs, int3 rhs) {
-  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_div_v3i32'
-  int3 const r = tint_div_v3i32(a, (b + b));
-                 ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.ir.msl
index f7d3fe3..17afe14 100644
--- a/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
+  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
+}
 kernel void f() {
   uint3 a = uint3(1u, 2u, 3u);
   uint3 b = uint3(0u, 5u, 0u);
   uint3 const r = tint_div_v3u32(a, (b + b));
 }
-uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
-  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
-}
-program_source:7:19: error: use of undeclared identifier 'tint_div_v3u32'
-  uint3 const r = tint_div_v3u32(a, (b + b));
-                  ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.ir.msl
index b984663..bb9b678 100644
--- a/test/tint/expressions/binary/div_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int tint_div_i32(int lhs, int rhs) {
+  return (lhs / select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1)))));
+}
 kernel void f() {
   int a = 1;
   int b = 0;
   int const r = tint_div_i32(a, b);
 }
-int tint_div_i32(int lhs, int rhs) {
-  return (lhs / select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1)))));
-}
-program_source:7:17: error: use of undeclared identifier 'tint_div_i32'
-  int const r = tint_div_i32(a, b);
-                ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.ir.msl
index 9ca11a6..f9ddb08 100644
--- a/test/tint/expressions/binary/div_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint tint_div_u32(uint lhs, uint rhs) {
+  return (lhs / select(rhs, 1u, (rhs == 0u)));
+}
 kernel void f() {
   uint a = 1u;
   uint b = 0u;
   uint const r = tint_div_u32(a, b);
 }
-uint tint_div_u32(uint lhs, uint rhs) {
-  return (lhs / select(rhs, 1u, (rhs == 0u)));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_div_u32'
-  uint const r = tint_div_u32(a, b);
-                 ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.ir.msl
index e2bbf6c..c50922d 100644
--- a/test/tint/expressions/binary/div_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_div_v3i32(int3 lhs, int3 rhs) {
+  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
+}
 kernel void f() {
   int a = 4;
   int3 b = int3(0, 2, 0);
   int3 const v = b;
   int3 const r = tint_div_v3i32(int3(a), v);
 }
-int3 tint_div_v3i32(int3 lhs, int3 rhs) {
-  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
-}
-program_source:8:18: error: use of undeclared identifier 'tint_div_v3i32'
-  int3 const r = tint_div_v3i32(int3(a), v);
-                 ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.ir.msl
index a3bfb50..e264dbb 100644
--- a/test/tint/expressions/binary/div_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
+  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
+}
 kernel void f() {
   uint a = 4u;
   uint3 b = uint3(0u, 2u, 0u);
   uint3 const v = b;
   uint3 const r = tint_div_v3u32(uint3(a), v);
 }
-uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
-  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
-}
-program_source:8:19: error: use of undeclared identifier 'tint_div_v3u32'
-  uint3 const r = tint_div_v3u32(uint3(a), v);
-                  ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.ir.msl
index 4e52ba9..369fa0d 100644
--- a/test/tint/expressions/binary/div_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_div_v3i32(int3 lhs, int3 rhs) {
+  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
+}
 kernel void f() {
   int3 a = int3(1, 2, 3);
   int b = 0;
   int3 const v = a;
   int3 const r = tint_div_v3i32(v, int3(b));
 }
-int3 tint_div_v3i32(int3 lhs, int3 rhs) {
-  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
-}
-program_source:8:18: error: use of undeclared identifier 'tint_div_v3i32'
-  int3 const r = tint_div_v3i32(v, int3(b));
-                 ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.ir.msl
index ea03857..717fc92 100644
--- a/test/tint/expressions/binary/div_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
+  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
+}
 kernel void f() {
   uint3 a = uint3(1u, 2u, 3u);
   uint b = 0u;
   uint3 const v = a;
   uint3 const r = tint_div_v3u32(v, uint3(b));
 }
-uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
-  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
-}
-program_source:8:19: error: use of undeclared identifier 'tint_div_v3u32'
-  uint3 const r = tint_div_v3u32(v, uint3(b));
-                  ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.ir.msl
index f604b15..54474a8 100644
--- a/test/tint/expressions/binary/div_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_div_v3i32(int3 lhs, int3 rhs) {
+  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
+}
 kernel void f() {
   int3 a = int3(1, 2, 3);
   int3 b = int3(0, 5, 0);
   int3 const r = tint_div_v3i32(a, b);
 }
-int3 tint_div_v3i32(int3 lhs, int3 rhs) {
-  return (lhs / select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1))))));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_div_v3i32'
-  int3 const r = tint_div_v3i32(a, b);
-                 ^
-
diff --git a/test/tint/expressions/binary/div_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/div_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.ir.msl
index 5b8af54..0d4d5f0 100644
--- a/test/tint/expressions/binary/div_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.ir.msl
@@ -1,17 +1,11 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
+  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
+}
 kernel void f() {
   uint3 a = uint3(1u, 2u, 3u);
   uint3 b = uint3(0u, 5u, 0u);
   uint3 const r = tint_div_v3u32(a, b);
 }
-uint3 tint_div_v3u32(uint3 lhs, uint3 rhs) {
-  return (lhs / select(rhs, uint3(1u), (rhs == uint3(0u))));
-}
-program_source:7:19: error: use of undeclared identifier 'tint_div_v3u32'
-  uint3 const r = tint_div_v3u32(a, b);
-                  ^
-
diff --git a/test/tint/expressions/binary/mod/scalar-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/scalar-scalar/i32.wgsl.expected.ir.msl
index 09b1ef2..e706a8b 100644
--- a/test/tint/expressions/binary/mod/scalar-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/scalar-scalar/i32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int tint_mod_i32(int lhs, int rhs) {
+  int const v = select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1))));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   int const a = 1;
   int const b = 2;
   int const r = tint_mod_i32(a, b);
 }
-int tint_mod_i32(int lhs, int rhs) {
-  int const v = select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1))));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:17: error: use of undeclared identifier 'tint_mod_i32'
-  int const r = tint_mod_i32(a, b);
-                ^
-
diff --git a/test/tint/expressions/binary/mod/scalar-scalar/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/scalar-scalar/u32.wgsl.expected.ir.msl
index b1b4a6f..4523947 100644
--- a/test/tint/expressions/binary/mod/scalar-scalar/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/scalar-scalar/u32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint tint_mod_u32(uint lhs, uint rhs) {
+  uint const v = select(rhs, 1u, (rhs == 0u));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   uint const a = 1u;
   uint const b = 2u;
   uint const r = tint_mod_u32(a, b);
 }
-uint tint_mod_u32(uint lhs, uint rhs) {
-  uint const v = select(rhs, 1u, (rhs == 0u));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_mod_u32'
-  uint const r = tint_mod_u32(a, b);
-                 ^
-
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.ir.msl
index d0ad189..9ce7e38 100644
--- a/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/scalar-vec3/i32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
+  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   int const a = 4;
   int3 const b = int3(1, 2, 3);
   int3 const r = tint_mod_v3i32(int3(a), b);
 }
-int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
-  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_mod_v3i32'
-  int3 const r = tint_mod_v3i32(int3(a), b);
-                 ^
-
diff --git a/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.ir.msl
index c0e4d0b..47bbf31 100644
--- a/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/scalar-vec3/u32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
+  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   uint const a = 4u;
   uint3 const b = uint3(1u, 2u, 3u);
   uint3 const r = tint_mod_v3u32(uint3(a), b);
 }
-uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
-  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:19: error: use of undeclared identifier 'tint_mod_v3u32'
-  uint3 const r = tint_mod_v3u32(uint3(a), b);
-                  ^
-
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.ir.msl
index 8ac3f73..86ab832 100644
--- a/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/vec3-scalar/i32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
+  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   int3 const a = int3(1, 2, 3);
   int const b = 4;
   int3 const r = tint_mod_v3i32(a, int3(b));
 }
-int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
-  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_mod_v3i32'
-  int3 const r = tint_mod_v3i32(a, int3(b));
-                 ^
-
diff --git a/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.ir.msl
index f815457..7132fef 100644
--- a/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/vec3-scalar/u32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
+  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   uint3 const a = uint3(1u, 2u, 3u);
   uint const b = 4u;
   uint3 const r = tint_mod_v3u32(a, uint3(b));
 }
-uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
-  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:19: error: use of undeclared identifier 'tint_mod_v3u32'
-  uint3 const r = tint_mod_v3u32(a, uint3(b));
-                  ^
-
diff --git a/test/tint/expressions/binary/mod/vec3-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/vec3-vec3/i32.wgsl.expected.ir.msl
index 51bb608..1bd8a4d 100644
--- a/test/tint/expressions/binary/mod/vec3-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/vec3-vec3/i32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
+  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   int3 const a = int3(1, 2, 3);
   int3 const b = int3(4, 5, 6);
   int3 const r = tint_mod_v3i32(a, b);
 }
-int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
-  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_mod_v3i32'
-  int3 const r = tint_mod_v3i32(a, b);
-                 ^
-
diff --git a/test/tint/expressions/binary/mod/vec3-vec3/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod/vec3-vec3/u32.wgsl.expected.ir.msl
index 59cfca7..0e9071e 100644
--- a/test/tint/expressions/binary/mod/vec3-vec3/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod/vec3-vec3/u32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
+  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   uint3 const a = uint3(1u, 2u, 3u);
   uint3 const b = uint3(4u, 5u, 6u);
   uint3 const r = tint_mod_v3u32(a, b);
 }
-uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
-  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:19: error: use of undeclared identifier 'tint_mod_v3u32'
-  uint3 const r = tint_mod_v3u32(a, b);
-                  ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.ir.msl
index 39862e7..33cc550 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/i32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int tint_mod_i32(int lhs, int rhs) {
+  int const v = select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1))));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   int const a = 1;
   int const b = 0;
   int const r = tint_mod_i32(a, b);
 }
-int tint_mod_i32(int lhs, int rhs) {
-  int const v = select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1))));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:17: error: use of undeclared identifier 'tint_mod_i32'
-  int const r = tint_mod_i32(a, b);
-                ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.ir.msl
index d3132b1..d7c8be1 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-scalar/u32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint tint_mod_u32(uint lhs, uint rhs) {
+  uint const v = select(rhs, 1u, (rhs == 0u));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   uint const a = 1u;
   uint const b = 0u;
   uint const r = tint_mod_u32(a, b);
 }
-uint tint_mod_u32(uint lhs, uint rhs) {
-  uint const v = select(rhs, 1u, (rhs == 0u));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_mod_u32'
-  uint const r = tint_mod_u32(a, b);
-                 ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.ir.msl
index 42118c6..b547642 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-vec3/i32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
+  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   int const a = 4;
   int3 const b = int3(0, 2, 0);
   int3 const r = tint_mod_v3i32(int3(a), b);
 }
-int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
-  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_mod_v3i32'
-  int3 const r = tint_mod_v3i32(int3(a), b);
-                 ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.ir.msl
index 175880b..5407b2c 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_constant/scalar-vec3/u32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
+  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   uint const a = 4u;
   uint3 const b = uint3(0u, 2u, 0u);
   uint3 const r = tint_mod_v3u32(uint3(a), b);
 }
-uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
-  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:19: error: use of undeclared identifier 'tint_mod_v3u32'
-  uint3 const r = tint_mod_v3u32(uint3(a), b);
-                  ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.ir.msl
index 34063d1..a9e877b 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-scalar/i32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
+  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   int3 const a = int3(1, 2, 3);
   int const b = 0;
   int3 const r = tint_mod_v3i32(a, int3(b));
 }
-int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
-  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_mod_v3i32'
-  int3 const r = tint_mod_v3i32(a, int3(b));
-                 ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.ir.msl
index ae9e090..8e7c960 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-scalar/u32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
+  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   uint3 const a = uint3(1u, 2u, 3u);
   uint const b = 0u;
   uint3 const r = tint_mod_v3u32(a, uint3(b));
 }
-uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
-  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:19: error: use of undeclared identifier 'tint_mod_v3u32'
-  uint3 const r = tint_mod_v3u32(a, uint3(b));
-                  ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.ir.msl
index f7be78a..520e56f 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/i32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
+  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   int3 const a = int3(1, 2, 3);
   int3 const b = int3(0, 5, 0);
   int3 const r = tint_mod_v3i32(a, b);
 }
-int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
-  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_mod_v3i32'
-  int3 const r = tint_mod_v3i32(a, b);
-                 ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.ir.msl
index 8e25d40..a60f38d 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_constant/vec3-vec3/u32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
+  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   uint3 const a = uint3(1u, 2u, 3u);
   uint3 const b = uint3(0u, 5u, 0u);
   uint3 const r = tint_mod_v3u32(a, b);
 }
-uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
-  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:19: error: use of undeclared identifier 'tint_mod_v3u32'
-  uint3 const r = tint_mod_v3u32(a, b);
-                  ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.ir.msl
index 42df4ee..843f145 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/i32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int tint_mod_i32(int lhs, int rhs) {
+  int const v = select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1))));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   int a = 1;
   int b = 0;
   int const r = tint_mod_i32(a, (b + b));
 }
-int tint_mod_i32(int lhs, int rhs) {
-  int const v = select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1))));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:17: error: use of undeclared identifier 'tint_mod_i32'
-  int const r = tint_mod_i32(a, (b + b));
-                ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.ir.msl
index 5f6496a..a4a1ad9 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-scalar/u32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint tint_mod_u32(uint lhs, uint rhs) {
+  uint const v = select(rhs, 1u, (rhs == 0u));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   uint a = 1u;
   uint b = 0u;
   uint const r = tint_mod_u32(a, (b + b));
 }
-uint tint_mod_u32(uint lhs, uint rhs) {
-  uint const v = select(rhs, 1u, (rhs == 0u));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_mod_u32'
-  uint const r = tint_mod_u32(a, (b + b));
-                 ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.ir.msl
index 4395b07..67af20b 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.ir.msl
@@ -1,19 +1,13 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
+  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   int a = 4;
   int3 b = int3(0, 2, 0);
-  int3 const v = (b + b);
-  int3 const r = tint_mod_v3i32(int3(a), v);
+  int3 const v_1 = (b + b);
+  int3 const r = tint_mod_v3i32(int3(a), v_1);
 }
-int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
-  int3 const v_1 = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
-  return (lhs - ((lhs / v_1) * v_1));
-}
-program_source:8:18: error: use of undeclared identifier 'tint_mod_v3i32'
-  int3 const r = tint_mod_v3i32(int3(a), v);
-                 ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.ir.msl
index 23fa120..4dfb69b 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/u32.wgsl.expected.ir.msl
@@ -1,19 +1,13 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
+  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   uint a = 4u;
   uint3 b = uint3(0u, 2u, 0u);
-  uint3 const v = (b + b);
-  uint3 const r = tint_mod_v3u32(uint3(a), v);
+  uint3 const v_1 = (b + b);
+  uint3 const r = tint_mod_v3u32(uint3(a), v_1);
 }
-uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
-  uint3 const v_1 = select(rhs, uint3(1u), (rhs == uint3(0u)));
-  return (lhs - ((lhs / v_1) * v_1));
-}
-program_source:8:19: error: use of undeclared identifier 'tint_mod_v3u32'
-  uint3 const r = tint_mod_v3u32(uint3(a), v);
-                  ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.ir.msl
index b8ab3e9..fdd4a00 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-scalar/i32.wgsl.expected.ir.msl
@@ -1,19 +1,13 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
+  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   int3 a = int3(1, 2, 3);
   int b = 0;
-  int3 const v = a;
-  int3 const r = tint_mod_v3i32(v, int3((b + b)));
+  int3 const v_1 = a;
+  int3 const r = tint_mod_v3i32(v_1, int3((b + b)));
 }
-int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
-  int3 const v_1 = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
-  return (lhs - ((lhs / v_1) * v_1));
-}
-program_source:8:18: error: use of undeclared identifier 'tint_mod_v3i32'
-  int3 const r = tint_mod_v3i32(v, int3((b + b)));
-                 ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.ir.msl
index c63c2cb..ad990c1 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-scalar/u32.wgsl.expected.ir.msl
@@ -1,19 +1,13 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
+  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   uint3 a = uint3(1u, 2u, 3u);
   uint b = 0u;
-  uint3 const v = a;
-  uint3 const r = tint_mod_v3u32(v, uint3((b + b)));
+  uint3 const v_1 = a;
+  uint3 const r = tint_mod_v3u32(v_1, uint3((b + b)));
 }
-uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
-  uint3 const v_1 = select(rhs, uint3(1u), (rhs == uint3(0u)));
-  return (lhs - ((lhs / v_1) * v_1));
-}
-program_source:8:19: error: use of undeclared identifier 'tint_mod_v3u32'
-  uint3 const r = tint_mod_v3u32(v, uint3((b + b)));
-                  ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.ir.msl
index a9eca17..980b4a1 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
+  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   int3 a = int3(1, 2, 3);
   int3 b = int3(0, 5, 0);
   int3 const r = tint_mod_v3i32(a, (b + b));
 }
-int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
-  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_mod_v3i32'
-  int3 const r = tint_mod_v3i32(a, (b + b));
-                 ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.ir.msl
index 8b388fa..3d18b4a 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/u32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
+  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   uint3 a = uint3(1u, 2u, 3u);
   uint3 b = uint3(0u, 5u, 0u);
   uint3 const r = tint_mod_v3u32(a, (b + b));
 }
-uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
-  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:19: error: use of undeclared identifier 'tint_mod_v3u32'
-  uint3 const r = tint_mod_v3u32(a, (b + b));
-                  ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.ir.msl
index e6fd9e8..56b4181 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/i32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int tint_mod_i32(int lhs, int rhs) {
+  int const v = select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1))));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   int a = 1;
   int b = 0;
   int const r = tint_mod_i32(a, b);
 }
-int tint_mod_i32(int lhs, int rhs) {
-  int const v = select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1))));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:17: error: use of undeclared identifier 'tint_mod_i32'
-  int const r = tint_mod_i32(a, b);
-                ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.ir.msl
index bbd1107..2acb8b6 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-scalar/u32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint tint_mod_u32(uint lhs, uint rhs) {
+  uint const v = select(rhs, 1u, (rhs == 0u));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   uint a = 1u;
   uint b = 0u;
   uint const r = tint_mod_u32(a, b);
 }
-uint tint_mod_u32(uint lhs, uint rhs) {
-  uint const v = select(rhs, 1u, (rhs == 0u));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_mod_u32'
-  uint const r = tint_mod_u32(a, b);
-                 ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.ir.msl
index 2231a73..50eb686 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/i32.wgsl.expected.ir.msl
@@ -1,19 +1,13 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
+  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   int a = 4;
   int3 b = int3(0, 2, 0);
-  int3 const v = b;
-  int3 const r = tint_mod_v3i32(int3(a), v);
+  int3 const v_1 = b;
+  int3 const r = tint_mod_v3i32(int3(a), v_1);
 }
-int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
-  int3 const v_1 = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
-  return (lhs - ((lhs / v_1) * v_1));
-}
-program_source:8:18: error: use of undeclared identifier 'tint_mod_v3i32'
-  int3 const r = tint_mod_v3i32(int3(a), v);
-                 ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.ir.msl
index 5b1ca58..f2fd2dc 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_identifier/scalar-vec3/u32.wgsl.expected.ir.msl
@@ -1,19 +1,13 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
+  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   uint a = 4u;
   uint3 b = uint3(0u, 2u, 0u);
-  uint3 const v = b;
-  uint3 const r = tint_mod_v3u32(uint3(a), v);
+  uint3 const v_1 = b;
+  uint3 const r = tint_mod_v3u32(uint3(a), v_1);
 }
-uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
-  uint3 const v_1 = select(rhs, uint3(1u), (rhs == uint3(0u)));
-  return (lhs - ((lhs / v_1) * v_1));
-}
-program_source:8:19: error: use of undeclared identifier 'tint_mod_v3u32'
-  uint3 const r = tint_mod_v3u32(uint3(a), v);
-                  ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.ir.msl
index a6017c7..20cef62 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/i32.wgsl.expected.ir.msl
@@ -1,19 +1,13 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
+  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   int3 a = int3(1, 2, 3);
   int b = 0;
-  int3 const v = a;
-  int3 const r = tint_mod_v3i32(v, int3(b));
+  int3 const v_1 = a;
+  int3 const r = tint_mod_v3i32(v_1, int3(b));
 }
-int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
-  int3 const v_1 = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
-  return (lhs - ((lhs / v_1) * v_1));
-}
-program_source:8:18: error: use of undeclared identifier 'tint_mod_v3i32'
-  int3 const r = tint_mod_v3i32(v, int3(b));
-                 ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.ir.msl
index 96355d9..d7533fa 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-scalar/u32.wgsl.expected.ir.msl
@@ -1,19 +1,13 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
+  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   uint3 a = uint3(1u, 2u, 3u);
   uint b = 0u;
-  uint3 const v = a;
-  uint3 const r = tint_mod_v3u32(v, uint3(b));
+  uint3 const v_1 = a;
+  uint3 const r = tint_mod_v3u32(v_1, uint3(b));
 }
-uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
-  uint3 const v_1 = select(rhs, uint3(1u), (rhs == uint3(0u)));
-  return (lhs - ((lhs / v_1) * v_1));
-}
-program_source:8:19: error: use of undeclared identifier 'tint_mod_v3u32'
-  uint3 const r = tint_mod_v3u32(v, uint3(b));
-                  ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.ir.msl
index a7c66f5..c204531 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/i32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
+  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   int3 a = int3(1, 2, 3);
   int3 b = int3(0, 5, 0);
   int3 const r = tint_mod_v3i32(a, b);
 }
-int3 tint_mod_v3i32(int3 lhs, int3 rhs) {
-  int3 const v = select(rhs, int3(1), ((rhs == int3(0)) | ((lhs == int3((-2147483647 - 1))) & (rhs == int3(-1)))));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:18: error: use of undeclared identifier 'tint_mod_v3i32'
-  int3 const r = tint_mod_v3i32(a, b);
-                 ^
-
diff --git a/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.ir.msl
index 60ddeeb..e040c27 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_identifier/vec3-vec3/u32.wgsl.expected.ir.msl
@@ -1,18 +1,12 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
+  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
+  return (lhs - ((lhs / v) * v));
+}
 kernel void f() {
   uint3 a = uint3(1u, 2u, 3u);
   uint3 b = uint3(0u, 5u, 0u);
   uint3 const r = tint_mod_v3u32(a, b);
 }
-uint3 tint_mod_v3u32(uint3 lhs, uint3 rhs) {
-  uint3 const v = select(rhs, uint3(1u), (rhs == uint3(0u)));
-  return (lhs - ((lhs / v) * v));
-}
-program_source:7:19: error: use of undeclared identifier 'tint_mod_v3u32'
-  uint3 const r = tint_mod_v3u32(a, b);
-                  ^
-
diff --git a/test/tint/statements/compound_assign/function.wgsl.expected.ir.msl b/test/tint/statements/compound_assign/function.wgsl.expected.ir.msl
index e7169f7..400e900 100644
--- a/test/tint/statements/compound_assign/function.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/function.wgsl.expected.ir.msl
@@ -1,8 +1,9 @@
-SKIP: FAILED
-
 #include <metal_stdlib>
 using namespace metal;
 
+int tint_div_i32(int lhs, int rhs) {
+  return (lhs / select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1)))));
+}
 void foo() {
   int a = 0;
   float4 b = 0.0f;
@@ -11,10 +12,3 @@
   b = (b * float4x4(float4(0.0f), float4(0.0f), float4(0.0f), float4(0.0f)));
   c = (c * 2.0f);
 }
-int tint_div_i32(int lhs, int rhs) {
-  return (lhs / select(rhs, 1, ((rhs == 0) | ((lhs == (-2147483647 - 1)) & (rhs == -1)))));
-}
-program_source:8:7: error: use of undeclared identifier 'tint_div_i32'
-  a = tint_div_i32(a, 2);
-      ^
-