diff --git a/naga/src/back/msl/mod.rs b/naga/src/back/msl/mod.rs
index 28e99acc5f..2694392dde 100644
--- a/naga/src/back/msl/mod.rs
+++ b/naga/src/back/msl/mod.rs
@@ -148,6 +148,8 @@ pub enum Error {
     UnsupportedRayTracing,
     #[error("overrides should not be present at this stage")]
     Override,
+    #[error("bitcasting to {0:?} is not supported")]
+    UnsupportedBitCast(crate::TypeInner),
 }
 
 #[derive(Clone, Debug, PartialEq, thiserror::Error)]
diff --git a/naga/src/back/msl/writer.rs b/naga/src/back/msl/writer.rs
index 4589d39892..cdfd62129b 100644
--- a/naga/src/back/msl/writer.rs
+++ b/naga/src/back/msl/writer.rs
@@ -1357,7 +1357,7 @@ impl<W: Write> Writer<W> {
         Ok(())
     }
 
-    /// Emit code for the sign(i32) expression.
+    /// Emit code for the isign expression.
     ///
     fn put_isign(
         &mut self,
@@ -1365,18 +1365,23 @@ impl<W: Write> Writer<W> {
         context: &ExpressionContext,
     ) -> BackendResult {
         write!(self.out, "{NAMESPACE}::select({NAMESPACE}::select(")?;
+        let scalar = context
+            .resolve_type(arg)
+            .scalar()
+            .expect("put_isign should only be called for args which have an integer scalar type")
+            .to_msl_name();
         match context.resolve_type(arg) {
             &crate::TypeInner::Vector { size, .. } => {
                 let size = back::vector_size_str(size);
-                write!(self.out, "int{size}(-1), int{size}(1)")?;
+                write!(self.out, "{scalar}{size}(-1), {scalar}{size}(1)")?;
             }
             _ => {
-                write!(self.out, "-1, 1")?;
+                write!(self.out, "{scalar}(-1), {scalar}(1)")?;
             }
         }
         write!(self.out, ", (")?;
         self.put_expression(arg, context, true)?;
-        write!(self.out, " > 0)), 0, (")?;
+        write!(self.out, " > 0)), {scalar}(0), (")?;
         self.put_expression(arg, context, true)?;
         write!(self.out, " == 0))")?;
         Ok(())
@@ -1605,7 +1610,12 @@ impl<W: Write> Writer<W> {
                 vector,
                 pattern,
             } => {
-                self.put_wrapped_expression_for_packed_vec3_access(vector, context, false)?;
+                self.put_wrapped_expression_for_packed_vec3_access(
+                    vector,
+                    context,
+                    false,
+                    &Self::put_expression,
+                )?;
                 write!(self.out, ".")?;
                 for &sc in pattern[..size as usize].iter() {
                     write!(self.out, "{}", back::COMPONENTS[sc as usize])?;
@@ -1748,7 +1758,6 @@ impl<W: Write> Writer<W> {
                 write!(self.out, ")")?;
             }
             crate::Expression::Binary { op, left, right } => {
-                let op_str = back::binary_operation_str(op);
                 let kind = context
                     .resolve_type(left)
                     .scalar_kind()
@@ -1773,38 +1782,56 @@ impl<W: Write> Writer<W> {
                     write!(self.out, ", ")?;
                     self.put_expression(right, context, true)?;
                     write!(self.out, ")")?;
+                } else if (op == crate::BinaryOperator::Add
+                    || op == crate::BinaryOperator::Subtract
+                    || op == crate::BinaryOperator::Multiply)
+                    && kind == crate::ScalarKind::Sint
+                {
+                    let to_unsigned = |ty: &crate::TypeInner| match *ty {
+                        crate::TypeInner::Scalar(scalar) => {
+                            Ok(crate::TypeInner::Scalar(crate::Scalar {
+                                kind: crate::ScalarKind::Uint,
+                                ..scalar
+                            }))
+                        }
+                        crate::TypeInner::Vector { size, scalar } => Ok(crate::TypeInner::Vector {
+                            size,
+                            scalar: crate::Scalar {
+                                kind: crate::ScalarKind::Uint,
+                                ..scalar
+                            },
+                        }),
+                        _ => Err(Error::UnsupportedBitCast(ty.clone())),
+                    };
+
+                    // Avoid undefined behaviour due to overflowing signed
+                    // integer arithmetic. Cast the operands to unsigned prior
+                    // to performing the operation, then cast the result back
+                    // to signed.
+                    self.put_bitcasted_expression(
+                        context.resolve_type(expr_handle),
+                        context,
+                        &|writer, context, is_scoped| {
+                            writer.put_binop(
+                                op,
+                                left,
+                                right,
+                                context,
+                                is_scoped,
+                                &|writer, expr, context, _is_scoped| {
+                                    writer.put_bitcasted_expression(
+                                        &to_unsigned(context.resolve_type(expr))?,
+                                        context,
+                                        &|writer, context, is_scoped| {
+                                            writer.put_expression(expr, context, is_scoped)
+                                        },
+                                    )
+                                },
+                            )
+                        },
+                    )?;
                 } else {
-                    if !is_scoped {
-                        write!(self.out, "(")?;
-                    }
-
-                    // Cast packed vector if necessary
-                    // Packed vector - matrix multiplications are not supported in MSL
-                    if op == crate::BinaryOperator::Multiply
-                        && matches!(
-                            context.resolve_type(right),
-                            &crate::TypeInner::Matrix { .. }
-                        )
-                    {
-                        self.put_wrapped_expression_for_packed_vec3_access(left, context, false)?;
-                    } else {
-                        self.put_expression(left, context, false)?;
-                    }
-
-                    write!(self.out, " {op_str} ")?;
-
-                    // See comment above
-                    if op == crate::BinaryOperator::Multiply
-                        && matches!(context.resolve_type(left), &crate::TypeInner::Matrix { .. })
-                    {
-                        self.put_wrapped_expression_for_packed_vec3_access(right, context, false)?;
-                    } else {
-                        self.put_expression(right, context, false)?;
-                    }
-
-                    if !is_scoped {
-                        write!(self.out, ")")?;
-                    }
+                    self.put_binop(op, left, right, context, is_scoped, &Self::put_expression)?;
                 }
             }
             crate::Expression::Select {
@@ -2325,23 +2352,114 @@ impl<W: Write> Writer<W> {
         Ok(())
     }
 
+    /// Emits code for a binary operation, using the provided callback to emit
+    /// the left and right operands.
+    fn put_binop<F>(
+        &mut self,
+        op: crate::BinaryOperator,
+        left: Handle<crate::Expression>,
+        right: Handle<crate::Expression>,
+        context: &ExpressionContext,
+        is_scoped: bool,
+        put_expression: &F,
+    ) -> BackendResult
+    where
+        F: Fn(&mut Self, Handle<crate::Expression>, &ExpressionContext, bool) -> BackendResult,
+    {
+        let op_str = back::binary_operation_str(op);
+
+        if !is_scoped {
+            write!(self.out, "(")?;
+        }
+
+        // Cast packed vector if necessary
+        // Packed vector - matrix multiplications are not supported in MSL
+        if op == crate::BinaryOperator::Multiply
+            && matches!(
+                context.resolve_type(right),
+                &crate::TypeInner::Matrix { .. }
+            )
+        {
+            self.put_wrapped_expression_for_packed_vec3_access(
+                left,
+                context,
+                false,
+                put_expression,
+            )?;
+        } else {
+            put_expression(self, left, context, false)?;
+        }
+
+        write!(self.out, " {op_str} ")?;
+
+        // See comment above
+        if op == crate::BinaryOperator::Multiply
+            && matches!(context.resolve_type(left), &crate::TypeInner::Matrix { .. })
+        {
+            self.put_wrapped_expression_for_packed_vec3_access(
+                right,
+                context,
+                false,
+                put_expression,
+            )?;
+        } else {
+            put_expression(self, right, context, false)?;
+        }
+
+        if !is_scoped {
+            write!(self.out, ")")?;
+        }
+
+        Ok(())
+    }
+
     /// Used by expressions like Swizzle and Binary since they need packed_vec3's to be casted to a vec3
-    fn put_wrapped_expression_for_packed_vec3_access(
+    fn put_wrapped_expression_for_packed_vec3_access<F>(
         &mut self,
         expr_handle: Handle<crate::Expression>,
         context: &ExpressionContext,
         is_scoped: bool,
-    ) -> BackendResult {
+        put_expression: &F,
+    ) -> BackendResult
+    where
+        F: Fn(&mut Self, Handle<crate::Expression>, &ExpressionContext, bool) -> BackendResult,
+    {
         if let Some(scalar) = context.get_packed_vec_kind(expr_handle) {
             write!(self.out, "{}::{}3(", NAMESPACE, scalar.to_msl_name())?;
-            self.put_expression(expr_handle, context, is_scoped)?;
+            put_expression(self, expr_handle, context, is_scoped)?;
             write!(self.out, ")")?;
         } else {
-            self.put_expression(expr_handle, context, is_scoped)?;
+            put_expression(self, expr_handle, context, is_scoped)?;
         }
         Ok(())
     }
 
+    /// Emits code for an expression using the provided callback, wrapping the
+    /// result in a bitcast to the type `cast_to`.
+    fn put_bitcasted_expression<F>(
+        &mut self,
+        cast_to: &crate::TypeInner,
+        context: &ExpressionContext,
+        put_expression: &F,
+    ) -> BackendResult
+    where
+        F: Fn(&mut Self, &ExpressionContext, bool) -> BackendResult,
+    {
+        write!(self.out, "as_type<")?;
+        match *cast_to {
+            crate::TypeInner::Scalar(scalar) => put_numeric_type(&mut self.out, scalar, &[])?,
+            crate::TypeInner::Vector { size, scalar } => {
+                put_numeric_type(&mut self.out, scalar, &[size])?
+            }
+            _ => return Err(Error::UnsupportedBitCast(cast_to.clone())),
+        };
+        write!(self.out, ">(")?;
+        put_expression(self, context, true)?;
+        write!(self.out, ")")?;
+
+        Ok(())
+    }
+
     /// Write a `GuardedIndex` as a Metal expression.
     fn put_index(
         &mut self,
diff --git a/naga/tests/out/msl/abstract-types-operators.msl b/naga/tests/out/msl/abstract-types-operators.msl
index cec86a53dc..54fd234d29 100644
--- a/naga/tests/out/msl/abstract-types-operators.msl
+++ b/naga/tests/out/msl/abstract-types-operators.msl
@@ -73,12 +73,12 @@ void runtime_values(
     float _e27 = f;
     plus_f_f_f = _e26 + _e27;
     int _e31 = i;
-    plus_iai_i = 1 + _e31;
+    plus_iai_i = as_type<int>(as_type<uint>(1) + as_type<uint>(_e31));
     int _e35 = i;
-    plus_i_iai = _e35 + 2;
+    plus_i_iai = as_type<int>(as_type<uint>(_e35) + as_type<uint>(2));
     int _e39 = i;
     int _e40 = i;
-    plus_i_i_i = _e39 + _e40;
+    plus_i_i_i = as_type<int>(as_type<uint>(_e39) + as_type<uint>(_e40));
     uint _e44 = u;
     plus_uai_u = 1u + _e44;
     uint _e48 = u;
@@ -97,5 +97,5 @@ void wgpu_4445_(
 void wgpu_4435_(
     threadgroup type_3& a
 ) {
-    uint y = a.inner[1 - 1];
+    uint y = a.inner[as_type<int>(as_type<uint>(1) - as_type<uint>(1))];
 }
diff --git a/naga/tests/out/msl/access.msl b/naga/tests/out/msl/access.msl
index 2103a8d099..041923ae29 100644
--- a/naga/tests/out/msl/access.msl
+++ b/naga/tests/out/msl/access.msl
@@ -70,7 +70,7 @@ void test_matrix_within_struct_accesses(
     int idx = 1;
     Baz t = Baz {metal::float3x2(metal::float2(1.0), metal::float2(2.0), metal::float2(3.0))};
     int _e3 = idx;
-    idx = _e3 - 1;
+    idx = as_type<int>(as_type<uint>(_e3) - as_type<uint>(1));
     metal::float3x2 l0_ = baz.m;
     metal::float2 l1_ = baz.m[0];
     int _e14 = idx;
@@ -84,7 +84,7 @@ void test_matrix_within_struct_accesses(
     int _e38 = idx;
     float l6_ = baz.m[_e36][_e38];
     int _e51 = idx;
-    idx = _e51 + 1;
+    idx = as_type<int>(as_type<uint>(_e51) + as_type<uint>(1));
     t.m = metal::float3x2(metal::float2(6.0), metal::float2(5.0), metal::float2(4.0));
     t.m[0] = metal::float2(9.0);
     int _e66 = idx;
@@ -106,7 +106,7 @@ void test_matrix_within_array_within_struct_accesses(
     int idx_1 = 1;
     MatCx2InArray t_1 = MatCx2InArray {type_15 {}};
     int _e3 = idx_1;
-    idx_1 = _e3 - 1;
+    idx_1 = as_type<int>(as_type<uint>(_e3) - as_type<uint>(1));
     type_15 l0_1 = nested_mat_cx2_.am;
     metal::float4x2 l1_1 = nested_mat_cx2_.am.inner[0];
     metal::float2 l2_1 = nested_mat_cx2_.am.inner[0][0];
@@ -121,7 +121,7 @@ void test_matrix_within_array_within_struct_accesses(
     int _e48 = idx_1;
     float l7_ = nested_mat_cx2_.am.inner[0][_e46][_e48];
     int _e55 = idx_1;
-    idx_1 = _e55 + 1;
+    idx_1 = as_type<int>(as_type<uint>(_e55) + as_type<uint>(1));
     t_1.am = type_15 {};
     t_1.am.inner[0] = metal::float4x2(metal::float2(8.0), metal::float2(7.0), metal::float2(6.0), metal::float2(5.0));
     t_1.am.inner[0][0] = metal::float2(9.0);
diff --git a/naga/tests/out/msl/boids.msl b/naga/tests/out/msl/boids.msl
index 07acd7cf62..a4d709bd09 100644
--- a/naga/tests/out/msl/boids.msl
+++ b/naga/tests/out/msl/boids.msl
@@ -84,7 +84,7 @@ kernel void main_(
             metal::float2 _e61 = pos;
             cMass = _e60 + _e61;
             int _e63 = cMassCount;
-            cMassCount = _e63 + 1;
+            cMassCount = as_type<int>(as_type<uint>(_e63) + as_type<uint>(1));
         }
         metal::float2 _e66 = pos;
         metal::float2 _e67 = vPos;
@@ -103,7 +103,7 @@ kernel void main_(
             metal::float2 _e86 = vel;
             cVel = _e85 + _e86;
             int _e88 = cVelCount;
-            cVelCount = _e88 + 1;
+            cVelCount = as_type<int>(as_type<uint>(_e88) + as_type<uint>(1));
         }
 #define LOOP_IS_BOUNDED { volatile bool unpredictable_break_from_loop = false; if (unpredictable_break_from_loop) break; }
         LOOP_IS_BOUNDED
diff --git a/naga/tests/out/msl/empty-global-name.msl b/naga/tests/out/msl/empty-global-name.msl
index 01cac3f6e0..5db7fdd932 100644
--- a/naga/tests/out/msl/empty-global-name.msl
+++ b/naga/tests/out/msl/empty-global-name.msl
@@ -12,7 +12,7 @@ void function(
     device type_1& unnamed
 ) {
     int _e3 = unnamed.member;
-    unnamed.member = _e3 + 1;
+    unnamed.member = as_type<int>(as_type<uint>(_e3) + as_type<uint>(1));
     return;
 }
 
diff --git a/naga/tests/out/msl/image.msl b/naga/tests/out/msl/image.msl
index 114ed36553..1a8567a0ca 100644
--- a/naga/tests/out/msl/image.msl
+++ b/naga/tests/out/msl/image.msl
@@ -21,14 +21,14 @@ kernel void main_(
     metal::uint4 value1_ = image_mipmapped_src.read(metal::uint2(itc), static_cast<int>(local_id.z));
     metal::uint4 value2_ = image_multisampled_src.read(metal::uint2(itc), static_cast<int>(local_id.z));
     metal::uint4 value4_ = image_storage_src.read(metal::uint2(itc));
-    metal::uint4 value5_ = image_array_src.read(metal::uint2(itc), local_id.z, static_cast<int>(local_id.z) + 1);
-    metal::uint4 value6_ = image_array_src.read(metal::uint2(itc), static_cast<int>(local_id.z), static_cast<int>(local_id.z) + 1);
+    metal::uint4 value5_ = image_array_src.read(metal::uint2(itc), local_id.z, as_type<int>(as_type<uint>(static_cast<int>(local_id.z)) + as_type<uint>(1)));
+    metal::uint4 value6_ = image_array_src.read(metal::uint2(itc), static_cast<int>(local_id.z), as_type<int>(as_type<uint>(static_cast<int>(local_id.z)) + as_type<uint>(1)));
     metal::uint4 value7_ = image_1d_src.read(uint(static_cast<int>(local_id.x)));
     metal::uint4 value1u = image_mipmapped_src.read(metal::uint2(static_cast<metal::uint2>(itc)), static_cast<int>(local_id.z));
     metal::uint4 value2u = image_multisampled_src.read(metal::uint2(static_cast<metal::uint2>(itc)), static_cast<int>(local_id.z));
     metal::uint4 value4u = image_storage_src.read(metal::uint2(static_cast<metal::uint2>(itc)));
-    metal::uint4 value5u = image_array_src.read(metal::uint2(static_cast<metal::uint2>(itc)), local_id.z, static_cast<int>(local_id.z) + 1);
-    metal::uint4 value6u = image_array_src.read(metal::uint2(static_cast<metal::uint2>(itc)), static_cast<int>(local_id.z), static_cast<int>(local_id.z) + 1);
+    metal::uint4 value5u = image_array_src.read(metal::uint2(static_cast<metal::uint2>(itc)), local_id.z, as_type<int>(as_type<uint>(static_cast<int>(local_id.z)) + as_type<uint>(1)));
+    metal::uint4 value6u = image_array_src.read(metal::uint2(static_cast<metal::uint2>(itc)), static_cast<int>(local_id.z), as_type<int>(as_type<uint>(static_cast<int>(local_id.z)) + as_type<uint>(1)));
     metal::uint4 value7u = image_1d_src.read(uint(static_cast<uint>(local_id.x)));
     image_dst.write((((value1_ + value2_) + value4_) + value5_) + value6_, uint(itc.x));
     image_dst.write((((value1u + value2u) + value4u) + value5u) + value6u, uint(static_cast<uint>(itc.x)));
diff --git a/naga/tests/out/msl/int64.msl b/naga/tests/out/msl/int64.msl
index 79a3c06889..369eb01c6f 100644
--- a/naga/tests/out/msl/int64.msl
+++ b/naga/tests/out/msl/int64.msl
@@ -44,76 +44,76 @@ long int64_function(
 ) {
     long val = 20L;
     long _e8 = val;
-    val = _e8 + ((31L - 1002003004005006L) + -9223372036854775807L);
+    val = as_type<long>(as_type<ulong>(_e8) + as_type<ulong>(as_type<long>(as_type<ulong>(as_type<long>(as_type<ulong>(31L) - as_type<ulong>(1002003004005006L))) + as_type<ulong>(-9223372036854775807L))));
     long _e10 = val;
     long _e13 = val;
-    val = _e13 + (_e10 + 5L);
+    val = as_type<long>(as_type<ulong>(_e13) + as_type<ulong>(as_type<long>(as_type<ulong>(_e10) + as_type<ulong>(5L))));
     uint _e17 = input_uniform.val_u32_;
     long _e18 = val;
     long _e22 = val;
-    val = _e22 + static_cast<long>(_e17 + static_cast<uint>(_e18));
+    val = as_type<long>(as_type<ulong>(_e22) + as_type<ulong>(static_cast<long>(_e17 + static_cast<uint>(_e18))));
     int _e26 = input_uniform.val_i32_;
     long _e27 = val;
     long _e31 = val;
-    val = _e31 + static_cast<long>(_e26 + static_cast<int>(_e27));
+    val = as_type<long>(as_type<ulong>(_e31) + as_type<ulong>(static_cast<long>(as_type<int>(as_type<uint>(_e26) + as_type<uint>(static_cast<int>(_e27))))));
     float _e35 = input_uniform.val_f32_;
     long _e36 = val;
     long _e40 = val;
-    val = _e40 + static_cast<long>(_e35 + static_cast<float>(_e36));
+    val = as_type<long>(as_type<ulong>(_e40) + as_type<ulong>(static_cast<long>(_e35 + static_cast<float>(_e36))));
     long _e44 = input_uniform.val_i64_;
     long _e47 = val;
-    val = _e47 + metal::long3(_e44).z;
+    val = as_type<long>(as_type<ulong>(_e47) + as_type<ulong>(metal::long3(_e44).z));
     ulong _e51 = input_uniform.val_u64_;
     long _e53 = val;
-    val = _e53 + as_type<long>(_e51);
+    val = as_type<long>(as_type<ulong>(_e53) + as_type<ulong>(as_type<long>(_e51)));
     metal::ulong2 _e57 = input_uniform.val_u64_2_;
     long _e60 = val;
-    val = _e60 + as_type<metal::long2>(_e57).y;
+    val = as_type<long>(as_type<ulong>(_e60) + as_type<ulong>(as_type<metal::long2>(_e57).y));
     metal::ulong3 _e64 = input_uniform.val_u64_3_;
     long _e67 = val;
-    val = _e67 + as_type<metal::long3>(_e64).z;
+    val = as_type<long>(as_type<ulong>(_e67) + as_type<ulong>(as_type<metal::long3>(_e64).z));
     metal::ulong4 _e71 = input_uniform.val_u64_4_;
     long _e74 = val;
-    val = _e74 + as_type<metal::long4>(_e71).w;
+    val = as_type<long>(as_type<ulong>(_e74) + as_type<ulong>(as_type<metal::long4>(_e71).w));
     long _e80 = input_uniform.val_i64_;
     long _e83 = input_storage.val_i64_;
-    output.val_i64_ = _e80 + _e83;
+    output.val_i64_ = as_type<long>(as_type<ulong>(_e80) + as_type<ulong>(_e83));
     metal::long2 _e89 = input_uniform.val_i64_2_;
     metal::long2 _e92 = input_storage.val_i64_2_;
-    output.val_i64_2_ = _e89 + _e92;
+    output.val_i64_2_ = as_type<metal::long2>(as_type<metal::ulong2>(_e89) + as_type<metal::ulong2>(_e92));
     metal::long3 _e98 = input_uniform.val_i64_3_;
     metal::long3 _e101 = input_storage.val_i64_3_;
-    output.val_i64_3_ = _e98 + _e101;
+    output.val_i64_3_ = as_type<metal::long3>(as_type<metal::ulong3>(_e98) + as_type<metal::ulong3>(_e101));
     metal::long4 _e107 = input_uniform.val_i64_4_;
     metal::long4 _e110 = input_storage.val_i64_4_;
-    output.val_i64_4_ = _e107 + _e110;
+    output.val_i64_4_ = as_type<metal::long4>(as_type<metal::ulong4>(_e107) + as_type<metal::ulong4>(_e110));
     type_12 _e116 = input_arrays.val_i64_array_2_;
     output_arrays.val_i64_array_2_ = _e116;
     long _e117 = val;
     long _e119 = val;
-    val = _e119 + metal::abs(_e117);
+    val = as_type<long>(as_type<ulong>(_e119) + as_type<ulong>(metal::abs(_e117)));
     long _e121 = val;
     long _e122 = val;
     long _e123 = val;
     long _e125 = val;
-    val = _e125 + metal::clamp(_e121, _e122, _e123);
+    val = as_type<long>(as_type<ulong>(_e125) + as_type<ulong>(metal::clamp(_e121, _e122, _e123)));
     long _e127 = val;
     metal::long2 _e128 = metal::long2(_e127);
     long _e129 = val;
     metal::long2 _e130 = metal::long2(_e129);
     long _e132 = val;
-    val = _e132 + ( + _e128.x * _e130.x + _e128.y * _e130.y);
+    val = as_type<long>(as_type<ulong>(_e132) + as_type<ulong>(( + _e128.x * _e130.x + _e128.y * _e130.y)));
     long _e134 = val;
     long _e135 = val;
     long _e137 = val;
-    val = _e137 + metal::max(_e134, _e135);
+    val = as_type<long>(as_type<ulong>(_e137) + as_type<ulong>(metal::max(_e134, _e135)));
     long _e139 = val;
     long _e140 = val;
     long _e142 = val;
-    val = _e142 + metal::min(_e139, _e140);
+    val = as_type<long>(as_type<ulong>(_e142) + as_type<ulong>(metal::min(_e139, _e140)));
     long _e144 = val;
     long _e146 = val;
-    val = _e146 + metal::select(metal::select(-1, 1, (_e144 > 0)), 0, (_e144 == 0));
+    val = as_type<long>(as_type<ulong>(_e146) + as_type<ulong>(metal::select(metal::select(long(-1), long(1), (_e144 > 0)), long(0), (_e144 == 0))));
     long _e148 = val;
     return _e148;
 }
@@ -139,7 +139,7 @@ ulong uint64_function(
     int _e26 = input_uniform.val_i32_;
     ulong _e27 = val_1;
     ulong _e31 = val_1;
-    val_1 = _e31 + static_cast<ulong>(_e26 + static_cast<int>(_e27));
+    val_1 = _e31 + static_cast<ulong>(as_type<int>(as_type<uint>(_e26) + as_type<uint>(static_cast<int>(_e27))));
     float _e35 = input_uniform.val_f32_;
     ulong _e36 = val_1;
     ulong _e40 = val_1;
diff --git a/naga/tests/out/msl/operators.msl b/naga/tests/out/msl/operators.msl
index 85fba28c33..da483ed794 100644
--- a/naga/tests/out/msl/operators.msl
+++ b/naga/tests/out/msl/operators.msl
@@ -19,7 +19,7 @@ metal::float4 builtins(
     float b1_ = as_type<float>(1);
     metal::float4 b2_ = as_type<metal::float4>(v_i32_one);
     metal::int4 v_i32_zero = metal::int4(0, 0, 0, 0);
-    return ((((static_cast<metal::float4>(metal::int4(s1_) + v_i32_zero) + s2_) + m1_) + m2_) + metal::float4(b1_)) + b2_;
+    return ((((static_cast<metal::float4>(as_type<metal::int4>(as_type<metal::uint4>(metal::int4(s1_)) + as_type<metal::uint4>(v_i32_zero))) + s2_) + m1_) + m2_) + metal::float4(b1_)) + b2_;
 }
 
 metal::float4 splat(
@@ -68,22 +68,22 @@ void arithmetic(
     float neg0_1 = -(1.0);
     metal::int2 neg1_1 = -(metal::int2(1));
     metal::float2 neg2_ = -(metal::float2(1.0));
-    int add0_ = 2 + 1;
+    int add0_ = as_type<int>(as_type<uint>(2) + as_type<uint>(1));
     uint add1_ = 2u + 1u;
     float add2_ = 2.0 + 1.0;
-    metal::int2 add3_ = metal::int2(2) + metal::int2(1);
+    metal::int2 add3_ = as_type<metal::int2>(as_type<metal::uint2>(metal::int2(2)) + as_type<metal::uint2>(metal::int2(1)));
     metal::uint3 add4_ = metal::uint3(2u) + metal::uint3(1u);
     metal::float4 add5_ = metal::float4(2.0) + metal::float4(1.0);
-    int sub0_ = 2 - 1;
+    int sub0_ = as_type<int>(as_type<uint>(2) - as_type<uint>(1));
     uint sub1_ = 2u - 1u;
     float sub2_ = 2.0 - 1.0;
-    metal::int2 sub3_ = metal::int2(2) - metal::int2(1);
+    metal::int2 sub3_ = as_type<metal::int2>(as_type<metal::uint2>(metal::int2(2)) - as_type<metal::uint2>(metal::int2(1)));
     metal::uint3 sub4_ = metal::uint3(2u) - metal::uint3(1u);
     metal::float4 sub5_ = metal::float4(2.0) - metal::float4(1.0);
-    int mul0_ = 2 * 1;
+    int mul0_ = as_type<int>(as_type<uint>(2) * as_type<uint>(1));
     uint mul1_ = 2u * 1u;
     float mul2_ = 2.0 * 1.0;
-    metal::int2 mul3_ = metal::int2(2) * metal::int2(1);
+    metal::int2 mul3_ = as_type<metal::int2>(as_type<metal::uint2>(metal::int2(2)) * as_type<metal::uint2>(metal::int2(1)));
     metal::uint3 mul4_ = metal::uint3(2u) * metal::uint3(1u);
     metal::float4 mul5_ = metal::float4(2.0) * metal::float4(1.0);
     int div0_ = 2 / 1;
@@ -99,20 +99,20 @@ void arithmetic(
     metal::uint3 rem4_ = metal::uint3(2u) % metal::uint3(1u);
     metal::float4 rem5_ = metal::fmod(metal::float4(2.0), metal::float4(1.0));
     {
-        metal::int2 add0_1 = metal::int2(2) + metal::int2(1);
-        metal::int2 add1_1 = metal::int2(2) + metal::int2(1);
+        metal::int2 add0_1 = as_type<metal::int2>(as_type<metal::uint2>(metal::int2(2)) + as_type<metal::uint2>(metal::int2(1)));
+        metal::int2 add1_1 = as_type<metal::int2>(as_type<metal::uint2>(metal::int2(2)) + as_type<metal::uint2>(metal::int2(1)));
         metal::uint2 add2_1 = metal::uint2(2u) + metal::uint2(1u);
         metal::uint2 add3_1 = metal::uint2(2u) + metal::uint2(1u);
         metal::float2 add4_1 = metal::float2(2.0) + metal::float2(1.0);
         metal::float2 add5_1 = metal::float2(2.0) + metal::float2(1.0);
-        metal::int2 sub0_1 = metal::int2(2) - metal::int2(1);
-        metal::int2 sub1_1 = metal::int2(2) - metal::int2(1);
+        metal::int2 sub0_1 = as_type<metal::int2>(as_type<metal::uint2>(metal::int2(2)) - as_type<metal::uint2>(metal::int2(1)));
+        metal::int2 sub1_1 = as_type<metal::int2>(as_type<metal::uint2>(metal::int2(2)) - as_type<metal::uint2>(metal::int2(1)));
         metal::uint2 sub2_1 = metal::uint2(2u) - metal::uint2(1u);
         metal::uint2 sub3_1 = metal::uint2(2u) - metal::uint2(1u);
         metal::float2 sub4_1 = metal::float2(2.0) - metal::float2(1.0);
         metal::float2 sub5_1 = metal::float2(2.0) - metal::float2(1.0);
-        metal::int2 mul0_1 = metal::int2(2) * 1;
-        metal::int2 mul1_1 = 2 * metal::int2(1);
+        metal::int2 mul0_1 = as_type<metal::int2>(as_type<metal::uint2>(metal::int2(2)) * as_type<uint>(1));
+        metal::int2 mul1_1 = as_type<metal::int2>(as_type<uint>(2) * as_type<metal::uint2>(metal::int2(1)));
         metal::uint2 mul2_1 = metal::uint2(2u) * 1u;
         metal::uint2 mul3_1 = 2u * metal::uint2(1u);
         metal::float2 mul4_1 = metal::float2(2.0) * 1.0;
@@ -213,12 +213,12 @@ void assignment(
     metal::int3 vec0_ = metal::int3 {};
     a_1 = 1;
     int _e5 = a_1;
-    a_1 = _e5 + 1;
+    a_1 = as_type<int>(as_type<uint>(_e5) + as_type<uint>(1));
     int _e7 = a_1;
-    a_1 = _e7 - 1;
+    a_1 = as_type<int>(as_type<uint>(_e7) - as_type<uint>(1));
     int _e9 = a_1;
     int _e10 = a_1;
-    a_1 = _e10 * _e9;
+    a_1 = as_type<int>(as_type<uint>(_e10) * as_type<uint>(_e9));
     int _e12 = a_1;
     int _e13 = a_1;
     a_1 = _e13 / _e12;
@@ -235,13 +235,13 @@ void assignment(
     int _e25 = a_1;
     a_1 = _e25 >> 1u;
     int _e28 = a_1;
-    a_1 = _e28 + 1;
+    a_1 = as_type<int>(as_type<uint>(_e28) + as_type<uint>(1));
     int _e31 = a_1;
-    a_1 = _e31 - 1;
+    a_1 = as_type<int>(as_type<uint>(_e31) - as_type<uint>(1));
     int _e37 = vec0_[1];
-    vec0_[1] = _e37 + 1;
+    vec0_[1] = as_type<int>(as_type<uint>(_e37) + as_type<uint>(1));
     int _e41 = vec0_[1];
-    vec0_[1] = _e41 - 1;
+    vec0_[1] = as_type<int>(as_type<uint>(_e41) - as_type<uint>(1));
     return;
 }