[naga msl-out] Avoid undefined behaviour due to signed integer overflow (#6959)

Signed integer overflow is undefined behaviour in MSL. However, signed
integers are defined to be two's complement. This allows us to cast
signed values to their corresponding unsigned type, perform the
arithmetic on these unsigned values (which has defined overflow
behaviour) then cast the result back to signed.

Care must be taken when emitting the isign polyfill, which uses
metal::select(). We must ensure the -1, 0, and 1 literals used as
inputs to select() have the correct width, else bitcasting the output
of select() will fail due to mismatched widths.

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)]

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,

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))];
 }

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);

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

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;
 }
 

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)));

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;

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;
 }