[msl] Polyfill signed integer arithmetic Avoid UB caused by signed integer overflow by bitcasting to unsigned integers and back again. Bug: 42251016 Change-Id: I98078a9a4e9603b3ab6faa4dbc0fc0812a5dcc3e Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/205235 Reviewed-by: dan sinclair <dsinclair@chromium.org> Commit-Queue: James Price <jrprice@google.com>

diff --git a/test/tint/access/ptr.wgsl.expected.ir.msl b/test/tint/access/ptr.wgsl.expected.ir.msl
index 53bee9e..23072a5 100644
--- a/test/tint/access/ptr.wgsl.expected.ir.msl
+++ b/test/tint/access/ptr.wgsl.expected.ir.msl

@@ -21,16 +21,16 @@
 
 int accept_ptr_deref_call_func(thread int* const val) {
   int const v = (*val);
-  return (v + accept_value((*val)));
+  return as_type<int>((as_type<uint>(v) + as_type<uint>(accept_value((*val)))));
 }
 
 int accept_ptr_deref_pass_through(thread int* const val) {
   int const v_1 = (*val);
-  return (v_1 + accept_ptr_deref_call_func(val));
+  return as_type<int>((as_type<uint>(v_1) + as_type<uint>(accept_ptr_deref_call_func(val))));
 }
 
 int accept_ptr_to_struct_and_access(thread S* const val) {
-  return ((*val).a + (*val).b);
+  return as_type<int>((as_type<uint>((*val).a) + as_type<uint>((*val).b)));
 }
 
 int accept_ptr_to_struct_access_pass_ptr(thread S* const val) {
@@ -63,11 +63,11 @@
   float3 v4 = float3(0.0f);
   int const t1 = atomic_load_explicit(tint_module_vars.g1, memory_order_relaxed);
   int const v_2 = accept_ptr_deref_pass_through((&v1));
-  int const v_3 = (v_2 + accept_ptr_to_struct_and_access((&v2)));
-  int const v_4 = (v_3 + accept_ptr_to_struct_and_access(v3));
-  int const v_5 = (v_4 + accept_ptr_vec_access_elements((&v4)));
-  int const v_6 = (v_5 + accept_ptr_to_struct_access_pass_ptr((&v2)));
-  (*tint_module_vars.s) = ((v_6 + call_builtin_with_mod_scope_ptr(tint_module_vars)) + t1);
+  int const v_3 = as_type<int>((as_type<uint>(v_2) + as_type<uint>(accept_ptr_to_struct_and_access((&v2)))));
+  int const v_4 = as_type<int>((as_type<uint>(v_3) + as_type<uint>(accept_ptr_to_struct_and_access(v3))));
+  int const v_5 = as_type<int>((as_type<uint>(v_4) + as_type<uint>(accept_ptr_vec_access_elements((&v4)))));
+  int const v_6 = as_type<int>((as_type<uint>(v_5) + as_type<uint>(accept_ptr_to_struct_access_pass_ptr((&v2)))));
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_6) + as_type<uint>(call_builtin_with_mod_scope_ptr(tint_module_vars))))) + as_type<uint>(t1)));
 }
 
 kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], device int* s [[buffer(0)]], threadgroup tint_symbol_2* v_7 [[threadgroup(0)]]) {

diff --git a/test/tint/array/assign_to_subexpr.wgsl.expected.ir.msl b/test/tint/array/assign_to_subexpr.wgsl.expected.ir.msl
index 265bd44..c7b0190 100644
--- a/test/tint/array/assign_to_subexpr.wgsl.expected.ir.msl
+++ b/test/tint/array/assign_to_subexpr.wgsl.expected.ir.msl

@@ -34,7 +34,7 @@
   (*dst_ptr) = src;
   (*dst_struct_ptr).arr = src;
   (*dst_array_ptr)[0] = src;
-  return (((*dst_ptr)[0] + (*dst_struct_ptr).arr[0]) + (*dst_array_ptr)[0][0]);
+  return as_type<int>((as_type<uint>(as_type<int>((as_type<uint>((*dst_ptr)[0]) + as_type<uint>((*dst_struct_ptr).arr[0])))) + as_type<uint>((*dst_array_ptr)[0][0])));
 }
 
 kernel void tint_symbol(device int* s [[buffer(0)]]) {

diff --git a/test/tint/array/type_initializer.wgsl.expected.ir.msl b/test/tint/array/type_initializer.wgsl.expected.ir.msl
index 1f0ec3b..bebcd67 100644
--- a/test/tint/array/type_initializer.wgsl.expected.ir.msl
+++ b/test/tint/array/type_initializer.wgsl.expected.ir.msl

@@ -22,16 +22,16 @@
   int const x = 42;
   tint_array<int, 4> const empty = tint_array<int, 4>{};
   tint_array<int, 4> const nonempty = tint_array<int, 4>{1, 2, 3, 4};
-  tint_array<int, 4> const nonempty_with_expr = tint_array<int, 4>{1, x, (x + 1), nonempty[3]};
+  tint_array<int, 4> const nonempty_with_expr = tint_array<int, 4>{1, x, as_type<int>((as_type<uint>(x) + as_type<uint>(1))), nonempty[3]};
   tint_array<tint_array<tint_array<int, 4>, 3>, 2> const nested_empty = tint_array<tint_array<tint_array<int, 4>, 3>, 2>{};
   tint_array<tint_array<tint_array<int, 4>, 3>, 2> const nested_nonempty = tint_array<tint_array<tint_array<int, 4>, 3>, 2>{tint_array<tint_array<int, 4>, 3>{tint_array<int, 4>{1, 2, 3, 4}, tint_array<int, 4>{5, 6, 7, 8}, tint_array<int, 4>{9, 10, 11, 12}}, tint_array<tint_array<int, 4>, 3>{tint_array<int, 4>{13, 14, 15, 16}, tint_array<int, 4>{17, 18, 19, 20}, tint_array<int, 4>{21, 22, 23, 24}}};
-  tint_array<int, 4> const v = tint_array<int, 4>{1, 2, x, (x + 1)};
-  tint_array<tint_array<tint_array<int, 4>, 3>, 2> const nested_nonempty_with_expr = tint_array<tint_array<tint_array<int, 4>, 3>, 2>{tint_array<tint_array<int, 4>, 3>{v, tint_array<int, 4>{5, 6, nonempty[2], (nonempty[3] + 1)}, nonempty}, nested_nonempty[1]};
+  tint_array<int, 4> const v = tint_array<int, 4>{1, 2, x, as_type<int>((as_type<uint>(x) + as_type<uint>(1)))};
+  tint_array<tint_array<tint_array<int, 4>, 3>, 2> const nested_nonempty_with_expr = tint_array<tint_array<tint_array<int, 4>, 3>, 2>{tint_array<tint_array<int, 4>, 3>{v, tint_array<int, 4>{5, 6, nonempty[2], as_type<int>((as_type<uint>(nonempty[3]) + as_type<uint>(1)))}, nonempty}, nested_nonempty[1]};
   int const subexpr_empty = 0;
   int const subexpr_nonempty = 3;
-  int const subexpr_nonempty_with_expr = tint_array<int, 4>{1, x, (x + 1), nonempty[3]}[2];
+  int const subexpr_nonempty_with_expr = tint_array<int, 4>{1, x, as_type<int>((as_type<uint>(x) + as_type<uint>(1))), nonempty[3]}[2];
   tint_array<int, 4> const subexpr_nested_empty = tint_array<int, 4>{};
   tint_array<int, 4> const subexpr_nested_nonempty = tint_array<int, 4>{5, 6, 7, 8};
-  tint_array<int, 4> const subexpr_nested_nonempty_with_expr = tint_array<tint_array<int, 4>, 2>{tint_array<int, 4>{1, x, (x + 1), nonempty[3]}, nested_nonempty[1][2]}[1];
-  (*tint_module_vars.s) = (((((((((((empty[0] + nonempty[0]) + nonempty_with_expr[0]) + nested_empty[0][0][0]) + nested_nonempty[0][0][0]) + nested_nonempty_with_expr[0][0][0]) + subexpr_empty) + subexpr_nonempty) + subexpr_nonempty_with_expr) + subexpr_nested_empty[0]) + subexpr_nested_nonempty[0]) + subexpr_nested_nonempty_with_expr[0]);
+  tint_array<int, 4> const subexpr_nested_nonempty_with_expr = tint_array<tint_array<int, 4>, 2>{tint_array<int, 4>{1, x, as_type<int>((as_type<uint>(x) + as_type<uint>(1))), nonempty[3]}, nested_nonempty[1][2]}[1];
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(empty[0]) + as_type<uint>(nonempty[0])))) + as_type<uint>(nonempty_with_expr[0])))) + as_type<uint>(nested_empty[0][0][0])))) + as_type<uint>(nested_nonempty[0][0][0])))) + as_type<uint>(nested_nonempty_with_expr[0][0][0])))) + as_type<uint>(subexpr_empty)))) + as_type<uint>(subexpr_nonempty)))) + as_type<uint>(subexpr_nonempty_with_expr)))) + as_type<uint>(subexpr_nested_empty[0])))) + as_type<uint>(subexpr_nested_nonempty[0])))) + as_type<uint>(subexpr_nested_nonempty_with_expr[0])));
 }

diff --git a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.ir.msl b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.ir.msl
index 657609b..3e98afa 100644
--- a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.ir.msl

@@ -99,24 +99,24 @@
   float4x3 const mat4x3_f32 = float4x3(v_7, v_8, v_9, float3(v_6[3u].packed));
   float4x4 const mat4x4_f32 = (*tint_module_vars.sb).arr[idx].mat4x4_f32;
   tint_array<float3, 2> const arr2_vec3_f32 = tint_load_array_packed_vec3((&(*tint_module_vars.sb).arr[idx].arr2_vec3_f32));
-  int const v_10 = (tint_f32_to_i32(scalar_f32) + scalar_i32);
-  int const v_11 = (v_10 + int(scalar_u32));
-  int const v_12 = ((v_11 + tint_f32_to_i32(vec2_f32[0u])) + vec2_i32[0u]);
-  int const v_13 = (v_12 + int(vec2_u32[0u]));
-  int const v_14 = ((v_13 + tint_f32_to_i32(vec3_f32[1u])) + vec3_i32[1u]);
-  int const v_15 = (v_14 + int(vec3_u32[1u]));
-  int const v_16 = ((v_15 + tint_f32_to_i32(vec4_f32[2u])) + vec4_i32[2u]);
-  int const v_17 = (v_16 + int(vec4_u32[2u]));
-  int const v_18 = (v_17 + tint_f32_to_i32(mat2x2_f32[0][0u]));
-  int const v_19 = (v_18 + tint_f32_to_i32(mat2x3_f32[0][0u]));
-  int const v_20 = (v_19 + tint_f32_to_i32(mat2x4_f32[0][0u]));
-  int const v_21 = (v_20 + tint_f32_to_i32(mat3x2_f32[0][0u]));
-  int const v_22 = (v_21 + tint_f32_to_i32(mat3x3_f32[0][0u]));
-  int const v_23 = (v_22 + tint_f32_to_i32(mat3x4_f32[0][0u]));
-  int const v_24 = (v_23 + tint_f32_to_i32(mat4x2_f32[0][0u]));
-  int const v_25 = (v_24 + tint_f32_to_i32(mat4x3_f32[0][0u]));
-  int const v_26 = (v_25 + tint_f32_to_i32(mat4x4_f32[0][0u]));
-  (*tint_module_vars.s) = (v_26 + tint_f32_to_i32(arr2_vec3_f32[0][0u]));
+  int const v_10 = as_type<int>((as_type<uint>(tint_f32_to_i32(scalar_f32)) + as_type<uint>(scalar_i32)));
+  int const v_11 = as_type<int>((as_type<uint>(v_10) + as_type<uint>(int(scalar_u32))));
+  int const v_12 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_11) + as_type<uint>(tint_f32_to_i32(vec2_f32[0u]))))) + as_type<uint>(vec2_i32[0u])));
+  int const v_13 = as_type<int>((as_type<uint>(v_12) + as_type<uint>(int(vec2_u32[0u]))));
+  int const v_14 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_13) + as_type<uint>(tint_f32_to_i32(vec3_f32[1u]))))) + as_type<uint>(vec3_i32[1u])));
+  int const v_15 = as_type<int>((as_type<uint>(v_14) + as_type<uint>(int(vec3_u32[1u]))));
+  int const v_16 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_15) + as_type<uint>(tint_f32_to_i32(vec4_f32[2u]))))) + as_type<uint>(vec4_i32[2u])));
+  int const v_17 = as_type<int>((as_type<uint>(v_16) + as_type<uint>(int(vec4_u32[2u]))));
+  int const v_18 = as_type<int>((as_type<uint>(v_17) + as_type<uint>(tint_f32_to_i32(mat2x2_f32[0][0u]))));
+  int const v_19 = as_type<int>((as_type<uint>(v_18) + as_type<uint>(tint_f32_to_i32(mat2x3_f32[0][0u]))));
+  int const v_20 = as_type<int>((as_type<uint>(v_19) + as_type<uint>(tint_f32_to_i32(mat2x4_f32[0][0u]))));
+  int const v_21 = as_type<int>((as_type<uint>(v_20) + as_type<uint>(tint_f32_to_i32(mat3x2_f32[0][0u]))));
+  int const v_22 = as_type<int>((as_type<uint>(v_21) + as_type<uint>(tint_f32_to_i32(mat3x3_f32[0][0u]))));
+  int const v_23 = as_type<int>((as_type<uint>(v_22) + as_type<uint>(tint_f32_to_i32(mat3x4_f32[0][0u]))));
+  int const v_24 = as_type<int>((as_type<uint>(v_23) + as_type<uint>(tint_f32_to_i32(mat4x2_f32[0][0u]))));
+  int const v_25 = as_type<int>((as_type<uint>(v_24) + as_type<uint>(tint_f32_to_i32(mat4x3_f32[0][0u]))));
+  int const v_26 = as_type<int>((as_type<uint>(v_25) + as_type<uint>(tint_f32_to_i32(mat4x4_f32[0][0u]))));
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(v_26) + as_type<uint>(tint_f32_to_i32(arr2_vec3_f32[0][0u]))));
 }
 
 kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], const device S_packed_vec3* sb [[buffer(1)]], device int* s [[buffer(0)]]) {

diff --git a/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.ir.msl b/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.ir.msl
index 8f43b3e..97e35d4 100644
--- a/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/dynamic_index/read_f16.wgsl.expected.ir.msl

@@ -149,38 +149,38 @@
   half4x4 const mat4x4_f16 = (*tint_module_vars.sb).arr[idx].mat4x4_f16;
   tint_array<float3, 2> const arr2_vec3_f32 = tint_load_array_packed_vec3((&(*tint_module_vars.sb).arr[idx].arr2_vec3_f32));
   tint_array<half4x2, 2> const arr2_mat4x2_f16 = (*tint_module_vars.sb).arr[idx].arr2_mat4x2_f16;
-  int const v_19 = (tint_f32_to_i32(scalar_f32) + scalar_i32);
-  int const v_20 = (v_19 + int(scalar_u32));
-  int const v_21 = (v_20 + tint_f16_to_i32(scalar_f16));
-  int const v_22 = ((v_21 + tint_f32_to_i32(vec2_f32[0u])) + vec2_i32[0u]);
-  int const v_23 = (v_22 + int(vec2_u32[0u]));
-  int const v_24 = (v_23 + tint_f16_to_i32(vec2_f16[0u]));
-  int const v_25 = ((v_24 + tint_f32_to_i32(vec3_f32[1u])) + vec3_i32[1u]);
-  int const v_26 = (v_25 + int(vec3_u32[1u]));
-  int const v_27 = (v_26 + tint_f16_to_i32(vec3_f16[1u]));
-  int const v_28 = ((v_27 + tint_f32_to_i32(vec4_f32[2u])) + vec4_i32[2u]);
-  int const v_29 = (v_28 + int(vec4_u32[2u]));
-  int const v_30 = (v_29 + tint_f16_to_i32(vec4_f16[2u]));
-  int const v_31 = (v_30 + tint_f32_to_i32(mat2x2_f32[0][0u]));
-  int const v_32 = (v_31 + tint_f32_to_i32(mat2x3_f32[0][0u]));
-  int const v_33 = (v_32 + tint_f32_to_i32(mat2x4_f32[0][0u]));
-  int const v_34 = (v_33 + tint_f32_to_i32(mat3x2_f32[0][0u]));
-  int const v_35 = (v_34 + tint_f32_to_i32(mat3x3_f32[0][0u]));
-  int const v_36 = (v_35 + tint_f32_to_i32(mat3x4_f32[0][0u]));
-  int const v_37 = (v_36 + tint_f32_to_i32(mat4x2_f32[0][0u]));
-  int const v_38 = (v_37 + tint_f32_to_i32(mat4x3_f32[0][0u]));
-  int const v_39 = (v_38 + tint_f32_to_i32(mat4x4_f32[0][0u]));
-  int const v_40 = (v_39 + tint_f16_to_i32(mat2x2_f16[0][0u]));
-  int const v_41 = (v_40 + tint_f16_to_i32(mat2x3_f16[0][0u]));
-  int const v_42 = (v_41 + tint_f16_to_i32(mat2x4_f16[0][0u]));
-  int const v_43 = (v_42 + tint_f16_to_i32(mat3x2_f16[0][0u]));
-  int const v_44 = (v_43 + tint_f16_to_i32(mat3x3_f16[0][0u]));
-  int const v_45 = (v_44 + tint_f16_to_i32(mat3x4_f16[0][0u]));
-  int const v_46 = (v_45 + tint_f16_to_i32(mat4x2_f16[0][0u]));
-  int const v_47 = (v_46 + tint_f16_to_i32(mat4x3_f16[0][0u]));
-  int const v_48 = (v_47 + tint_f16_to_i32(mat4x4_f16[0][0u]));
-  int const v_49 = (v_48 + tint_f16_to_i32(arr2_mat4x2_f16[0][0][0u]));
-  (*tint_module_vars.s) = (v_49 + tint_f32_to_i32(arr2_vec3_f32[0][0u]));
+  int const v_19 = as_type<int>((as_type<uint>(tint_f32_to_i32(scalar_f32)) + as_type<uint>(scalar_i32)));
+  int const v_20 = as_type<int>((as_type<uint>(v_19) + as_type<uint>(int(scalar_u32))));
+  int const v_21 = as_type<int>((as_type<uint>(v_20) + as_type<uint>(tint_f16_to_i32(scalar_f16))));
+  int const v_22 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_21) + as_type<uint>(tint_f32_to_i32(vec2_f32[0u]))))) + as_type<uint>(vec2_i32[0u])));
+  int const v_23 = as_type<int>((as_type<uint>(v_22) + as_type<uint>(int(vec2_u32[0u]))));
+  int const v_24 = as_type<int>((as_type<uint>(v_23) + as_type<uint>(tint_f16_to_i32(vec2_f16[0u]))));
+  int const v_25 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_24) + as_type<uint>(tint_f32_to_i32(vec3_f32[1u]))))) + as_type<uint>(vec3_i32[1u])));
+  int const v_26 = as_type<int>((as_type<uint>(v_25) + as_type<uint>(int(vec3_u32[1u]))));
+  int const v_27 = as_type<int>((as_type<uint>(v_26) + as_type<uint>(tint_f16_to_i32(vec3_f16[1u]))));
+  int const v_28 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_27) + as_type<uint>(tint_f32_to_i32(vec4_f32[2u]))))) + as_type<uint>(vec4_i32[2u])));
+  int const v_29 = as_type<int>((as_type<uint>(v_28) + as_type<uint>(int(vec4_u32[2u]))));
+  int const v_30 = as_type<int>((as_type<uint>(v_29) + as_type<uint>(tint_f16_to_i32(vec4_f16[2u]))));
+  int const v_31 = as_type<int>((as_type<uint>(v_30) + as_type<uint>(tint_f32_to_i32(mat2x2_f32[0][0u]))));
+  int const v_32 = as_type<int>((as_type<uint>(v_31) + as_type<uint>(tint_f32_to_i32(mat2x3_f32[0][0u]))));
+  int const v_33 = as_type<int>((as_type<uint>(v_32) + as_type<uint>(tint_f32_to_i32(mat2x4_f32[0][0u]))));
+  int const v_34 = as_type<int>((as_type<uint>(v_33) + as_type<uint>(tint_f32_to_i32(mat3x2_f32[0][0u]))));
+  int const v_35 = as_type<int>((as_type<uint>(v_34) + as_type<uint>(tint_f32_to_i32(mat3x3_f32[0][0u]))));
+  int const v_36 = as_type<int>((as_type<uint>(v_35) + as_type<uint>(tint_f32_to_i32(mat3x4_f32[0][0u]))));
+  int const v_37 = as_type<int>((as_type<uint>(v_36) + as_type<uint>(tint_f32_to_i32(mat4x2_f32[0][0u]))));
+  int const v_38 = as_type<int>((as_type<uint>(v_37) + as_type<uint>(tint_f32_to_i32(mat4x3_f32[0][0u]))));
+  int const v_39 = as_type<int>((as_type<uint>(v_38) + as_type<uint>(tint_f32_to_i32(mat4x4_f32[0][0u]))));
+  int const v_40 = as_type<int>((as_type<uint>(v_39) + as_type<uint>(tint_f16_to_i32(mat2x2_f16[0][0u]))));
+  int const v_41 = as_type<int>((as_type<uint>(v_40) + as_type<uint>(tint_f16_to_i32(mat2x3_f16[0][0u]))));
+  int const v_42 = as_type<int>((as_type<uint>(v_41) + as_type<uint>(tint_f16_to_i32(mat2x4_f16[0][0u]))));
+  int const v_43 = as_type<int>((as_type<uint>(v_42) + as_type<uint>(tint_f16_to_i32(mat3x2_f16[0][0u]))));
+  int const v_44 = as_type<int>((as_type<uint>(v_43) + as_type<uint>(tint_f16_to_i32(mat3x3_f16[0][0u]))));
+  int const v_45 = as_type<int>((as_type<uint>(v_44) + as_type<uint>(tint_f16_to_i32(mat3x4_f16[0][0u]))));
+  int const v_46 = as_type<int>((as_type<uint>(v_45) + as_type<uint>(tint_f16_to_i32(mat4x2_f16[0][0u]))));
+  int const v_47 = as_type<int>((as_type<uint>(v_46) + as_type<uint>(tint_f16_to_i32(mat4x3_f16[0][0u]))));
+  int const v_48 = as_type<int>((as_type<uint>(v_47) + as_type<uint>(tint_f16_to_i32(mat4x4_f16[0][0u]))));
+  int const v_49 = as_type<int>((as_type<uint>(v_48) + as_type<uint>(tint_f16_to_i32(arr2_mat4x2_f16[0][0][0u]))));
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(v_49) + as_type<uint>(tint_f32_to_i32(arr2_vec3_f32[0][0u]))));
 }
 
 kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], const device S_packed_vec3* sb [[buffer(1)]], device int* s [[buffer(0)]]) {

diff --git a/test/tint/buffer/storage/static_index/read.wgsl.expected.ir.msl b/test/tint/buffer/storage/static_index/read.wgsl.expected.ir.msl
index 6a6c7b6..6bfd190 100644
--- a/test/tint/buffer/storage/static_index/read.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/static_index/read.wgsl.expected.ir.msl

@@ -106,22 +106,22 @@
   tint_array<float3, 2> const arr2_vec3_f32 = tint_load_array_packed_vec3((&(*tint_module_vars.sb).arr2_vec3_f32));
   Inner const struct_inner = (*tint_module_vars.sb).struct_inner;
   tint_array<Inner, 4> const array_struct_inner = (*tint_module_vars.sb).array_struct_inner;
-  int const v_10 = (tint_f32_to_i32(scalar_f32) + scalar_i32);
-  int const v_11 = (v_10 + int(scalar_u32));
-  int const v_12 = ((v_11 + tint_f32_to_i32(vec2_f32[0u])) + vec2_i32[0u]);
-  int const v_13 = (v_12 + int(vec2_u32[0u]));
-  int const v_14 = ((v_13 + tint_f32_to_i32(vec3_f32[1u])) + vec3_i32[1u]);
-  int const v_15 = (v_14 + int(vec3_u32[1u]));
-  int const v_16 = ((v_15 + tint_f32_to_i32(vec4_f32[2u])) + vec4_i32[2u]);
-  int const v_17 = (v_16 + int(vec4_u32[2u]));
-  int const v_18 = (v_17 + tint_f32_to_i32(mat2x2_f32[0][0u]));
-  int const v_19 = (v_18 + tint_f32_to_i32(mat2x3_f32[0][0u]));
-  int const v_20 = (v_19 + tint_f32_to_i32(mat2x4_f32[0][0u]));
-  int const v_21 = (v_20 + tint_f32_to_i32(mat3x2_f32[0][0u]));
-  int const v_22 = (v_21 + tint_f32_to_i32(mat3x3_f32[0][0u]));
-  int const v_23 = (v_22 + tint_f32_to_i32(mat3x4_f32[0][0u]));
-  int const v_24 = (v_23 + tint_f32_to_i32(mat4x2_f32[0][0u]));
-  int const v_25 = (v_24 + tint_f32_to_i32(mat4x3_f32[0][0u]));
-  int const v_26 = (v_25 + tint_f32_to_i32(mat4x4_f32[0][0u]));
-  (*tint_module_vars.s) = (((v_26 + tint_f32_to_i32(arr2_vec3_f32[0][0u])) + struct_inner.scalar_i32) + array_struct_inner[0].scalar_i32);
+  int const v_10 = as_type<int>((as_type<uint>(tint_f32_to_i32(scalar_f32)) + as_type<uint>(scalar_i32)));
+  int const v_11 = as_type<int>((as_type<uint>(v_10) + as_type<uint>(int(scalar_u32))));
+  int const v_12 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_11) + as_type<uint>(tint_f32_to_i32(vec2_f32[0u]))))) + as_type<uint>(vec2_i32[0u])));
+  int const v_13 = as_type<int>((as_type<uint>(v_12) + as_type<uint>(int(vec2_u32[0u]))));
+  int const v_14 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_13) + as_type<uint>(tint_f32_to_i32(vec3_f32[1u]))))) + as_type<uint>(vec3_i32[1u])));
+  int const v_15 = as_type<int>((as_type<uint>(v_14) + as_type<uint>(int(vec3_u32[1u]))));
+  int const v_16 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_15) + as_type<uint>(tint_f32_to_i32(vec4_f32[2u]))))) + as_type<uint>(vec4_i32[2u])));
+  int const v_17 = as_type<int>((as_type<uint>(v_16) + as_type<uint>(int(vec4_u32[2u]))));
+  int const v_18 = as_type<int>((as_type<uint>(v_17) + as_type<uint>(tint_f32_to_i32(mat2x2_f32[0][0u]))));
+  int const v_19 = as_type<int>((as_type<uint>(v_18) + as_type<uint>(tint_f32_to_i32(mat2x3_f32[0][0u]))));
+  int const v_20 = as_type<int>((as_type<uint>(v_19) + as_type<uint>(tint_f32_to_i32(mat2x4_f32[0][0u]))));
+  int const v_21 = as_type<int>((as_type<uint>(v_20) + as_type<uint>(tint_f32_to_i32(mat3x2_f32[0][0u]))));
+  int const v_22 = as_type<int>((as_type<uint>(v_21) + as_type<uint>(tint_f32_to_i32(mat3x3_f32[0][0u]))));
+  int const v_23 = as_type<int>((as_type<uint>(v_22) + as_type<uint>(tint_f32_to_i32(mat3x4_f32[0][0u]))));
+  int const v_24 = as_type<int>((as_type<uint>(v_23) + as_type<uint>(tint_f32_to_i32(mat4x2_f32[0][0u]))));
+  int const v_25 = as_type<int>((as_type<uint>(v_24) + as_type<uint>(tint_f32_to_i32(mat4x3_f32[0][0u]))));
+  int const v_26 = as_type<int>((as_type<uint>(v_25) + as_type<uint>(tint_f32_to_i32(mat4x4_f32[0][0u]))));
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_26) + as_type<uint>(tint_f32_to_i32(arr2_vec3_f32[0][0u]))))) + as_type<uint>(struct_inner.scalar_i32)))) + as_type<uint>(array_struct_inner[0].scalar_i32)));
 }

diff --git a/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.ir.msl b/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.ir.msl
index 98e297b..08ef76f 100644
--- a/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.ir.msl
+++ b/test/tint/buffer/storage/static_index/read_f16.wgsl.expected.ir.msl

@@ -158,36 +158,36 @@
   tint_array<half4x2, 2> const arr2_mat4x2_f16 = (*tint_module_vars.sb).arr2_mat4x2_f16;
   Inner const struct_inner = (*tint_module_vars.sb).struct_inner;
   tint_array<Inner, 4> const array_struct_inner = (*tint_module_vars.sb).array_struct_inner;
-  int const v_19 = (tint_f32_to_i32(scalar_f32) + scalar_i32);
-  int const v_20 = (v_19 + int(scalar_u32));
-  int const v_21 = (v_20 + tint_f16_to_i32(scalar_f16));
-  int const v_22 = ((v_21 + tint_f32_to_i32(vec2_f32[0u])) + vec2_i32[0u]);
-  int const v_23 = (v_22 + int(vec2_u32[0u]));
-  int const v_24 = (v_23 + tint_f16_to_i32(vec2_f16[0u]));
-  int const v_25 = ((v_24 + tint_f32_to_i32(vec3_f32[1u])) + vec3_i32[1u]);
-  int const v_26 = (v_25 + int(vec3_u32[1u]));
-  int const v_27 = (v_26 + tint_f16_to_i32(vec3_f16[1u]));
-  int const v_28 = ((v_27 + tint_f32_to_i32(vec4_f32[2u])) + vec4_i32[2u]);
-  int const v_29 = (v_28 + int(vec4_u32[2u]));
-  int const v_30 = (v_29 + tint_f16_to_i32(vec4_f16[2u]));
-  int const v_31 = (v_30 + tint_f32_to_i32(mat2x2_f32[0][0u]));
-  int const v_32 = (v_31 + tint_f32_to_i32(mat2x3_f32[0][0u]));
-  int const v_33 = (v_32 + tint_f32_to_i32(mat2x4_f32[0][0u]));
-  int const v_34 = (v_33 + tint_f32_to_i32(mat3x2_f32[0][0u]));
-  int const v_35 = (v_34 + tint_f32_to_i32(mat3x3_f32[0][0u]));
-  int const v_36 = (v_35 + tint_f32_to_i32(mat3x4_f32[0][0u]));
-  int const v_37 = (v_36 + tint_f32_to_i32(mat4x2_f32[0][0u]));
-  int const v_38 = (v_37 + tint_f32_to_i32(mat4x3_f32[0][0u]));
-  int const v_39 = (v_38 + tint_f32_to_i32(mat4x4_f32[0][0u]));
-  int const v_40 = (v_39 + tint_f16_to_i32(mat2x2_f16[0][0u]));
-  int const v_41 = (v_40 + tint_f16_to_i32(mat2x3_f16[0][0u]));
-  int const v_42 = (v_41 + tint_f16_to_i32(mat2x4_f16[0][0u]));
-  int const v_43 = (v_42 + tint_f16_to_i32(mat3x2_f16[0][0u]));
-  int const v_44 = (v_43 + tint_f16_to_i32(mat3x3_f16[0][0u]));
-  int const v_45 = (v_44 + tint_f16_to_i32(mat3x4_f16[0][0u]));
-  int const v_46 = (v_45 + tint_f16_to_i32(mat4x2_f16[0][0u]));
-  int const v_47 = (v_46 + tint_f16_to_i32(mat4x3_f16[0][0u]));
-  int const v_48 = (v_47 + tint_f16_to_i32(mat4x4_f16[0][0u]));
-  int const v_49 = (v_48 + tint_f32_to_i32(arr2_vec3_f32[0][0u]));
-  (*tint_module_vars.s) = (((v_49 + tint_f16_to_i32(arr2_mat4x2_f16[0][0][0u])) + struct_inner.scalar_i32) + array_struct_inner[0].scalar_i32);
+  int const v_19 = as_type<int>((as_type<uint>(tint_f32_to_i32(scalar_f32)) + as_type<uint>(scalar_i32)));
+  int const v_20 = as_type<int>((as_type<uint>(v_19) + as_type<uint>(int(scalar_u32))));
+  int const v_21 = as_type<int>((as_type<uint>(v_20) + as_type<uint>(tint_f16_to_i32(scalar_f16))));
+  int const v_22 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_21) + as_type<uint>(tint_f32_to_i32(vec2_f32[0u]))))) + as_type<uint>(vec2_i32[0u])));
+  int const v_23 = as_type<int>((as_type<uint>(v_22) + as_type<uint>(int(vec2_u32[0u]))));
+  int const v_24 = as_type<int>((as_type<uint>(v_23) + as_type<uint>(tint_f16_to_i32(vec2_f16[0u]))));
+  int const v_25 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_24) + as_type<uint>(tint_f32_to_i32(vec3_f32[1u]))))) + as_type<uint>(vec3_i32[1u])));
+  int const v_26 = as_type<int>((as_type<uint>(v_25) + as_type<uint>(int(vec3_u32[1u]))));
+  int const v_27 = as_type<int>((as_type<uint>(v_26) + as_type<uint>(tint_f16_to_i32(vec3_f16[1u]))));
+  int const v_28 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_27) + as_type<uint>(tint_f32_to_i32(vec4_f32[2u]))))) + as_type<uint>(vec4_i32[2u])));
+  int const v_29 = as_type<int>((as_type<uint>(v_28) + as_type<uint>(int(vec4_u32[2u]))));
+  int const v_30 = as_type<int>((as_type<uint>(v_29) + as_type<uint>(tint_f16_to_i32(vec4_f16[2u]))));
+  int const v_31 = as_type<int>((as_type<uint>(v_30) + as_type<uint>(tint_f32_to_i32(mat2x2_f32[0][0u]))));
+  int const v_32 = as_type<int>((as_type<uint>(v_31) + as_type<uint>(tint_f32_to_i32(mat2x3_f32[0][0u]))));
+  int const v_33 = as_type<int>((as_type<uint>(v_32) + as_type<uint>(tint_f32_to_i32(mat2x4_f32[0][0u]))));
+  int const v_34 = as_type<int>((as_type<uint>(v_33) + as_type<uint>(tint_f32_to_i32(mat3x2_f32[0][0u]))));
+  int const v_35 = as_type<int>((as_type<uint>(v_34) + as_type<uint>(tint_f32_to_i32(mat3x3_f32[0][0u]))));
+  int const v_36 = as_type<int>((as_type<uint>(v_35) + as_type<uint>(tint_f32_to_i32(mat3x4_f32[0][0u]))));
+  int const v_37 = as_type<int>((as_type<uint>(v_36) + as_type<uint>(tint_f32_to_i32(mat4x2_f32[0][0u]))));
+  int const v_38 = as_type<int>((as_type<uint>(v_37) + as_type<uint>(tint_f32_to_i32(mat4x3_f32[0][0u]))));
+  int const v_39 = as_type<int>((as_type<uint>(v_38) + as_type<uint>(tint_f32_to_i32(mat4x4_f32[0][0u]))));
+  int const v_40 = as_type<int>((as_type<uint>(v_39) + as_type<uint>(tint_f16_to_i32(mat2x2_f16[0][0u]))));
+  int const v_41 = as_type<int>((as_type<uint>(v_40) + as_type<uint>(tint_f16_to_i32(mat2x3_f16[0][0u]))));
+  int const v_42 = as_type<int>((as_type<uint>(v_41) + as_type<uint>(tint_f16_to_i32(mat2x4_f16[0][0u]))));
+  int const v_43 = as_type<int>((as_type<uint>(v_42) + as_type<uint>(tint_f16_to_i32(mat3x2_f16[0][0u]))));
+  int const v_44 = as_type<int>((as_type<uint>(v_43) + as_type<uint>(tint_f16_to_i32(mat3x3_f16[0][0u]))));
+  int const v_45 = as_type<int>((as_type<uint>(v_44) + as_type<uint>(tint_f16_to_i32(mat3x4_f16[0][0u]))));
+  int const v_46 = as_type<int>((as_type<uint>(v_45) + as_type<uint>(tint_f16_to_i32(mat4x2_f16[0][0u]))));
+  int const v_47 = as_type<int>((as_type<uint>(v_46) + as_type<uint>(tint_f16_to_i32(mat4x3_f16[0][0u]))));
+  int const v_48 = as_type<int>((as_type<uint>(v_47) + as_type<uint>(tint_f16_to_i32(mat4x4_f16[0][0u]))));
+  int const v_49 = as_type<int>((as_type<uint>(v_48) + as_type<uint>(tint_f32_to_i32(arr2_vec3_f32[0][0u]))));
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_49) + as_type<uint>(tint_f16_to_i32(arr2_mat4x2_f16[0][0][0u]))))) + as_type<uint>(struct_inner.scalar_i32)))) + as_type<uint>(array_struct_inner[0].scalar_i32)));
 }

diff --git a/test/tint/buffer/uniform/dynamic_index/read.wgsl.expected.ir.msl b/test/tint/buffer/uniform/dynamic_index/read.wgsl.expected.ir.msl
index 94a3bab..092f6b3 100644
--- a/test/tint/buffer/uniform/dynamic_index/read.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/dynamic_index/read.wgsl.expected.ir.msl

@@ -99,24 +99,24 @@
   float4x3 const mat4x3_f32 = float4x3(v_7, v_8, v_9, float3(v_6[3u].packed));
   float4x4 const mat4x4_f32 = (*tint_module_vars.ub).arr[idx].mat4x4_f32;
   tint_array<float3, 2> const arr2_vec3_f32 = tint_load_array_packed_vec3((&(*tint_module_vars.ub).arr[idx].arr2_vec3_f32));
-  int const v_10 = (tint_f32_to_i32(scalar_f32) + scalar_i32);
-  int const v_11 = (v_10 + int(scalar_u32));
-  int const v_12 = ((v_11 + tint_f32_to_i32(vec2_f32[0u])) + vec2_i32[0u]);
-  int const v_13 = (v_12 + int(vec2_u32[0u]));
-  int const v_14 = ((v_13 + tint_f32_to_i32(vec3_f32[1u])) + vec3_i32[1u]);
-  int const v_15 = (v_14 + int(vec3_u32[1u]));
-  int const v_16 = ((v_15 + tint_f32_to_i32(vec4_f32[2u])) + vec4_i32[2u]);
-  int const v_17 = (v_16 + int(vec4_u32[2u]));
-  int const v_18 = (v_17 + tint_f32_to_i32(mat2x2_f32[0][0u]));
-  int const v_19 = (v_18 + tint_f32_to_i32(mat2x3_f32[0][0u]));
-  int const v_20 = (v_19 + tint_f32_to_i32(mat2x4_f32[0][0u]));
-  int const v_21 = (v_20 + tint_f32_to_i32(mat3x2_f32[0][0u]));
-  int const v_22 = (v_21 + tint_f32_to_i32(mat3x3_f32[0][0u]));
-  int const v_23 = (v_22 + tint_f32_to_i32(mat3x4_f32[0][0u]));
-  int const v_24 = (v_23 + tint_f32_to_i32(mat4x2_f32[0][0u]));
-  int const v_25 = (v_24 + tint_f32_to_i32(mat4x3_f32[0][0u]));
-  int const v_26 = (v_25 + tint_f32_to_i32(mat4x4_f32[0][0u]));
-  (*tint_module_vars.s) = (v_26 + tint_f32_to_i32(arr2_vec3_f32[0][0u]));
+  int const v_10 = as_type<int>((as_type<uint>(tint_f32_to_i32(scalar_f32)) + as_type<uint>(scalar_i32)));
+  int const v_11 = as_type<int>((as_type<uint>(v_10) + as_type<uint>(int(scalar_u32))));
+  int const v_12 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_11) + as_type<uint>(tint_f32_to_i32(vec2_f32[0u]))))) + as_type<uint>(vec2_i32[0u])));
+  int const v_13 = as_type<int>((as_type<uint>(v_12) + as_type<uint>(int(vec2_u32[0u]))));
+  int const v_14 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_13) + as_type<uint>(tint_f32_to_i32(vec3_f32[1u]))))) + as_type<uint>(vec3_i32[1u])));
+  int const v_15 = as_type<int>((as_type<uint>(v_14) + as_type<uint>(int(vec3_u32[1u]))));
+  int const v_16 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_15) + as_type<uint>(tint_f32_to_i32(vec4_f32[2u]))))) + as_type<uint>(vec4_i32[2u])));
+  int const v_17 = as_type<int>((as_type<uint>(v_16) + as_type<uint>(int(vec4_u32[2u]))));
+  int const v_18 = as_type<int>((as_type<uint>(v_17) + as_type<uint>(tint_f32_to_i32(mat2x2_f32[0][0u]))));
+  int const v_19 = as_type<int>((as_type<uint>(v_18) + as_type<uint>(tint_f32_to_i32(mat2x3_f32[0][0u]))));
+  int const v_20 = as_type<int>((as_type<uint>(v_19) + as_type<uint>(tint_f32_to_i32(mat2x4_f32[0][0u]))));
+  int const v_21 = as_type<int>((as_type<uint>(v_20) + as_type<uint>(tint_f32_to_i32(mat3x2_f32[0][0u]))));
+  int const v_22 = as_type<int>((as_type<uint>(v_21) + as_type<uint>(tint_f32_to_i32(mat3x3_f32[0][0u]))));
+  int const v_23 = as_type<int>((as_type<uint>(v_22) + as_type<uint>(tint_f32_to_i32(mat3x4_f32[0][0u]))));
+  int const v_24 = as_type<int>((as_type<uint>(v_23) + as_type<uint>(tint_f32_to_i32(mat4x2_f32[0][0u]))));
+  int const v_25 = as_type<int>((as_type<uint>(v_24) + as_type<uint>(tint_f32_to_i32(mat4x3_f32[0][0u]))));
+  int const v_26 = as_type<int>((as_type<uint>(v_25) + as_type<uint>(tint_f32_to_i32(mat4x4_f32[0][0u]))));
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(v_26) + as_type<uint>(tint_f32_to_i32(arr2_vec3_f32[0][0u]))));
 }
 
 kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], const constant S_packed_vec3* ub [[buffer(0)]], device int* s [[buffer(1)]]) {

diff --git a/test/tint/buffer/uniform/dynamic_index/read_f16.wgsl.expected.ir.msl b/test/tint/buffer/uniform/dynamic_index/read_f16.wgsl.expected.ir.msl
index d8794aa..0245d85 100644
--- a/test/tint/buffer/uniform/dynamic_index/read_f16.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/dynamic_index/read_f16.wgsl.expected.ir.msl

@@ -149,38 +149,38 @@
   half4x4 const mat4x4_f16 = (*tint_module_vars.ub).arr[idx].mat4x4_f16;
   tint_array<float3, 2> const arr2_vec3_f32 = tint_load_array_packed_vec3((&(*tint_module_vars.ub).arr[idx].arr2_vec3_f32));
   tint_array<half4x2, 2> const arr2_mat4x2_f16 = (*tint_module_vars.ub).arr[idx].arr2_mat4x2_f16;
-  int const v_19 = (tint_f32_to_i32(scalar_f32) + scalar_i32);
-  int const v_20 = (v_19 + int(scalar_u32));
-  int const v_21 = (v_20 + tint_f16_to_i32(scalar_f16));
-  int const v_22 = ((v_21 + tint_f32_to_i32(vec2_f32[0u])) + vec2_i32[0u]);
-  int const v_23 = (v_22 + int(vec2_u32[0u]));
-  int const v_24 = (v_23 + tint_f16_to_i32(vec2_f16[0u]));
-  int const v_25 = ((v_24 + tint_f32_to_i32(vec3_f32[1u])) + vec3_i32[1u]);
-  int const v_26 = (v_25 + int(vec3_u32[1u]));
-  int const v_27 = (v_26 + tint_f16_to_i32(vec3_f16[1u]));
-  int const v_28 = ((v_27 + tint_f32_to_i32(vec4_f32[2u])) + vec4_i32[2u]);
-  int const v_29 = (v_28 + int(vec4_u32[2u]));
-  int const v_30 = (v_29 + tint_f16_to_i32(vec4_f16[2u]));
-  int const v_31 = (v_30 + tint_f32_to_i32(mat2x2_f32[0][0u]));
-  int const v_32 = (v_31 + tint_f32_to_i32(mat2x3_f32[0][0u]));
-  int const v_33 = (v_32 + tint_f32_to_i32(mat2x4_f32[0][0u]));
-  int const v_34 = (v_33 + tint_f32_to_i32(mat3x2_f32[0][0u]));
-  int const v_35 = (v_34 + tint_f32_to_i32(mat3x3_f32[0][0u]));
-  int const v_36 = (v_35 + tint_f32_to_i32(mat3x4_f32[0][0u]));
-  int const v_37 = (v_36 + tint_f32_to_i32(mat4x2_f32[0][0u]));
-  int const v_38 = (v_37 + tint_f32_to_i32(mat4x3_f32[0][0u]));
-  int const v_39 = (v_38 + tint_f32_to_i32(mat4x4_f32[0][0u]));
-  int const v_40 = (v_39 + tint_f16_to_i32(mat2x2_f16[0][0u]));
-  int const v_41 = (v_40 + tint_f16_to_i32(mat2x3_f16[0][0u]));
-  int const v_42 = (v_41 + tint_f16_to_i32(mat2x4_f16[0][0u]));
-  int const v_43 = (v_42 + tint_f16_to_i32(mat3x2_f16[0][0u]));
-  int const v_44 = (v_43 + tint_f16_to_i32(mat3x3_f16[0][0u]));
-  int const v_45 = (v_44 + tint_f16_to_i32(mat3x4_f16[0][0u]));
-  int const v_46 = (v_45 + tint_f16_to_i32(mat4x2_f16[0][0u]));
-  int const v_47 = (v_46 + tint_f16_to_i32(mat4x3_f16[0][0u]));
-  int const v_48 = (v_47 + tint_f16_to_i32(mat4x4_f16[0][0u]));
-  int const v_49 = (v_48 + tint_f32_to_i32(arr2_vec3_f32[0][0u]));
-  (*tint_module_vars.s) = (v_49 + tint_f16_to_i32(arr2_mat4x2_f16[0][0][0u]));
+  int const v_19 = as_type<int>((as_type<uint>(tint_f32_to_i32(scalar_f32)) + as_type<uint>(scalar_i32)));
+  int const v_20 = as_type<int>((as_type<uint>(v_19) + as_type<uint>(int(scalar_u32))));
+  int const v_21 = as_type<int>((as_type<uint>(v_20) + as_type<uint>(tint_f16_to_i32(scalar_f16))));
+  int const v_22 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_21) + as_type<uint>(tint_f32_to_i32(vec2_f32[0u]))))) + as_type<uint>(vec2_i32[0u])));
+  int const v_23 = as_type<int>((as_type<uint>(v_22) + as_type<uint>(int(vec2_u32[0u]))));
+  int const v_24 = as_type<int>((as_type<uint>(v_23) + as_type<uint>(tint_f16_to_i32(vec2_f16[0u]))));
+  int const v_25 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_24) + as_type<uint>(tint_f32_to_i32(vec3_f32[1u]))))) + as_type<uint>(vec3_i32[1u])));
+  int const v_26 = as_type<int>((as_type<uint>(v_25) + as_type<uint>(int(vec3_u32[1u]))));
+  int const v_27 = as_type<int>((as_type<uint>(v_26) + as_type<uint>(tint_f16_to_i32(vec3_f16[1u]))));
+  int const v_28 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_27) + as_type<uint>(tint_f32_to_i32(vec4_f32[2u]))))) + as_type<uint>(vec4_i32[2u])));
+  int const v_29 = as_type<int>((as_type<uint>(v_28) + as_type<uint>(int(vec4_u32[2u]))));
+  int const v_30 = as_type<int>((as_type<uint>(v_29) + as_type<uint>(tint_f16_to_i32(vec4_f16[2u]))));
+  int const v_31 = as_type<int>((as_type<uint>(v_30) + as_type<uint>(tint_f32_to_i32(mat2x2_f32[0][0u]))));
+  int const v_32 = as_type<int>((as_type<uint>(v_31) + as_type<uint>(tint_f32_to_i32(mat2x3_f32[0][0u]))));
+  int const v_33 = as_type<int>((as_type<uint>(v_32) + as_type<uint>(tint_f32_to_i32(mat2x4_f32[0][0u]))));
+  int const v_34 = as_type<int>((as_type<uint>(v_33) + as_type<uint>(tint_f32_to_i32(mat3x2_f32[0][0u]))));
+  int const v_35 = as_type<int>((as_type<uint>(v_34) + as_type<uint>(tint_f32_to_i32(mat3x3_f32[0][0u]))));
+  int const v_36 = as_type<int>((as_type<uint>(v_35) + as_type<uint>(tint_f32_to_i32(mat3x4_f32[0][0u]))));
+  int const v_37 = as_type<int>((as_type<uint>(v_36) + as_type<uint>(tint_f32_to_i32(mat4x2_f32[0][0u]))));
+  int const v_38 = as_type<int>((as_type<uint>(v_37) + as_type<uint>(tint_f32_to_i32(mat4x3_f32[0][0u]))));
+  int const v_39 = as_type<int>((as_type<uint>(v_38) + as_type<uint>(tint_f32_to_i32(mat4x4_f32[0][0u]))));
+  int const v_40 = as_type<int>((as_type<uint>(v_39) + as_type<uint>(tint_f16_to_i32(mat2x2_f16[0][0u]))));
+  int const v_41 = as_type<int>((as_type<uint>(v_40) + as_type<uint>(tint_f16_to_i32(mat2x3_f16[0][0u]))));
+  int const v_42 = as_type<int>((as_type<uint>(v_41) + as_type<uint>(tint_f16_to_i32(mat2x4_f16[0][0u]))));
+  int const v_43 = as_type<int>((as_type<uint>(v_42) + as_type<uint>(tint_f16_to_i32(mat3x2_f16[0][0u]))));
+  int const v_44 = as_type<int>((as_type<uint>(v_43) + as_type<uint>(tint_f16_to_i32(mat3x3_f16[0][0u]))));
+  int const v_45 = as_type<int>((as_type<uint>(v_44) + as_type<uint>(tint_f16_to_i32(mat3x4_f16[0][0u]))));
+  int const v_46 = as_type<int>((as_type<uint>(v_45) + as_type<uint>(tint_f16_to_i32(mat4x2_f16[0][0u]))));
+  int const v_47 = as_type<int>((as_type<uint>(v_46) + as_type<uint>(tint_f16_to_i32(mat4x3_f16[0][0u]))));
+  int const v_48 = as_type<int>((as_type<uint>(v_47) + as_type<uint>(tint_f16_to_i32(mat4x4_f16[0][0u]))));
+  int const v_49 = as_type<int>((as_type<uint>(v_48) + as_type<uint>(tint_f32_to_i32(arr2_vec3_f32[0][0u]))));
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(v_49) + as_type<uint>(tint_f16_to_i32(arr2_mat4x2_f16[0][0][0u]))));
 }
 
 kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], const constant S_packed_vec3* ub [[buffer(0)]], device int* s [[buffer(1)]]) {

diff --git a/test/tint/buffer/uniform/static_index/read.wgsl.expected.ir.msl b/test/tint/buffer/uniform/static_index/read.wgsl.expected.ir.msl
index dc396ed..a4eb1f7 100644
--- a/test/tint/buffer/uniform/static_index/read.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/static_index/read.wgsl.expected.ir.msl

@@ -107,22 +107,22 @@
   tint_array<float3, 2> const arr2_vec3_f32 = tint_load_array_packed_vec3((&(*tint_module_vars.ub).arr2_vec3_f32));
   Inner const struct_inner = (*tint_module_vars.ub).struct_inner;
   tint_array<Inner, 4> const array_struct_inner = (*tint_module_vars.ub).array_struct_inner;
-  int const v_10 = (tint_f32_to_i32(scalar_f32) + scalar_i32);
-  int const v_11 = (v_10 + int(scalar_u32));
-  int const v_12 = ((v_11 + tint_f32_to_i32(vec2_f32[0u])) + vec2_i32[0u]);
-  int const v_13 = (v_12 + int(vec2_u32[0u]));
-  int const v_14 = ((v_13 + tint_f32_to_i32(vec3_f32[1u])) + vec3_i32[1u]);
-  int const v_15 = (v_14 + int(vec3_u32[1u]));
-  int const v_16 = ((v_15 + tint_f32_to_i32(vec4_f32[2u])) + vec4_i32[2u]);
-  int const v_17 = (v_16 + int(vec4_u32[2u]));
-  int const v_18 = (v_17 + tint_f32_to_i32(mat2x2_f32[0][0u]));
-  int const v_19 = (v_18 + tint_f32_to_i32(mat2x3_f32[0][0u]));
-  int const v_20 = (v_19 + tint_f32_to_i32(mat2x4_f32[0][0u]));
-  int const v_21 = (v_20 + tint_f32_to_i32(mat3x2_f32[0][0u]));
-  int const v_22 = (v_21 + tint_f32_to_i32(mat3x3_f32[0][0u]));
-  int const v_23 = (v_22 + tint_f32_to_i32(mat3x4_f32[0][0u]));
-  int const v_24 = (v_23 + tint_f32_to_i32(mat4x2_f32[0][0u]));
-  int const v_25 = (v_24 + tint_f32_to_i32(mat4x3_f32[0][0u]));
-  int const v_26 = (v_25 + tint_f32_to_i32(mat4x4_f32[0][0u]));
-  (*tint_module_vars.s) = (((v_26 + tint_f32_to_i32(arr2_vec3_f32[0][0u])) + struct_inner.scalar_i32) + array_struct_inner[0].scalar_i32);
+  int const v_10 = as_type<int>((as_type<uint>(tint_f32_to_i32(scalar_f32)) + as_type<uint>(scalar_i32)));
+  int const v_11 = as_type<int>((as_type<uint>(v_10) + as_type<uint>(int(scalar_u32))));
+  int const v_12 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_11) + as_type<uint>(tint_f32_to_i32(vec2_f32[0u]))))) + as_type<uint>(vec2_i32[0u])));
+  int const v_13 = as_type<int>((as_type<uint>(v_12) + as_type<uint>(int(vec2_u32[0u]))));
+  int const v_14 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_13) + as_type<uint>(tint_f32_to_i32(vec3_f32[1u]))))) + as_type<uint>(vec3_i32[1u])));
+  int const v_15 = as_type<int>((as_type<uint>(v_14) + as_type<uint>(int(vec3_u32[1u]))));
+  int const v_16 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_15) + as_type<uint>(tint_f32_to_i32(vec4_f32[2u]))))) + as_type<uint>(vec4_i32[2u])));
+  int const v_17 = as_type<int>((as_type<uint>(v_16) + as_type<uint>(int(vec4_u32[2u]))));
+  int const v_18 = as_type<int>((as_type<uint>(v_17) + as_type<uint>(tint_f32_to_i32(mat2x2_f32[0][0u]))));
+  int const v_19 = as_type<int>((as_type<uint>(v_18) + as_type<uint>(tint_f32_to_i32(mat2x3_f32[0][0u]))));
+  int const v_20 = as_type<int>((as_type<uint>(v_19) + as_type<uint>(tint_f32_to_i32(mat2x4_f32[0][0u]))));
+  int const v_21 = as_type<int>((as_type<uint>(v_20) + as_type<uint>(tint_f32_to_i32(mat3x2_f32[0][0u]))));
+  int const v_22 = as_type<int>((as_type<uint>(v_21) + as_type<uint>(tint_f32_to_i32(mat3x3_f32[0][0u]))));
+  int const v_23 = as_type<int>((as_type<uint>(v_22) + as_type<uint>(tint_f32_to_i32(mat3x4_f32[0][0u]))));
+  int const v_24 = as_type<int>((as_type<uint>(v_23) + as_type<uint>(tint_f32_to_i32(mat4x2_f32[0][0u]))));
+  int const v_25 = as_type<int>((as_type<uint>(v_24) + as_type<uint>(tint_f32_to_i32(mat4x3_f32[0][0u]))));
+  int const v_26 = as_type<int>((as_type<uint>(v_25) + as_type<uint>(tint_f32_to_i32(mat4x4_f32[0][0u]))));
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_26) + as_type<uint>(tint_f32_to_i32(arr2_vec3_f32[0][0u]))))) + as_type<uint>(struct_inner.scalar_i32)))) + as_type<uint>(array_struct_inner[0].scalar_i32)));
 }

diff --git a/test/tint/buffer/uniform/static_index/read_f16.wgsl.expected.ir.msl b/test/tint/buffer/uniform/static_index/read_f16.wgsl.expected.ir.msl
index 9b13da0..739c2c6 100644
--- a/test/tint/buffer/uniform/static_index/read_f16.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/static_index/read_f16.wgsl.expected.ir.msl

@@ -157,36 +157,36 @@
   tint_array<half4x2, 2> const arr2_mat4x2_f16 = (*tint_module_vars.ub).arr2_mat4x2_f16;
   Inner const struct_inner = (*tint_module_vars.ub).struct_inner;
   tint_array<Inner, 4> const array_struct_inner = (*tint_module_vars.ub).array_struct_inner;
-  int const v_19 = (tint_f32_to_i32(scalar_f32) + scalar_i32);
-  int const v_20 = (v_19 + int(scalar_u32));
-  int const v_21 = (v_20 + tint_f16_to_i32(scalar_f16));
-  int const v_22 = ((v_21 + tint_f32_to_i32(vec2_f32[0u])) + vec2_i32[0u]);
-  int const v_23 = (v_22 + int(vec2_u32[0u]));
-  int const v_24 = (v_23 + tint_f16_to_i32(vec2_f16[0u]));
-  int const v_25 = ((v_24 + tint_f32_to_i32(vec3_f32[1u])) + vec3_i32[1u]);
-  int const v_26 = (v_25 + int(vec3_u32[1u]));
-  int const v_27 = (v_26 + tint_f16_to_i32(vec3_f16[1u]));
-  int const v_28 = ((v_27 + tint_f32_to_i32(vec4_f32[2u])) + vec4_i32[2u]);
-  int const v_29 = (v_28 + int(vec4_u32[2u]));
-  int const v_30 = (v_29 + tint_f16_to_i32(vec4_f16[2u]));
-  int const v_31 = (v_30 + tint_f32_to_i32(mat2x2_f32[0][0u]));
-  int const v_32 = (v_31 + tint_f32_to_i32(mat2x3_f32[0][0u]));
-  int const v_33 = (v_32 + tint_f32_to_i32(mat2x4_f32[0][0u]));
-  int const v_34 = (v_33 + tint_f32_to_i32(mat3x2_f32[0][0u]));
-  int const v_35 = (v_34 + tint_f32_to_i32(mat3x3_f32[0][0u]));
-  int const v_36 = (v_35 + tint_f32_to_i32(mat3x4_f32[0][0u]));
-  int const v_37 = (v_36 + tint_f32_to_i32(mat4x2_f32[0][0u]));
-  int const v_38 = (v_37 + tint_f32_to_i32(mat4x3_f32[0][0u]));
-  int const v_39 = (v_38 + tint_f32_to_i32(mat4x4_f32[0][0u]));
-  int const v_40 = (v_39 + tint_f16_to_i32(mat2x2_f16[0][0u]));
-  int const v_41 = (v_40 + tint_f16_to_i32(mat2x3_f16[0][0u]));
-  int const v_42 = (v_41 + tint_f16_to_i32(mat2x4_f16[0][0u]));
-  int const v_43 = (v_42 + tint_f16_to_i32(mat3x2_f16[0][0u]));
-  int const v_44 = (v_43 + tint_f16_to_i32(mat3x3_f16[0][0u]));
-  int const v_45 = (v_44 + tint_f16_to_i32(mat3x4_f16[0][0u]));
-  int const v_46 = (v_45 + tint_f16_to_i32(mat4x2_f16[0][0u]));
-  int const v_47 = (v_46 + tint_f16_to_i32(mat4x3_f16[0][0u]));
-  int const v_48 = (v_47 + tint_f16_to_i32(mat4x4_f16[0][0u]));
-  int const v_49 = (v_48 + tint_f32_to_i32(arr2_vec3_f32[0][0u]));
-  (*tint_module_vars.s) = (((v_49 + tint_f16_to_i32(arr2_mat4x2_f16[0][0][0u])) + struct_inner.scalar_i32) + array_struct_inner[0].scalar_i32);
+  int const v_19 = as_type<int>((as_type<uint>(tint_f32_to_i32(scalar_f32)) + as_type<uint>(scalar_i32)));
+  int const v_20 = as_type<int>((as_type<uint>(v_19) + as_type<uint>(int(scalar_u32))));
+  int const v_21 = as_type<int>((as_type<uint>(v_20) + as_type<uint>(tint_f16_to_i32(scalar_f16))));
+  int const v_22 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_21) + as_type<uint>(tint_f32_to_i32(vec2_f32[0u]))))) + as_type<uint>(vec2_i32[0u])));
+  int const v_23 = as_type<int>((as_type<uint>(v_22) + as_type<uint>(int(vec2_u32[0u]))));
+  int const v_24 = as_type<int>((as_type<uint>(v_23) + as_type<uint>(tint_f16_to_i32(vec2_f16[0u]))));
+  int const v_25 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_24) + as_type<uint>(tint_f32_to_i32(vec3_f32[1u]))))) + as_type<uint>(vec3_i32[1u])));
+  int const v_26 = as_type<int>((as_type<uint>(v_25) + as_type<uint>(int(vec3_u32[1u]))));
+  int const v_27 = as_type<int>((as_type<uint>(v_26) + as_type<uint>(tint_f16_to_i32(vec3_f16[1u]))));
+  int const v_28 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_27) + as_type<uint>(tint_f32_to_i32(vec4_f32[2u]))))) + as_type<uint>(vec4_i32[2u])));
+  int const v_29 = as_type<int>((as_type<uint>(v_28) + as_type<uint>(int(vec4_u32[2u]))));
+  int const v_30 = as_type<int>((as_type<uint>(v_29) + as_type<uint>(tint_f16_to_i32(vec4_f16[2u]))));
+  int const v_31 = as_type<int>((as_type<uint>(v_30) + as_type<uint>(tint_f32_to_i32(mat2x2_f32[0][0u]))));
+  int const v_32 = as_type<int>((as_type<uint>(v_31) + as_type<uint>(tint_f32_to_i32(mat2x3_f32[0][0u]))));
+  int const v_33 = as_type<int>((as_type<uint>(v_32) + as_type<uint>(tint_f32_to_i32(mat2x4_f32[0][0u]))));
+  int const v_34 = as_type<int>((as_type<uint>(v_33) + as_type<uint>(tint_f32_to_i32(mat3x2_f32[0][0u]))));
+  int const v_35 = as_type<int>((as_type<uint>(v_34) + as_type<uint>(tint_f32_to_i32(mat3x3_f32[0][0u]))));
+  int const v_36 = as_type<int>((as_type<uint>(v_35) + as_type<uint>(tint_f32_to_i32(mat3x4_f32[0][0u]))));
+  int const v_37 = as_type<int>((as_type<uint>(v_36) + as_type<uint>(tint_f32_to_i32(mat4x2_f32[0][0u]))));
+  int const v_38 = as_type<int>((as_type<uint>(v_37) + as_type<uint>(tint_f32_to_i32(mat4x3_f32[0][0u]))));
+  int const v_39 = as_type<int>((as_type<uint>(v_38) + as_type<uint>(tint_f32_to_i32(mat4x4_f32[0][0u]))));
+  int const v_40 = as_type<int>((as_type<uint>(v_39) + as_type<uint>(tint_f16_to_i32(mat2x2_f16[0][0u]))));
+  int const v_41 = as_type<int>((as_type<uint>(v_40) + as_type<uint>(tint_f16_to_i32(mat2x3_f16[0][0u]))));
+  int const v_42 = as_type<int>((as_type<uint>(v_41) + as_type<uint>(tint_f16_to_i32(mat2x4_f16[0][0u]))));
+  int const v_43 = as_type<int>((as_type<uint>(v_42) + as_type<uint>(tint_f16_to_i32(mat3x2_f16[0][0u]))));
+  int const v_44 = as_type<int>((as_type<uint>(v_43) + as_type<uint>(tint_f16_to_i32(mat3x3_f16[0][0u]))));
+  int const v_45 = as_type<int>((as_type<uint>(v_44) + as_type<uint>(tint_f16_to_i32(mat3x4_f16[0][0u]))));
+  int const v_46 = as_type<int>((as_type<uint>(v_45) + as_type<uint>(tint_f16_to_i32(mat4x2_f16[0][0u]))));
+  int const v_47 = as_type<int>((as_type<uint>(v_46) + as_type<uint>(tint_f16_to_i32(mat4x3_f16[0][0u]))));
+  int const v_48 = as_type<int>((as_type<uint>(v_47) + as_type<uint>(tint_f16_to_i32(mat4x4_f16[0][0u]))));
+  int const v_49 = as_type<int>((as_type<uint>(v_48) + as_type<uint>(tint_f32_to_i32(arr2_vec3_f32[0][0u]))));
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(v_49) + as_type<uint>(tint_f16_to_i32(arr2_mat4x2_f16[0][0][0u]))))) + as_type<uint>(struct_inner.scalar_i32)))) + as_type<uint>(array_struct_inner[0].scalar_i32)));
 }

diff --git a/test/tint/buffer/uniform/std140/array/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 69e2b39..bac1d35a 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -20,7 +20,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index df19f03..4995b32 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -25,7 +25,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/array/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 59d476c..ff6c9b7 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -25,7 +25,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index a274f01..1bf7983 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -20,7 +20,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/array/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 6da3ee8..d5b7298 100644
--- a/test/tint/buffer/uniform/std140/array/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -20,7 +20,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/array/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 052fd40..ed3e893 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -25,7 +25,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/array/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index e6ae94a..33ad6f1 100644
--- a/test/tint/buffer/uniform/std140/array/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -20,7 +20,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index c13fc35..3184a3b 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -20,7 +20,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/array/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index a655f95..1812c24 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -20,7 +20,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index f143f2f..b8d1a15 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -25,7 +25,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/array/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index dbbb2a9..f9e9bd6 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -25,7 +25,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index a654e59..944b525 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -20,7 +20,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/array/mat4x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/array/mat4x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 2e34e4f..7d63cb3 100644
--- a/test/tint/buffer/uniform/std140/array/mat4x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/array/mat4x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -20,7 +20,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index b70a649..e2265f9 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -28,7 +28,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 7c94e1f..9c04bdc 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -28,7 +28,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 74e55f2..7c9dd83 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -41,7 +41,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 5f72409..cc87658 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -41,7 +41,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index a15644d..665349c 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -28,7 +28,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 56775c4..8fd4f70 100644
--- a/test/tint/buffer/uniform/std140/struct/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -28,7 +28,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 6eb3d24..f228d84 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -28,7 +28,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 4f82715..66ab5d2 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -28,7 +28,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index a4e09f9..33bc7d5 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -41,7 +41,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 48c967d..6e5c531 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -41,7 +41,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index d3baa63..98ccff7 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -28,7 +28,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 27325bd..c72b84d 100644
--- a/test/tint/buffer/uniform/std140/struct/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -28,7 +28,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index c63d015..5cf83ad 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -28,7 +28,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index d29fc5c..216fe03 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -28,7 +28,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index efa0370..4c23dec 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -41,7 +41,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 387e13f..7adeaff 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -40,7 +40,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 3ce237e..6ef00d4 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -28,7 +28,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 434ba94..5ac5fca 100644
--- a/test/tint/buffer/uniform/std140/struct/mat4x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/struct/mat4x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -27,7 +27,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 53f4230..ceffef1 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/static_index_via_ptr.wgsl.expected.ir.msl
index 747490f..d112f60 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x2_f16/static_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 3c7f6c6..742667b 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.msl
index 030fe26..8330d3d 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x2_f32/static_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index c451754..0389270 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -24,7 +24,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/static_index_via_ptr.wgsl.expected.ir.msl
index 10f09b9..8bfaf43 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x3_f16/static_index_via_ptr.wgsl.expected.ir.msl

@@ -24,7 +24,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 29d973e..59519ad 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -24,7 +24,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.msl
index 7adeee3..61305e2 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x3_f32/static_index_via_ptr.wgsl.expected.ir.msl

@@ -24,7 +24,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index a60f37f..2afc726 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/static_index_via_ptr.wgsl.expected.ir.msl
index 09fb364..4ec5c4d 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x4_f16/static_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 5f57cb2..0d7e19b 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.msl
index 6ae7177..101fdb7 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat2x4_f32/static_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 2d73c16..ddb7409 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/static_index_via_ptr.wgsl.expected.ir.msl
index 8249b4e..de3c21c 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x2_f16/static_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 7ed1a49..d464543 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/static_index_via_ptr.wgsl.expected.ir.msl
index 6e7c979..9188e96 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x2_f32/static_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 9bc4253..b68c821 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -24,7 +24,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/static_index_via_ptr.wgsl.expected.ir.msl
index 3e6307e..430d5f8 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x3_f16/static_index_via_ptr.wgsl.expected.ir.msl

@@ -24,7 +24,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index fbfe824..3ab8e87 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -24,7 +24,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/static_index_via_ptr.wgsl.expected.ir.msl
index 6d66244..229bff8 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x3_f32/static_index_via_ptr.wgsl.expected.ir.msl

@@ -24,7 +24,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 12108b3..af67891 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/static_index_via_ptr.wgsl.expected.ir.msl
index a8ce52e..35d4780 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x4_f16/static_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 47a7264..55daaa0 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/static_index_via_ptr.wgsl.expected.ir.msl
index 5a01d59..59010b3 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat3x4_f32/static_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index d1ea899..c0c721a 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/static_index_via_ptr.wgsl.expected.ir.msl
index 69e2374..ad5a6cf 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x2_f16/static_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 725ecbc..be9a251 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/static_index_via_ptr.wgsl.expected.ir.msl
index 00a2e78..25e560c 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x2_f32/static_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index c40a478..9c9357f 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -24,7 +24,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/static_index_via_ptr.wgsl.expected.ir.msl
index ffcfbac..5f69499 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x3_f16/static_index_via_ptr.wgsl.expected.ir.msl

@@ -24,7 +24,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index c0b9594..254ad839 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -24,7 +24,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.msl
index 7173bca..1ef4dd9 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x3_f32/static_index_via_ptr.wgsl.expected.ir.msl

@@ -24,7 +24,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 7d172b3..7dc67e4 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/static_index_via_ptr.wgsl.expected.ir.msl
index 05c1206..6447028 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x4_f16/static_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
index 8eff1fd..b0c6478 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/dynamic_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/static_index_via_ptr.wgsl.expected.ir.msl b/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/static_index_via_ptr.wgsl.expected.ir.msl
index 7dc00bc..d487cc2 100644
--- a/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/static_index_via_ptr.wgsl.expected.ir.msl
+++ b/test/tint/buffer/uniform/std140/unnested/mat4x4_f32/static_index_via_ptr.wgsl.expected.ir.msl

@@ -7,7 +7,7 @@
 };
 
 int i(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 

diff --git a/test/tint/bug/chromium/344265982.wgsl.expected.ir.msl b/test/tint/bug/chromium/344265982.wgsl.expected.ir.msl
index 5fc61ca..66c2544 100644
--- a/test/tint/bug/chromium/344265982.wgsl.expected.ir.msl
+++ b/test/tint/bug/chromium/344265982.wgsl.expected.ir.msl

@@ -33,7 +33,7 @@
         case 1:
         {
           {
-            i = (i + 1);
+            i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
           }
           continue;
         }
@@ -44,7 +44,7 @@
         }
       }
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.ir.msl
index 31c8097..69b47a5 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.ir.msl

@@ -25,7 +25,7 @@
       (*tint_module_vars.v4u)[i] = 1u;
       (*tint_module_vars.v2b)[i] = true;
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }
@@ -47,7 +47,7 @@
       }
       foo(tint_module_vars);
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.ir.msl
index d2cc280..4eb9cf4 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.ir.msl

@@ -35,7 +35,7 @@
       }
       foo(tint_module_vars);
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.ir.msl
index 3efa51b..34a3cd7 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.ir.msl

@@ -38,7 +38,7 @@
       v3b[i] = true;
       v4b[i] = true;
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.ir.msl
index 3c97ae1..3484c4e 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.ir.msl

@@ -30,7 +30,7 @@
       v4u_2[i] = 1u;
       v2b_2[i] = true;
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.ir.msl b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.ir.msl
index d779d9c..70f3124 100644
--- a/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.ir.msl
+++ b/test/tint/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.ir.msl

@@ -30,7 +30,7 @@
       v2u[i] = 1u;
       v2b[i] = true;
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/bug/tint/1121.wgsl.expected.ir.msl b/test/tint/bug/tint/1121.wgsl.expected.ir.msl
index 53964d3..b0ca2dd 100644
--- a/test/tint/bug/tint/1121.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1121.wgsl.expected.ir.msl

@@ -103,11 +103,11 @@
           } else {
             break;
           }
-          int2 tilePixel0Idx = int2((x * TILE_SIZE), (y * TILE_SIZE));
+          int2 tilePixel0Idx = int2(as_type<int>((as_type<uint>(x) * as_type<uint>(TILE_SIZE))), as_type<int>((as_type<uint>(y) * as_type<uint>(TILE_SIZE))));
           float2 const v_5 = (2.0f * float2(tilePixel0Idx));
           float2 floorCoord = ((v_5 / (*tint_module_vars.uniforms).fullScreenSize.xy) - float2(1.0f));
           int2 const v_6 = tilePixel0Idx;
-          float2 const v_7 = (2.0f * float2((v_6 + int2(TILE_SIZE))));
+          float2 const v_7 = (2.0f * float2(as_type<int2>((as_type<uint2>(v_6) + as_type<uint2>(int2(TILE_SIZE))))));
           float2 ceilCoord = ((v_7 / (*tint_module_vars.uniforms).fullScreenSize.xy) - float2(1.0f));
           float2 viewFloorCoord = float2((((-(viewNear) * floorCoord[0u]) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord[1u]) - (M[2][1] * viewNear)) / M[1][1]));
           float2 viewCeilCoord = float2((((-(viewNear) * ceilCoord[0u]) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * ceilCoord[1u]) - (M[2][1] * viewNear)) / M[1][1]));
@@ -149,7 +149,7 @@
             }
           }
           if ((dp >= 0.0f)) {
-            uint tileId = uint((x + (y * TILE_COUNT_X)));
+            uint tileId = uint(as_type<int>((as_type<uint>(x) + as_type<uint>(as_type<int>((as_type<uint>(y) * as_type<uint>(TILE_COUNT_X)))))));
             bool v_9 = false;
             if ((tileId < 0u)) {
               v_9 = true;
@@ -158,27 +158,27 @@
             }
             if (v_9) {
               {
-                x = (x + 1);
+                x = as_type<int>((as_type<uint>(x) + as_type<uint>(1)));
               }
               continue;
             }
             uint offset = atomic_fetch_add_explicit((&(*tint_module_vars.tileLightId).data[tileId].count), 1u, memory_order_relaxed);
             if ((offset >= (*tint_module_vars.config).numTileLightSlot)) {
               {
-                x = (x + 1);
+                x = as_type<int>((as_type<uint>(x) + as_type<uint>(1)));
               }
               continue;
             }
             (*tint_module_vars.tileLightId).data[tileId].lightId[offset] = GlobalInvocationID[0u];
           }
           {
-            x = (x + 1);
+            x = as_type<int>((as_type<uint>(x) + as_type<uint>(1)));
           }
           continue;
         }
       }
       {
-        y = (y + 1);
+        y = as_type<int>((as_type<uint>(y) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/bug/tint/1520.spvasm.expected.ir.msl b/test/tint/bug/tint/1520.spvasm.expected.ir.msl
index 9b6e79e..7103296 100644
--- a/test/tint/bug/tint/1520.spvasm.expected.ir.msl
+++ b/test/tint/bug/tint/1520.spvasm.expected.ir.msl

@@ -78,13 +78,13 @@
   ok = x_41;
   int4 const x_44 = int4(x_27);
   val = x_44;
-  int4 const x_47 = (x_44 + int4(1));
+  int4 const x_47 = as_type<int4>((as_type<uint4>(x_44) + as_type<uint4>(int4(1))));
   val = x_47;
-  int4 const x_48 = (x_47 - int4(1));
+  int4 const x_48 = as_type<int4>((as_type<uint4>(x_47) - as_type<uint4>(int4(1))));
   val = x_48;
-  int4 const x_49 = (x_48 + int4(1));
+  int4 const x_49 = as_type<int4>((as_type<uint4>(x_48) + as_type<uint4>(int4(1))));
   val = x_49;
-  int4 const x_50 = (x_49 - int4(1));
+  int4 const x_50 = as_type<int4>((as_type<uint4>(x_49) - as_type<uint4>(int4(1))));
   val = x_50;
   x_55 = false;
   if (x_41) {
@@ -92,11 +92,11 @@
     x_55 = x_54;
   }
   ok = x_55;
-  int4 const x_58 = (x_50 * int4(2));
+  int4 const x_58 = as_type<int4>((as_type<uint4>(x_50) * as_type<uint4>(int4(2))));
   val = x_58;
   int4 const x_59 = tint_div_v4i32(x_58, int4(2));
   val = x_59;
-  int4 const x_60 = (x_59 * int4(2));
+  int4 const x_60 = as_type<int4>((as_type<uint4>(x_59) * as_type<uint4>(int4(2))));
   val = x_60;
   int4 const x_61 = tint_div_v4i32(x_60, int4(2));
   val = x_61;

diff --git a/test/tint/bug/tint/1557.wgsl.expected.ir.msl b/test/tint/bug/tint/1557.wgsl.expected.ir.msl
index 81900a2..ad98ec0 100644
--- a/test/tint/bug/tint/1557.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1557.wgsl.expected.ir.msl

@@ -20,7 +20,7 @@
       if ((j >= 1)) {
         break;
       }
-      j = (j + 1);
+      j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)));
       int k = f();
       {
       }

diff --git a/test/tint/bug/tint/1605.wgsl.expected.ir.msl b/test/tint/bug/tint/1605.wgsl.expected.ir.msl
index f618569..11318e0f 100644
--- a/test/tint/bug/tint/1605.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1605.wgsl.expected.ir.msl

@@ -28,7 +28,7 @@
         }
       }
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/bug/tint/1664.wgsl.expected.ir.msl b/test/tint/bug/tint/1664.wgsl.expected.ir.msl
index a36872a..fee332e 100644
--- a/test/tint/bug/tint/1664.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1664.wgsl.expected.ir.msl

@@ -4,10 +4,10 @@
 kernel void f0() {
   int const a = 2147483647;
   int const b = 1;
-  int const c = (a + 1);
+  int const c = as_type<int>((as_type<uint>(a) + as_type<uint>(1)));
 }
 
 void f1() {
   int const a = 1;
-  int const b = ((-2147483647 - 1) - a);
+  int const b = as_type<int>((as_type<uint>((-2147483647 - 1)) - as_type<uint>(a)));
 }

diff --git a/test/tint/bug/tint/1677.wgsl.expected.ir.msl b/test/tint/bug/tint/1677.wgsl.expected.ir.msl
index 1b6f262..ac1a1ec 100644
--- a/test/tint/bug/tint/1677.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/1677.wgsl.expected.ir.msl

@@ -23,7 +23,7 @@
 };
 
 void tint_symbol_inner(uint3 id, tint_module_vars_struct tint_module_vars) {
-  int3 const pos = (int3((*tint_module_vars.input).position) - int3(0));
+  int3 const pos = as_type<int3>((as_type<uint3>(int3((*tint_module_vars.input).position)) - as_type<uint3>(int3(0))));
 }
 
 kernel void tint_symbol(uint3 id [[thread_position_in_grid]], const device Input_packed_vec3* input [[buffer(0)]]) {

diff --git a/test/tint/bug/tint/354627692.wgsl.expected.ir.msl b/test/tint/bug/tint/354627692.wgsl.expected.ir.msl
index cab996a..568f3b8 100644
--- a/test/tint/bug/tint/354627692.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/354627692.wgsl.expected.ir.msl

@@ -18,10 +18,10 @@
         {
           TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
             if ((i > 5)) {
-              i = (i * 2);
+              i = as_type<int>((as_type<uint>(i) * as_type<uint>(2)));
               break;
             } else {
-              i = (i * 2);
+              i = as_type<int>((as_type<uint>(i) * as_type<uint>(2)));
               break;
             }
             /* unreachable */

diff --git a/test/tint/bug/tint/948.wgsl.expected.ir.msl b/test/tint/bug/tint/948.wgsl.expected.ir.msl
index cdf8df29..5a500d2 100644
--- a/test/tint/bug/tint/948.wgsl.expected.ir.msl
+++ b/test/tint/bug/tint/948.wgsl.expected.ir.msl

@@ -236,7 +236,7 @@
       }
       {
         int const x_304 = i;
-        i = (x_304 + 1);
+        i = as_type<int>((as_type<uint>(x_304) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl
index 4f47486..efd369f 100644
--- a/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl
+++ b/test/tint/builtins/workgroupUniformLoad/for_loop.wgsl.expected.ir.msl

@@ -26,7 +26,7 @@
         threadgroup_barrier(mem_flags::mem_threadgroup);
         int const v_2 = (*tint_module_vars.b);
         threadgroup_barrier(mem_flags::mem_threadgroup);
-        i = (i + v_2);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(v_2)));
       }
       continue;
     }

diff --git a/test/tint/expressions/binary/add/scalar-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/add/scalar-scalar/i32.wgsl.expected.ir.msl
index 830879b..16f8aab 100644
--- a/test/tint/expressions/binary/add/scalar-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/add/scalar-scalar/i32.wgsl.expected.ir.msl

@@ -4,5 +4,5 @@
 kernel void f() {
   int const a = 1;
   int const b = 2;
-  int const r = (a + b);
+  int const r = as_type<int>((as_type<uint>(a) + as_type<uint>(b)));
 }

diff --git a/test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.ir.msl
index 3ae3f08..43d5947 100644
--- a/test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.ir.msl

@@ -4,5 +4,5 @@
 kernel void f() {
   int const a = 4;
   int3 const b = int3(1, 2, 3);
-  int3 const r = (a + b);
+  int3 const r = as_type<int3>((as_type<uint>(a) + as_type<uint3>(b)));
 }

diff --git a/test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.ir.msl
index cd7ba07..4044dab 100644
--- a/test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.ir.msl

@@ -4,5 +4,5 @@
 kernel void f() {
   int3 const a = int3(1, 2, 3);
   int const b = 4;
-  int3 const r = (a + b);
+  int3 const r = as_type<int3>((as_type<uint3>(a) + as_type<uint>(b)));
 }

diff --git a/test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.ir.msl
index 440380d..41a4b0c 100644
--- a/test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.ir.msl

@@ -4,5 +4,5 @@
 kernel void f() {
   int3 const a = int3(1, 2, 3);
   int3 const b = int3(4, 5, 6);
-  int3 const r = (a + b);
+  int3 const r = as_type<int3>((as_type<uint3>(a) + as_type<uint3>(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 80bf690..ebfd012 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

@@ -11,5 +11,5 @@
 kernel void f() {
   int a = 1;
   int b = 0;
-  int const r = tint_div_i32(a, (b + b));
+  int const r = tint_div_i32(a, as_type<int>((as_type<uint>(b) + as_type<uint>(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 07405f7..6579d98 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

@@ -11,6 +11,6 @@
 kernel void f() {
   int a = 4;
   int3 b = int3(0, 2, 0);
-  int3 const v_3 = (b + b);
+  int3 const v_3 = as_type<int3>((as_type<uint3>(b) + as_type<uint3>(b)));
   int3 const r = tint_div_v3i32(int3(a), v_3);
 }

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 875a9e2..13b980a 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

@@ -12,5 +12,5 @@
   int3 a = int3(1, 2, 3);
   int b = 0;
   int3 const v_3 = a;
-  int3 const r = tint_div_v3i32(v_3, int3((b + b)));
+  int3 const r = tint_div_v3i32(v_3, int3(as_type<int>((as_type<uint>(b) + as_type<uint>(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 c65386e..a499961 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

@@ -11,5 +11,5 @@
 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 const r = tint_div_v3i32(a, as_type<int3>((as_type<uint3>(b) + as_type<uint3>(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 d11d161..53bd8f6 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

@@ -6,7 +6,7 @@
   bool const v_1 = bool((v & uint((rhs == -1))));
   uint const v_2 = uint((rhs == 0));
   int const v_3 = select(rhs, 1, bool((v_2 | uint(v_1))));
-  return (lhs - ((lhs / v_3) * v_3));
+  return as_type<int>((as_type<uint>(lhs) - as_type<uint>(as_type<int>((as_type<uint>((lhs / v_3)) * as_type<uint>(v_3))))));
 }
 
 kernel void f() {

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 4eb5587..8f6eeb7 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

@@ -6,7 +6,7 @@
   bool3 const v_1 = bool3((v & uint3((rhs == int3(-1)))));
   uint3 const v_2 = uint3((rhs == int3(0)));
   int3 const v_3 = select(rhs, int3(1), bool3((v_2 | uint3(v_1))));
-  return (lhs - ((lhs / v_3) * v_3));
+  return as_type<int3>((as_type<uint3>(lhs) - as_type<uint3>(as_type<int3>((as_type<uint3>((lhs / v_3)) * as_type<uint3>(v_3))))));
 }
 
 kernel void f() {

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 aaa74db..f3d7dda 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

@@ -6,7 +6,7 @@
   bool3 const v_1 = bool3((v & uint3((rhs == int3(-1)))));
   uint3 const v_2 = uint3((rhs == int3(0)));
   int3 const v_3 = select(rhs, int3(1), bool3((v_2 | uint3(v_1))));
-  return (lhs - ((lhs / v_3) * v_3));
+  return as_type<int3>((as_type<uint3>(lhs) - as_type<uint3>(as_type<int3>((as_type<uint3>((lhs / v_3)) * as_type<uint3>(v_3))))));
 }
 
 kernel void f() {

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 453a6ae..7e204ad 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

@@ -6,7 +6,7 @@
   bool3 const v_1 = bool3((v & uint3((rhs == int3(-1)))));
   uint3 const v_2 = uint3((rhs == int3(0)));
   int3 const v_3 = select(rhs, int3(1), bool3((v_2 | uint3(v_1))));
-  return (lhs - ((lhs / v_3) * v_3));
+  return as_type<int3>((as_type<uint3>(lhs) - as_type<uint3>(as_type<int3>((as_type<uint3>((lhs / v_3)) * as_type<uint3>(v_3))))));
 }
 
 kernel void f() {

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 2efc410..be68c37 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

@@ -6,7 +6,7 @@
   bool const v_1 = bool((v & uint((rhs == -1))));
   uint const v_2 = uint((rhs == 0));
   int const v_3 = select(rhs, 1, bool((v_2 | uint(v_1))));
-  return (lhs - ((lhs / v_3) * v_3));
+  return as_type<int>((as_type<uint>(lhs) - as_type<uint>(as_type<int>((as_type<uint>((lhs / v_3)) * as_type<uint>(v_3))))));
 }
 
 kernel void f() {

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 48ea2d3..5ae267a 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

@@ -6,7 +6,7 @@
   bool3 const v_1 = bool3((v & uint3((rhs == int3(-1)))));
   uint3 const v_2 = uint3((rhs == int3(0)));
   int3 const v_3 = select(rhs, int3(1), bool3((v_2 | uint3(v_1))));
-  return (lhs - ((lhs / v_3) * v_3));
+  return as_type<int3>((as_type<uint3>(lhs) - as_type<uint3>(as_type<int3>((as_type<uint3>((lhs / v_3)) * as_type<uint3>(v_3))))));
 }
 
 kernel void f() {

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 7c73d60..a3cdf32 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

@@ -6,7 +6,7 @@
   bool3 const v_1 = bool3((v & uint3((rhs == int3(-1)))));
   uint3 const v_2 = uint3((rhs == int3(0)));
   int3 const v_3 = select(rhs, int3(1), bool3((v_2 | uint3(v_1))));
-  return (lhs - ((lhs / v_3) * v_3));
+  return as_type<int3>((as_type<uint3>(lhs) - as_type<uint3>(as_type<int3>((as_type<uint3>((lhs / v_3)) * as_type<uint3>(v_3))))));
 }
 
 kernel void f() {

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 4899742..2bb7167 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

@@ -6,7 +6,7 @@
   bool3 const v_1 = bool3((v & uint3((rhs == int3(-1)))));
   uint3 const v_2 = uint3((rhs == int3(0)));
   int3 const v_3 = select(rhs, int3(1), bool3((v_2 | uint3(v_1))));
-  return (lhs - ((lhs / v_3) * v_3));
+  return as_type<int3>((as_type<uint3>(lhs) - as_type<uint3>(as_type<int3>((as_type<uint3>((lhs / v_3)) * as_type<uint3>(v_3))))));
 }
 
 kernel void f() {

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 11f2288..028ae5b 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

@@ -6,11 +6,11 @@
   bool const v_1 = bool((v & uint((rhs == -1))));
   uint const v_2 = uint((rhs == 0));
   int const v_3 = select(rhs, 1, bool((v_2 | uint(v_1))));
-  return (lhs - ((lhs / v_3) * v_3));
+  return as_type<int>((as_type<uint>(lhs) - as_type<uint>(as_type<int>((as_type<uint>((lhs / v_3)) * as_type<uint>(v_3))))));
 }
 
 kernel void f() {
   int a = 1;
   int b = 0;
-  int const r = tint_mod_i32(a, (b + b));
+  int const r = tint_mod_i32(a, as_type<int>((as_type<uint>(b) + as_type<uint>(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 0f9428d..301922f 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

@@ -6,12 +6,12 @@
   bool3 const v_1 = bool3((v & uint3((rhs == int3(-1)))));
   uint3 const v_2 = uint3((rhs == int3(0)));
   int3 const v_3 = select(rhs, int3(1), bool3((v_2 | uint3(v_1))));
-  return (lhs - ((lhs / v_3) * v_3));
+  return as_type<int3>((as_type<uint3>(lhs) - as_type<uint3>(as_type<int3>((as_type<uint3>((lhs / v_3)) * as_type<uint3>(v_3))))));
 }
 
 kernel void f() {
   int a = 4;
   int3 b = int3(0, 2, 0);
-  int3 const v_4 = (b + b);
+  int3 const v_4 = as_type<int3>((as_type<uint3>(b) + as_type<uint3>(b)));
   int3 const r = tint_mod_v3i32(int3(a), v_4);
 }

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 5fb0c50..31b20e7 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

@@ -6,12 +6,12 @@
   bool3 const v_1 = bool3((v & uint3((rhs == int3(-1)))));
   uint3 const v_2 = uint3((rhs == int3(0)));
   int3 const v_3 = select(rhs, int3(1), bool3((v_2 | uint3(v_1))));
-  return (lhs - ((lhs / v_3) * v_3));
+  return as_type<int3>((as_type<uint3>(lhs) - as_type<uint3>(as_type<int3>((as_type<uint3>((lhs / v_3)) * as_type<uint3>(v_3))))));
 }
 
 kernel void f() {
   int3 a = int3(1, 2, 3);
   int b = 0;
   int3 const v_4 = a;
-  int3 const r = tint_mod_v3i32(v_4, int3((b + b)));
+  int3 const r = tint_mod_v3i32(v_4, int3(as_type<int>((as_type<uint>(b) + as_type<uint>(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 d989a1b..380216c 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

@@ -6,11 +6,11 @@
   bool3 const v_1 = bool3((v & uint3((rhs == int3(-1)))));
   uint3 const v_2 = uint3((rhs == int3(0)));
   int3 const v_3 = select(rhs, int3(1), bool3((v_2 | uint3(v_1))));
-  return (lhs - ((lhs / v_3) * v_3));
+  return as_type<int3>((as_type<uint3>(lhs) - as_type<uint3>(as_type<int3>((as_type<uint3>((lhs / v_3)) * as_type<uint3>(v_3))))));
 }
 
 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 const r = tint_mod_v3i32(a, as_type<int3>((as_type<uint3>(b) + as_type<uint3>(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 a655b85..678c780 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

@@ -6,7 +6,7 @@
   bool const v_1 = bool((v & uint((rhs == -1))));
   uint const v_2 = uint((rhs == 0));
   int const v_3 = select(rhs, 1, bool((v_2 | uint(v_1))));
-  return (lhs - ((lhs / v_3) * v_3));
+  return as_type<int>((as_type<uint>(lhs) - as_type<uint>(as_type<int>((as_type<uint>((lhs / v_3)) * as_type<uint>(v_3))))));
 }
 
 kernel void f() {

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 c17dd02..8839cd2 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

@@ -6,7 +6,7 @@
   bool3 const v_1 = bool3((v & uint3((rhs == int3(-1)))));
   uint3 const v_2 = uint3((rhs == int3(0)));
   int3 const v_3 = select(rhs, int3(1), bool3((v_2 | uint3(v_1))));
-  return (lhs - ((lhs / v_3) * v_3));
+  return as_type<int3>((as_type<uint3>(lhs) - as_type<uint3>(as_type<int3>((as_type<uint3>((lhs / v_3)) * as_type<uint3>(v_3))))));
 }
 
 kernel void f() {

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 c3291c2..b290081 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

@@ -6,7 +6,7 @@
   bool3 const v_1 = bool3((v & uint3((rhs == int3(-1)))));
   uint3 const v_2 = uint3((rhs == int3(0)));
   int3 const v_3 = select(rhs, int3(1), bool3((v_2 | uint3(v_1))));
-  return (lhs - ((lhs / v_3) * v_3));
+  return as_type<int3>((as_type<uint3>(lhs) - as_type<uint3>(as_type<int3>((as_type<uint3>((lhs / v_3)) * as_type<uint3>(v_3))))));
 }
 
 kernel void f() {

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 f36c4fa..7445159 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

@@ -6,7 +6,7 @@
   bool3 const v_1 = bool3((v & uint3((rhs == int3(-1)))));
   uint3 const v_2 = uint3((rhs == int3(0)));
   int3 const v_3 = select(rhs, int3(1), bool3((v_2 | uint3(v_1))));
-  return (lhs - ((lhs / v_3) * v_3));
+  return as_type<int3>((as_type<uint3>(lhs) - as_type<uint3>(as_type<int3>((as_type<uint3>((lhs / v_3)) * as_type<uint3>(v_3))))));
 }
 
 kernel void f() {

diff --git a/test/tint/expressions/binary/mul/scalar-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mul/scalar-scalar/i32.wgsl.expected.ir.msl
index f6dfd1a..24f928e 100644
--- a/test/tint/expressions/binary/mul/scalar-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mul/scalar-scalar/i32.wgsl.expected.ir.msl

@@ -4,5 +4,5 @@
 kernel void f() {
   int const a = 1;
   int const b = 2;
-  int const r = (a * b);
+  int const r = as_type<int>((as_type<uint>(a) * as_type<uint>(b)));
 }

diff --git a/test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.ir.msl
index 4bf0ea5..cf4ad8f 100644
--- a/test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.ir.msl

@@ -4,5 +4,5 @@
 kernel void f() {
   int const a = 4;
   int3 const b = int3(1, 2, 3);
-  int3 const r = (a * b);
+  int3 const r = as_type<int3>((as_type<uint>(a) * as_type<uint3>(b)));
 }

diff --git a/test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.ir.msl
index 6bd1175..d846463 100644
--- a/test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.ir.msl

@@ -4,5 +4,5 @@
 kernel void f() {
   int3 const a = int3(1, 2, 3);
   int const b = 4;
-  int3 const r = (a * b);
+  int3 const r = as_type<int3>((as_type<uint3>(a) * as_type<uint>(b)));
 }

diff --git a/test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.ir.msl
index 3d0cf97..cfecfc2 100644
--- a/test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.ir.msl

@@ -4,5 +4,5 @@
 kernel void f() {
   int3 const a = int3(1, 2, 3);
   int3 const b = int3(4, 5, 6);
-  int3 const r = (a * b);
+  int3 const r = as_type<int3>((as_type<uint3>(a) * as_type<uint3>(b)));
 }

diff --git a/test/tint/expressions/binary/sub/scalar-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/sub/scalar-scalar/i32.wgsl.expected.ir.msl
index af343ac..1a68402 100644
--- a/test/tint/expressions/binary/sub/scalar-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/sub/scalar-scalar/i32.wgsl.expected.ir.msl

@@ -4,5 +4,5 @@
 kernel void f() {
   int const a = 1;
   int const b = 2;
-  int const r = (a - b);
+  int const r = as_type<int>((as_type<uint>(a) - as_type<uint>(b)));
 }

diff --git a/test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.ir.msl
index ccea36e..5111ae2 100644
--- a/test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.ir.msl

@@ -4,5 +4,5 @@
 kernel void f() {
   int const a = 4;
   int3 const b = int3(1, 2, 3);
-  int3 const r = (a - b);
+  int3 const r = as_type<int3>((as_type<uint>(a) - as_type<uint3>(b)));
 }

diff --git a/test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.ir.msl
index 0ff7cc3..8061d78 100644
--- a/test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.ir.msl

@@ -4,5 +4,5 @@
 kernel void f() {
   int3 const a = int3(1, 2, 3);
   int const b = 4;
-  int3 const r = (a - b);
+  int3 const r = as_type<int3>((as_type<uint3>(a) - as_type<uint>(b)));
 }

diff --git a/test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.ir.msl b/test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.ir.msl
index 142cb13..d3af694 100644
--- a/test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.ir.msl
+++ b/test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.ir.msl

@@ -4,5 +4,5 @@
 kernel void f() {
   int3 const a = int3(1, 2, 3);
   int3 const b = int3(4, 5, 6);
-  int3 const r = (a - b);
+  int3 const r = as_type<int3>((as_type<uint3>(a) - as_type<uint3>(b)));
 }

diff --git a/test/tint/expressions/literals/intmin.wgsl.expected.ir.msl b/test/tint/expressions/literals/intmin.wgsl.expected.ir.msl
index fc97589..9c65f58 100644
--- a/test/tint/expressions/literals/intmin.wgsl.expected.ir.msl
+++ b/test/tint/expressions/literals/intmin.wgsl.expected.ir.msl

@@ -3,7 +3,7 @@
 
 int add_int_min_explicit() {
   int a = (-2147483647 - 1);
-  int b = (a + 1);
+  int b = as_type<int>((as_type<uint>(a) + as_type<uint>(1)));
   int c = -2147483647;
   return c;
 }

diff --git a/test/tint/expressions/user_call/multi_param_no_return.wgsl.expected.ir.msl b/test/tint/expressions/user_call/multi_param_no_return.wgsl.expected.ir.msl
index 8206a3c..7412ab3 100644
--- a/test/tint/expressions/user_call/multi_param_no_return.wgsl.expected.ir.msl
+++ b/test/tint/expressions/user_call/multi_param_no_return.wgsl.expected.ir.msl

@@ -2,8 +2,8 @@
 using namespace metal;
 
 void c(int x, int y, int z) {
-  int a = (((1 + x) + y) + z);
-  a = (a + 2);
+  int a = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(1) + as_type<uint>(x)))) + as_type<uint>(y)))) + as_type<uint>(z)));
+  a = as_type<int>((as_type<uint>(a) + as_type<uint>(2)));
 }
 
 void b() {

diff --git a/test/tint/expressions/user_call/multi_param_return.wgsl.expected.ir.msl b/test/tint/expressions/user_call/multi_param_return.wgsl.expected.ir.msl
index 030b1ce..b27ebdf 100644
--- a/test/tint/expressions/user_call/multi_param_return.wgsl.expected.ir.msl
+++ b/test/tint/expressions/user_call/multi_param_return.wgsl.expected.ir.msl

@@ -2,13 +2,13 @@
 using namespace metal;
 
 int c(int x, int y, int z) {
-  int a = (((1 + x) + y) + z);
-  a = (a + 2);
+  int a = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(1) + as_type<uint>(x)))) + as_type<uint>(y)))) + as_type<uint>(z)));
+  a = as_type<int>((as_type<uint>(a) + as_type<uint>(2)));
   return a;
 }
 
 void b() {
   int b_1 = c(2, 3, 4);
   int const v = c(3, 4, 5);
-  b_1 = (b_1 + v);
+  b_1 = as_type<int>((as_type<uint>(b_1) + as_type<uint>(v)));
 }

diff --git a/test/tint/expressions/user_call/no_params_no_return.wgsl.expected.ir.msl b/test/tint/expressions/user_call/no_params_no_return.wgsl.expected.ir.msl
index addeb9e..7ef63d7 100644
--- a/test/tint/expressions/user_call/no_params_no_return.wgsl.expected.ir.msl
+++ b/test/tint/expressions/user_call/no_params_no_return.wgsl.expected.ir.msl

@@ -3,7 +3,7 @@
 
 void c() {
   int a = 1;
-  a = (a + 2);
+  a = as_type<int>((as_type<uint>(a) + as_type<uint>(2)));
 }
 
 void b() {

diff --git a/test/tint/expressions/user_call/no_params_return.wgsl.expected.ir.msl b/test/tint/expressions/user_call/no_params_return.wgsl.expected.ir.msl
index 19ec178..13b4090 100644
--- a/test/tint/expressions/user_call/no_params_return.wgsl.expected.ir.msl
+++ b/test/tint/expressions/user_call/no_params_return.wgsl.expected.ir.msl

@@ -3,12 +3,12 @@
 
 int c() {
   int a = 1;
-  a = (a + 2);
+  a = as_type<int>((as_type<uint>(a) + as_type<uint>(2)));
   return a;
 }
 
 void b() {
   int b_1 = c();
   int const v = c();
-  b_1 = (b_1 + v);
+  b_1 = as_type<int>((as_type<uint>(b_1) + as_type<uint>(v)));
 }

diff --git a/test/tint/expressions/user_call/one_param_no_return.wgsl.expected.ir.msl b/test/tint/expressions/user_call/one_param_no_return.wgsl.expected.ir.msl
index 07d7809..21fd720 100644
--- a/test/tint/expressions/user_call/one_param_no_return.wgsl.expected.ir.msl
+++ b/test/tint/expressions/user_call/one_param_no_return.wgsl.expected.ir.msl

@@ -2,8 +2,8 @@
 using namespace metal;
 
 void c(int z) {
-  int a = (1 + z);
-  a = (a + 2);
+  int a = as_type<int>((as_type<uint>(1) + as_type<uint>(z)));
+  a = as_type<int>((as_type<uint>(a) + as_type<uint>(2)));
 }
 
 void b() {

diff --git a/test/tint/expressions/user_call/one_param_return.wgsl.expected.ir.msl b/test/tint/expressions/user_call/one_param_return.wgsl.expected.ir.msl
index 891ee99..53d439a 100644
--- a/test/tint/expressions/user_call/one_param_return.wgsl.expected.ir.msl
+++ b/test/tint/expressions/user_call/one_param_return.wgsl.expected.ir.msl

@@ -2,13 +2,13 @@
 using namespace metal;
 
 int c(int z) {
-  int a = (1 + z);
-  a = (a + 2);
+  int a = as_type<int>((as_type<uint>(1) + as_type<uint>(z)));
+  a = as_type<int>((as_type<uint>(a) + as_type<uint>(2)));
   return a;
 }
 
 void b() {
   int b_1 = c(2);
   int const v = c(3);
-  b_1 = (b_1 + v);
+  b_1 = as_type<int>((as_type<uint>(b_1) + as_type<uint>(v)));
 }

diff --git a/test/tint/identifiers/underscore/double/alias.wgsl.expected.ir.msl b/test/tint/identifiers/underscore/double/alias.wgsl.expected.ir.msl
index 7dad3f8..3e5c9c4 100644
--- a/test/tint/identifiers/underscore/double/alias.wgsl.expected.ir.msl
+++ b/test/tint/identifiers/underscore/double/alias.wgsl.expected.ir.msl

@@ -9,5 +9,5 @@
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s};
   int c = 0;
   int d = 0;
-  (*tint_module_vars.s) = (c + d);
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(c) + as_type<uint>(d)));
 }

diff --git a/test/tint/identifiers/underscore/double/let.wgsl.expected.ir.msl b/test/tint/identifiers/underscore/double/let.wgsl.expected.ir.msl
index d6f2cb2..cb92137 100644
--- a/test/tint/identifiers/underscore/double/let.wgsl.expected.ir.msl
+++ b/test/tint/identifiers/underscore/double/let.wgsl.expected.ir.msl

@@ -11,5 +11,5 @@
   int const a__ = a;
   int const b = a;
   int const b__ = a__;
-  (*tint_module_vars.s) = (((a + a__) + b) + b__);
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) + as_type<uint>(a__)))) + as_type<uint>(b)))) + as_type<uint>(b__)));
 }

diff --git a/test/tint/identifiers/underscore/double/struct.wgsl.expected.ir.msl b/test/tint/identifiers/underscore/double/struct.wgsl.expected.ir.msl
index ef545b1..89392d1 100644
--- a/test/tint/identifiers/underscore/double/struct.wgsl.expected.ir.msl
+++ b/test/tint/identifiers/underscore/double/struct.wgsl.expected.ir.msl

@@ -13,5 +13,5 @@
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s};
   a__ const c = a__{};
   int const d = c.b__;
-  (*tint_module_vars.s) = (c.b__ + d);
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(c.b__) + as_type<uint>(d)));
 }

diff --git a/test/tint/identifiers/underscore/double/var.wgsl.expected.ir.msl b/test/tint/identifiers/underscore/double/var.wgsl.expected.ir.msl
index eaca8c6..91f5402 100644
--- a/test/tint/identifiers/underscore/double/var.wgsl.expected.ir.msl
+++ b/test/tint/identifiers/underscore/double/var.wgsl.expected.ir.msl

@@ -13,5 +13,5 @@
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .a=(&a), .a__=(&a__)};
   int b = (*tint_module_vars.a);
   int b__ = (*tint_module_vars.a__);
-  (*tint_module_vars.s) = (b + b__);
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(b) + as_type<uint>(b__)));
 }

diff --git a/test/tint/identifiers/underscore/prefix/lower/alias.wgsl.expected.ir.msl b/test/tint/identifiers/underscore/prefix/lower/alias.wgsl.expected.ir.msl
index 7dad3f8..3e5c9c4 100644
--- a/test/tint/identifiers/underscore/prefix/lower/alias.wgsl.expected.ir.msl
+++ b/test/tint/identifiers/underscore/prefix/lower/alias.wgsl.expected.ir.msl

@@ -9,5 +9,5 @@
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s};
   int c = 0;
   int d = 0;
-  (*tint_module_vars.s) = (c + d);
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(c) + as_type<uint>(d)));
 }

diff --git a/test/tint/identifiers/underscore/prefix/lower/let.wgsl.expected.ir.msl b/test/tint/identifiers/underscore/prefix/lower/let.wgsl.expected.ir.msl
index 56b1a9f..4df0bef 100644
--- a/test/tint/identifiers/underscore/prefix/lower/let.wgsl.expected.ir.msl
+++ b/test/tint/identifiers/underscore/prefix/lower/let.wgsl.expected.ir.msl

@@ -11,5 +11,5 @@
   int const _a = a;
   int const b = a;
   int const _b = _a;
-  (*tint_module_vars.s) = (((a + _a) + b) + _b);
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) + as_type<uint>(_a)))) + as_type<uint>(b)))) + as_type<uint>(_b)));
 }

diff --git a/test/tint/identifiers/underscore/prefix/lower/struct.wgsl.expected.ir.msl b/test/tint/identifiers/underscore/prefix/lower/struct.wgsl.expected.ir.msl
index 7abad97..4236eee 100644
--- a/test/tint/identifiers/underscore/prefix/lower/struct.wgsl.expected.ir.msl
+++ b/test/tint/identifiers/underscore/prefix/lower/struct.wgsl.expected.ir.msl

@@ -13,5 +13,5 @@
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s};
   _a const c = _a{};
   int const d = c._b;
-  (*tint_module_vars.s) = (c._b + d);
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(c._b) + as_type<uint>(d)));
 }

diff --git a/test/tint/identifiers/underscore/prefix/lower/var.wgsl.expected.ir.msl b/test/tint/identifiers/underscore/prefix/lower/var.wgsl.expected.ir.msl
index 2dbd2d7..79b931d 100644
--- a/test/tint/identifiers/underscore/prefix/lower/var.wgsl.expected.ir.msl
+++ b/test/tint/identifiers/underscore/prefix/lower/var.wgsl.expected.ir.msl

@@ -13,5 +13,5 @@
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .a=(&a), ._a=(&_a)};
   int b = (*tint_module_vars.a);
   int _b = (*tint_module_vars._a);
-  (*tint_module_vars.s) = (b + _b);
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(b) + as_type<uint>(_b)));
 }

diff --git a/test/tint/identifiers/underscore/prefix/upper/alias.wgsl.expected.ir.msl b/test/tint/identifiers/underscore/prefix/upper/alias.wgsl.expected.ir.msl
index 7dad3f8..3e5c9c4 100644
--- a/test/tint/identifiers/underscore/prefix/upper/alias.wgsl.expected.ir.msl
+++ b/test/tint/identifiers/underscore/prefix/upper/alias.wgsl.expected.ir.msl

@@ -9,5 +9,5 @@
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s};
   int c = 0;
   int d = 0;
-  (*tint_module_vars.s) = (c + d);
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(c) + as_type<uint>(d)));
 }

diff --git a/test/tint/identifiers/underscore/prefix/upper/let.wgsl.expected.ir.msl b/test/tint/identifiers/underscore/prefix/upper/let.wgsl.expected.ir.msl
index deecd22..a43607c 100644
--- a/test/tint/identifiers/underscore/prefix/upper/let.wgsl.expected.ir.msl
+++ b/test/tint/identifiers/underscore/prefix/upper/let.wgsl.expected.ir.msl

@@ -11,5 +11,5 @@
   int const _A = 2;
   int const B = A;
   int const _B = _A;
-  (*tint_module_vars.s) = (((A + _A) + B) + _B);
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(A) + as_type<uint>(_A)))) + as_type<uint>(B)))) + as_type<uint>(_B)));
 }

diff --git a/test/tint/identifiers/underscore/prefix/upper/struct.wgsl.expected.ir.msl b/test/tint/identifiers/underscore/prefix/upper/struct.wgsl.expected.ir.msl
index e6f308c..ed7b695 100644
--- a/test/tint/identifiers/underscore/prefix/upper/struct.wgsl.expected.ir.msl
+++ b/test/tint/identifiers/underscore/prefix/upper/struct.wgsl.expected.ir.msl

@@ -13,5 +13,5 @@
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s};
   _A const c = _A{};
   int const d = c._B;
-  (*tint_module_vars.s) = (c._B + d);
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(c._B) + as_type<uint>(d)));
 }

diff --git a/test/tint/identifiers/underscore/prefix/upper/var.wgsl.expected.ir.msl b/test/tint/identifiers/underscore/prefix/upper/var.wgsl.expected.ir.msl
index a1383a9..1cfb128 100644
--- a/test/tint/identifiers/underscore/prefix/upper/var.wgsl.expected.ir.msl
+++ b/test/tint/identifiers/underscore/prefix/upper/var.wgsl.expected.ir.msl

@@ -13,5 +13,5 @@
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.s=s, .A=(&A), ._A=(&_A)};
   int B = (*tint_module_vars.A);
   int _B = (*tint_module_vars._A);
-  (*tint_module_vars.s) = (B + _B);
+  (*tint_module_vars.s) = as_type<int>((as_type<uint>(B) + as_type<uint>(_B)));
 }

diff --git a/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl b/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl
index b030c35..000994f 100644
--- a/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/continue_in_switch.wgsl.expected.ir.msl

@@ -17,7 +17,7 @@
         case 0:
         {
           {
-            i = (i + 1);
+            i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
           }
           continue;
         }
@@ -27,7 +27,7 @@
         }
       }
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl
index e5a8abd..5755d1f 100644
--- a/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl
+++ b/test/tint/loops/continue_in_switch_with_breakif.wgsl.expected.ir.msl

@@ -23,13 +23,13 @@
       }
       if (tint_continue) {
         {
-          i = (i + 1);
+          i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
           if ((i >= 4)) { break; }
         }
         continue;
       }
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
         if ((i >= 4)) { break; }
       }
       continue;

diff --git a/test/tint/loops/loop.wgsl.expected.ir.msl b/test/tint/loops/loop.wgsl.expected.ir.msl
index 4b11fde..5062fd6 100644
--- a/test/tint/loops/loop.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop.wgsl.expected.ir.msl

@@ -9,7 +9,7 @@
   int i = 0;
   {
     TINT_ISOLATE_UB(tint_volatile_true) while(true) {
-      i = (i + 1);
+      i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       if ((i > 4)) {
         return i;
       }

diff --git a/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl b/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl
index 5be8b31..6348b86 100644
--- a/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop_with_break_if.wgsl.expected.ir.msl

@@ -13,7 +13,7 @@
         return i;
       }
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
         if ((i == 4)) { break; }
       }
       continue;

diff --git a/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl b/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl
index ca6575b..6d7a2b3 100644
--- a/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl
+++ b/test/tint/loops/loop_with_continuing.wgsl.expected.ir.msl

@@ -13,7 +13,7 @@
         return i;
       }
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/loops/multiple_continues.wgsl.expected.ir.msl b/test/tint/loops/multiple_continues.wgsl.expected.ir.msl
index d75eccc..a6ac5c9 100644
--- a/test/tint/loops/multiple_continues.wgsl.expected.ir.msl
+++ b/test/tint/loops/multiple_continues.wgsl.expected.ir.msl

@@ -17,21 +17,21 @@
         case 0:
         {
           {
-            i = (i + 1);
+            i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
           }
           continue;
         }
         case 1:
         {
           {
-            i = (i + 1);
+            i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
           }
           continue;
         }
         case 2:
         {
           {
-            i = (i + 1);
+            i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
           }
           continue;
         }
@@ -41,7 +41,7 @@
         }
       }
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/loops/multiple_switch.wgsl.expected.ir.msl b/test/tint/loops/multiple_switch.wgsl.expected.ir.msl
index 24b152c..7605424 100644
--- a/test/tint/loops/multiple_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/multiple_switch.wgsl.expected.ir.msl

@@ -18,7 +18,7 @@
         case 0:
         {
           {
-            i = (i + 1);
+            i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
           }
           continue;
         }
@@ -31,7 +31,7 @@
         case 0:
         {
           {
-            i = (i + 1);
+            i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
           }
           continue;
         }
@@ -41,7 +41,7 @@
         }
       }
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl
index 6309932..0935e8e 100644
--- a/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_loop_switch.wgsl.expected.ir.msl

@@ -24,7 +24,7 @@
             case 0:
             {
               {
-                j = (j + 2);
+                j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
               }
               continue;
             }
@@ -34,13 +34,13 @@
             }
           }
           {
-            j = (j + 2);
+            j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
           }
           continue;
         }
       }
       {
-        i = (i + 2);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
       }
       continue;
     }

diff --git a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl
index 7e16bb1..368eff5 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch.wgsl.expected.ir.msl

@@ -27,7 +27,7 @@
                 case 0:
                 {
                   {
-                    j = (j + 2);
+                    j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
                   }
                   continue;
                 }
@@ -37,13 +37,13 @@
                 }
               }
               {
-                j = (j + 2);
+                j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
               }
               continue;
             }
           }
           {
-            i = (i + 2);
+            i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
           }
           continue;
         }
@@ -53,7 +53,7 @@
         }
       }
       {
-        i = (i + 2);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
       }
       continue;
     }

diff --git a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl
index 73f3fab..82ceefb 100644
--- a/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_loop_switch_switch.wgsl.expected.ir.msl

@@ -28,7 +28,7 @@
                 case 0:
                 {
                   {
-                    j = (j + 2);
+                    j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
                   }
                   continue;
                 }
@@ -38,7 +38,7 @@
                     case 0:
                     {
                       {
-                        j = (j + 2);
+                        j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
                       }
                       continue;
                     }
@@ -55,13 +55,13 @@
                 }
               }
               {
-                j = (j + 2);
+                j = as_type<int>((as_type<uint>(j) + as_type<uint>(2)));
               }
               continue;
             }
           }
           {
-            i = (i + 2);
+            i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
           }
           continue;
         }
@@ -71,7 +71,7 @@
         }
       }
       {
-        i = (i + 2);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
       }
       continue;
     }

diff --git a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl
index c48d971..3946cf3 100644
--- a/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loop_switch_switch.wgsl.expected.ir.msl

@@ -21,7 +21,7 @@
             case 0:
             {
               {
-                i = (i + 2);
+                i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
               }
               continue;
             }
@@ -38,7 +38,7 @@
         }
       }
       {
-        i = (i + 2);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(2)));
       }
       continue;
     }

diff --git a/test/tint/loops/nested_loops.wgsl.expected.ir.msl b/test/tint/loops/nested_loops.wgsl.expected.ir.msl
index fe853a5..fe13807 100644
--- a/test/tint/loops/nested_loops.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loops.wgsl.expected.ir.msl

@@ -10,13 +10,13 @@
   int j = 0;
   {
     TINT_ISOLATE_UB(tint_volatile_true) while(true) {
-      i = (i + 1);
+      i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       if ((i > 4)) {
         return 1;
       }
       {
         TINT_ISOLATE_UB(tint_volatile_true_1) while(true) {
-          j = (j + 1);
+          j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)));
           if ((j > 4)) {
             return 2;
           }

diff --git a/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl b/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl
index 71435a3..1e711ce 100644
--- a/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl
+++ b/test/tint/loops/nested_loops_with_continuing.wgsl.expected.ir.msl

@@ -19,7 +19,7 @@
             return 2;
           }
           {
-            j = (j + 1);
+            j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)));
           }
           continue;
         }

diff --git a/test/tint/loops/single_continue.wgsl.expected.ir.msl b/test/tint/loops/single_continue.wgsl.expected.ir.msl
index 6f505df..ef52f7b 100644
--- a/test/tint/loops/single_continue.wgsl.expected.ir.msl
+++ b/test/tint/loops/single_continue.wgsl.expected.ir.msl

@@ -17,7 +17,7 @@
         case 0:
         {
           {
-            i = (i + 1);
+            i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
           }
           continue;
         }
@@ -27,7 +27,7 @@
         }
       }
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/loops/while.wgsl.expected.ir.msl b/test/tint/loops/while.wgsl.expected.ir.msl
index dbebdfb..a2e7eb1 100644
--- a/test/tint/loops/while.wgsl.expected.ir.msl
+++ b/test/tint/loops/while.wgsl.expected.ir.msl

@@ -13,7 +13,7 @@
       } else {
         break;
       }
-      i = (i + 1);
+      i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       {
       }
       continue;

diff --git a/test/tint/loops/while_with_continue.wgsl.expected.ir.msl b/test/tint/loops/while_with_continue.wgsl.expected.ir.msl
index dbebdfb..a2e7eb1 100644
--- a/test/tint/loops/while_with_continue.wgsl.expected.ir.msl
+++ b/test/tint/loops/while_with_continue.wgsl.expected.ir.msl

@@ -13,7 +13,7 @@
       } else {
         break;
       }
-      i = (i + 1);
+      i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       {
       }
       continue;

diff --git a/test/tint/ptr_ref/load/global/i32.spvasm.expected.ir.msl b/test/tint/ptr_ref/load/global/i32.spvasm.expected.ir.msl
index 84bd5c4..87d115f 100644
--- a/test/tint/ptr_ref/load/global/i32.spvasm.expected.ir.msl
+++ b/test/tint/ptr_ref/load/global/i32.spvasm.expected.ir.msl

@@ -6,7 +6,7 @@
 };
 
 void main_1(tint_module_vars_struct tint_module_vars) {
-  int const x_11 = ((*tint_module_vars.I) + 1);
+  int const x_11 = as_type<int>((as_type<uint>((*tint_module_vars.I)) + as_type<uint>(1)));
 }
 
 kernel void tint_symbol() {

diff --git a/test/tint/ptr_ref/load/global/i32.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/global/i32.wgsl.expected.ir.msl
index 3fbee5d..bdddc66 100644
--- a/test/tint/ptr_ref/load/global/i32.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/global/i32.wgsl.expected.ir.msl

@@ -9,5 +9,5 @@
   thread int I = 0;
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.I=(&I)};
   int const i = (*tint_module_vars.I);
-  int const u = (i + 1);
+  int const u = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
 }

diff --git a/test/tint/ptr_ref/load/local/i32.spvasm.expected.ir.msl b/test/tint/ptr_ref/load/local/i32.spvasm.expected.ir.msl
index cca18cf..f1096fd 100644
--- a/test/tint/ptr_ref/load/local/i32.spvasm.expected.ir.msl
+++ b/test/tint/ptr_ref/load/local/i32.spvasm.expected.ir.msl

@@ -4,7 +4,7 @@
 void main_1() {
   int i = 0;
   i = 123;
-  int const x_12 = (i + 1);
+  int const x_12 = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
 }
 
 kernel void tint_symbol() {

diff --git a/test/tint/ptr_ref/load/local/i32.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/local/i32.wgsl.expected.ir.msl
index 51e2ba4..3eaf723 100644
--- a/test/tint/ptr_ref/load/local/i32.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/local/i32.wgsl.expected.ir.msl

@@ -3,5 +3,5 @@
 
 kernel void tint_symbol() {
   int i = 123;
-  int const u = (i + 1);
+  int const u = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
 }

diff --git a/test/tint/ptr_ref/load/local/ptr_function.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/local/ptr_function.wgsl.expected.ir.msl
index 3c5aef3..2670a93 100644
--- a/test/tint/ptr_ref/load/local/ptr_function.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/local/ptr_function.wgsl.expected.ir.msl

@@ -4,5 +4,5 @@
 kernel void tint_symbol() {
   int i = 123;
   thread int* const p = (&i);
-  int const u = ((*p) + 1);
+  int const u = as_type<int>((as_type<uint>((*p)) + as_type<uint>(1)));
 }

diff --git a/test/tint/ptr_ref/load/local/ptr_private.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/local/ptr_private.wgsl.expected.ir.msl
index abc43ed..bfaf75c 100644
--- a/test/tint/ptr_ref/load/local/ptr_private.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/local/ptr_private.wgsl.expected.ir.msl

@@ -9,5 +9,5 @@
   thread int i = 123;
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.i=(&i)};
   thread int* const p = tint_module_vars.i;
-  int const u = ((*p) + 1);
+  int const u = as_type<int>((as_type<uint>((*p)) + as_type<uint>(1)));
 }

diff --git a/test/tint/ptr_ref/load/local/ptr_storage.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/local/ptr_storage.wgsl.expected.ir.msl
index 150eb1c..bbb20e2 100644
--- a/test/tint/ptr_ref/load/local/ptr_storage.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/local/ptr_storage.wgsl.expected.ir.msl

@@ -12,5 +12,5 @@
 kernel void tint_symbol(device S* v [[buffer(0)]]) {
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v=v};
   device int* const p = (&(*tint_module_vars.v).a);
-  int const u = ((*p) + 1);
+  int const u = as_type<int>((as_type<uint>((*p)) + as_type<uint>(1)));
 }

diff --git a/test/tint/ptr_ref/load/local/ptr_uniform.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/local/ptr_uniform.wgsl.expected.ir.msl
index 153f6c8..7398546 100644
--- a/test/tint/ptr_ref/load/local/ptr_uniform.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/local/ptr_uniform.wgsl.expected.ir.msl

@@ -12,5 +12,5 @@
 kernel void tint_symbol(const constant S* v [[buffer(0)]]) {
   tint_module_vars_struct const tint_module_vars = tint_module_vars_struct{.v=v};
   const constant int* const p = (&(*tint_module_vars.v).a);
-  int const u = ((*p) + 1);
+  int const u = as_type<int>((as_type<uint>((*p)) + as_type<uint>(1)));
 }

diff --git a/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.ir.msl
index 4f29644..d5b0976 100644
--- a/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/local/ptr_workgroup.wgsl.expected.ir.msl

@@ -16,7 +16,7 @@
   threadgroup_barrier(mem_flags::mem_threadgroup);
   (*tint_module_vars.i) = 123;
   threadgroup int* const p = tint_module_vars.i;
-  int const u = ((*p) + 1);
+  int const u = as_type<int>((as_type<uint>((*p)) + as_type<uint>(1)));
 }
 
 kernel void tint_symbol(uint tint_local_index [[thread_index_in_threadgroup]], threadgroup tint_symbol_2* v [[threadgroup(0)]]) {

diff --git a/test/tint/ptr_ref/load/param/ptr.spvasm.expected.ir.msl b/test/tint/ptr_ref/load/param/ptr.spvasm.expected.ir.msl
index 50dcbc8..756596b 100644
--- a/test/tint/ptr_ref/load/param/ptr.spvasm.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/ptr.spvasm.expected.ir.msl

@@ -3,7 +3,7 @@
 
 int func(int value, thread int* const pointer) {
   int const x_9 = (*pointer);
-  return (value + x_9);
+  return as_type<int>((as_type<uint>(value) + as_type<uint>(x_9)));
 }
 
 void main_1() {

diff --git a/test/tint/ptr_ref/load/param/ptr.wgsl.expected.ir.msl b/test/tint/ptr_ref/load/param/ptr.wgsl.expected.ir.msl
index 2240365..fa1b10f 100644
--- a/test/tint/ptr_ref/load/param/ptr.wgsl.expected.ir.msl
+++ b/test/tint/ptr_ref/load/param/ptr.wgsl.expected.ir.msl

@@ -2,7 +2,7 @@
 using namespace metal;
 
 int func(int value, thread int* const pointer) {
-  return (value + (*pointer));
+  return as_type<int>((as_type<uint>(value) + as_type<uint>((*pointer))));
 }
 
 kernel void tint_symbol() {

diff --git a/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.ir.msl b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.ir.msl
index c91bed2..9d528ba 100644
--- a/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.ir.msl
+++ b/test/tint/ptr_sugar/compound_assign_index.wgsl.expected.ir.msl

@@ -4,25 +4,25 @@
 void deref() {
   int3 a = 0;
   thread int3* const p = (&a);
-  (*p)[0] = ((*p)[0] + 42);
+  (*p)[0] = as_type<int>((as_type<uint>((*p)[0]) + as_type<uint>(42)));
 }
 
 void no_deref() {
   int3 a = 0;
   thread int3* const p = (&a);
-  (*p)[0] = ((*p)[0] + 42);
+  (*p)[0] = as_type<int>((as_type<uint>((*p)[0]) + as_type<uint>(42)));
 }
 
 void deref_inc() {
   int3 a = 0;
   thread int3* const p = (&a);
-  (*p)[0] = ((*p)[0] + 1);
+  (*p)[0] = as_type<int>((as_type<uint>((*p)[0]) + as_type<uint>(1)));
 }
 
 void no_deref_inc() {
   int3 a = 0;
   thread int3* const p = (&a);
-  (*p)[0] = ((*p)[0] + 1);
+  (*p)[0] = as_type<int>((as_type<uint>((*p)[0]) + as_type<uint>(1)));
 }
 
 kernel void tint_symbol() {

diff --git a/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.ir.msl b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.ir.msl
index 9d5b32d..ae0d826 100644
--- a/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.ir.msl
+++ b/test/tint/ptr_sugar/compound_assign_member.wgsl.expected.ir.msl

@@ -4,13 +4,13 @@
 void deref() {
   int3 a = 0;
   thread int3* const p = (&a);
-  (*p)[0u] = ((*p)[0u] + 42);
+  (*p)[0u] = as_type<int>((as_type<uint>((*p)[0u]) + as_type<uint>(42)));
 }
 
 void no_deref() {
   int3 a = 0;
   thread int3* const p = (&a);
-  (*p)[0u] = ((*p)[0u] + 42);
+  (*p)[0u] = as_type<int>((as_type<uint>((*p)[0u]) + as_type<uint>(42)));
 }
 
 kernel void tint_symbol() {

diff --git a/test/tint/samples/compute_boids.wgsl.expected.ir.msl b/test/tint/samples/compute_boids.wgsl.expected.ir.msl
index 0afa2a4..394da95 100644
--- a/test/tint/samples/compute_boids.wgsl.expected.ir.msl
+++ b/test/tint/samples/compute_boids.wgsl.expected.ir.msl

@@ -101,7 +101,7 @@
       float const v_3 = distance(pos, vPos);
       if ((v_3 < (*tint_module_vars.params).rule1Distance)) {
         cMass = (cMass + pos);
-        cMassCount = (cMassCount + 1);
+        cMassCount = as_type<int>((as_type<uint>(cMassCount) + as_type<uint>(1)));
       }
       float const v_4 = distance(pos, vPos);
       if ((v_4 < (*tint_module_vars.params).rule2Distance)) {
@@ -110,7 +110,7 @@
       float const v_5 = distance(pos, vPos);
       if ((v_5 < (*tint_module_vars.params).rule3Distance)) {
         cVel = (cVel + vel);
-        cVelCount = (cVelCount + 1);
+        cVelCount = as_type<int>((as_type<uint>(cVelCount) + as_type<uint>(1)));
       }
       {
         i = (i + 1u);

diff --git a/test/tint/shadowing/loop.wgsl.expected.ir.msl b/test/tint/shadowing/loop.wgsl.expected.ir.msl
index f34d31d..32f654a 100644
--- a/test/tint/shadowing/loop.wgsl.expected.ir.msl
+++ b/test/tint/shadowing/loop.wgsl.expected.ir.msl

@@ -29,7 +29,7 @@
       int x = (*tint_module_vars.output)[i];
       {
         int x = (*tint_module_vars.output)[x];
-        i = (i + x);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(x)));
         if ((i > 10)) { break; }
       }
       continue;

diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.ir.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.ir.msl
index 632ed2a..c4065b5 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.ir.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_body.wgsl.expected.ir.msl

@@ -46,7 +46,7 @@
       }
       s1.a1[(*tint_module_vars.uniforms).i] = v;
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.ir.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.ir.msl
index 28e960a..5c483ed 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_continuing.wgsl.expected.ir.msl

@@ -44,7 +44,7 @@
       } else {
         break;
       }
-      i = (i + 1);
+      i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       {
         s1.a1[(*tint_module_vars.uniforms).i] = v;
       }

diff --git a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.ir.msl b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.ir.msl
index 5f983f3..1226caa 100644
--- a/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.ir.msl
+++ b/test/tint/statements/assign/indexed_assign_to_array_in_struct/in_for_loop_init.wgsl.expected.ir.msl

@@ -46,7 +46,7 @@
         break;
       }
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/statements/assign/phony/call.wgsl.expected.ir.msl b/test/tint/statements/assign/phony/call.wgsl.expected.ir.msl
index 56f7731..f027333 100644
--- a/test/tint/statements/assign/phony/call.wgsl.expected.ir.msl
+++ b/test/tint/statements/assign/phony/call.wgsl.expected.ir.msl

@@ -2,7 +2,7 @@
 using namespace metal;
 
 int f(int a, int b, int c) {
-  return ((a * b) + c);
+  return as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) * as_type<uint>(b)))) + as_type<uint>(c)));
 }
 
 kernel void tint_symbol() {

diff --git a/test/tint/statements/assign/phony/multiple_side_effects.wgsl.expected.ir.msl b/test/tint/statements/assign/phony/multiple_side_effects.wgsl.expected.ir.msl
index 94223e3..5475ee2 100644
--- a/test/tint/statements/assign/phony/multiple_side_effects.wgsl.expected.ir.msl
+++ b/test/tint/statements/assign/phony/multiple_side_effects.wgsl.expected.ir.msl

@@ -2,11 +2,11 @@
 using namespace metal;
 
 int f(int a, int b, int c) {
-  return ((a * b) + c);
+  return as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) * as_type<uint>(b)))) + as_type<uint>(c)));
 }
 
 kernel void tint_symbol() {
   int const v = f(1, 2, 3);
   int const v_1 = f(4, 5, 6);
-  int const v_2 = (v + (v_1 * f(7, f(8, 9, 10), 11)));
+  as_type<int>((as_type<uint>(v) + as_type<uint>(as_type<int>((as_type<uint>(v_1) * as_type<uint>(f(7, f(8, 9, 10), 11)))))));
 }

diff --git a/test/tint/statements/compound_assign/complex_lhs.wgsl.expected.ir.msl b/test/tint/statements/compound_assign/complex_lhs.wgsl.expected.ir.msl
index 29f7eb7..f925068 100644
--- a/test/tint/statements/compound_assign/complex_lhs.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/complex_lhs.wgsl.expected.ir.msl

@@ -22,12 +22,12 @@
 };
 
 int foo(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 1);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(1)));
   return (*tint_module_vars.counter);
 }
 
 int bar(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.counter) = ((*tint_module_vars.counter) + 2);
+  (*tint_module_vars.counter) = as_type<int>((as_type<uint>((*tint_module_vars.counter)) + as_type<uint>(2)));
   return (*tint_module_vars.counter);
 }
 
@@ -36,5 +36,5 @@
   thread S* const p = (&x);
   thread int4* const v = (&(*p).a[foo(tint_module_vars)]);
   int const v_1 = bar(tint_module_vars);
-  (*v)[v_1] = ((*v)[v_1] + 5);
+  (*v)[v_1] = as_type<int>((as_type<uint>((*v)[v_1]) + as_type<uint>(5)));
 }

diff --git a/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.ir.msl b/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.ir.msl
index 29c27e7..c81a918 100644
--- a/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/divide_by_zero.wgsl.expected.ir.msl

@@ -11,7 +11,7 @@
   bool const v_1 = bool((v & uint((rhs == -1))));
   uint const v_2 = uint((rhs == 0));
   int const v_3 = select(rhs, 1, bool((v_2 | uint(v_1))));
-  return (lhs - ((lhs / v_3) * v_3));
+  return as_type<int>((as_type<uint>(lhs) - as_type<uint>(as_type<int>((as_type<uint>((lhs / v_3)) * as_type<uint>(v_3))))));
 }
 
 int tint_div_i32(int lhs, int rhs) {

diff --git a/test/tint/statements/compound_assign/scalar/minus.wgsl.expected.ir.msl b/test/tint/statements/compound_assign/scalar/minus.wgsl.expected.ir.msl
index d9c1a73..23d93a1 100644
--- a/test/tint/statements/compound_assign/scalar/minus.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/scalar/minus.wgsl.expected.ir.msl

@@ -10,5 +10,5 @@
 };
 
 void foo(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.v).a = ((*tint_module_vars.v).a - 2);
+  (*tint_module_vars.v).a = as_type<int>((as_type<uint>((*tint_module_vars.v).a) - as_type<uint>(2)));
 }

diff --git a/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.ir.msl b/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.ir.msl
index 13c96b9..679bca5 100644
--- a/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/scalar/modulo.wgsl.expected.ir.msl

@@ -14,7 +14,7 @@
   bool const v_2 = bool((v_1 & uint((rhs == -1))));
   uint const v_3 = uint((rhs == 0));
   int const v_4 = select(rhs, 1, bool((v_3 | uint(v_2))));
-  return (lhs - ((lhs / v_4) * v_4));
+  return as_type<int>((as_type<uint>(lhs) - as_type<uint>(as_type<int>((as_type<uint>((lhs / v_4)) * as_type<uint>(v_4))))));
 }
 
 void foo(tint_module_vars_struct tint_module_vars) {

diff --git a/test/tint/statements/compound_assign/scalar/plus.wgsl.expected.ir.msl b/test/tint/statements/compound_assign/scalar/plus.wgsl.expected.ir.msl
index 9c321cc..ce656cb 100644
--- a/test/tint/statements/compound_assign/scalar/plus.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/scalar/plus.wgsl.expected.ir.msl

@@ -10,5 +10,5 @@
 };
 
 void foo(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.v).a = ((*tint_module_vars.v).a + 2);
+  (*tint_module_vars.v).a = as_type<int>((as_type<uint>((*tint_module_vars.v).a) + as_type<uint>(2)));
 }

diff --git a/test/tint/statements/compound_assign/scalar/times.wgsl.expected.ir.msl b/test/tint/statements/compound_assign/scalar/times.wgsl.expected.ir.msl
index aa0f3d6..08fbca9 100644
--- a/test/tint/statements/compound_assign/scalar/times.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/scalar/times.wgsl.expected.ir.msl

@@ -10,5 +10,5 @@
 };
 
 void foo(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.v).a = ((*tint_module_vars.v).a * 2);
+  (*tint_module_vars.v).a = as_type<int>((as_type<uint>((*tint_module_vars.v).a) * as_type<uint>(2)));
 }

diff --git a/test/tint/statements/compound_assign/vector/minus.wgsl.expected.ir.msl b/test/tint/statements/compound_assign/vector/minus.wgsl.expected.ir.msl
index 37ca9b6..8d1d1a1 100644
--- a/test/tint/statements/compound_assign/vector/minus.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/vector/minus.wgsl.expected.ir.msl

@@ -10,5 +10,5 @@
 };
 
 void foo(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.v).a = ((*tint_module_vars.v).a - int4(2));
+  (*tint_module_vars.v).a = as_type<int4>((as_type<uint4>((*tint_module_vars.v).a) - as_type<uint4>(int4(2))));
 }

diff --git a/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.ir.msl b/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.ir.msl
index 095716d..24529d8 100644
--- a/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/vector/modulo-scalar.wgsl.expected.ir.msl

@@ -14,7 +14,7 @@
   bool4 const v_2 = bool4((v_1 & uint4((rhs == int4(-1)))));
   uint4 const v_3 = uint4((rhs == int4(0)));
   int4 const v_4 = select(rhs, int4(1), bool4((v_3 | uint4(v_2))));
-  return (lhs - ((lhs / v_4) * v_4));
+  return as_type<int4>((as_type<uint4>(lhs) - as_type<uint4>(as_type<int4>((as_type<uint4>((lhs / v_4)) * as_type<uint4>(v_4))))));
 }
 
 void foo(tint_module_vars_struct tint_module_vars) {

diff --git a/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.ir.msl b/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.ir.msl
index 3783e0f..8d79171 100644
--- a/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/vector/modulo.wgsl.expected.ir.msl

@@ -14,7 +14,7 @@
   bool4 const v_2 = bool4((v_1 & uint4((rhs == int4(-1)))));
   uint4 const v_3 = uint4((rhs == int4(0)));
   int4 const v_4 = select(rhs, int4(1), bool4((v_3 | uint4(v_2))));
-  return (lhs - ((lhs / v_4) * v_4));
+  return as_type<int4>((as_type<uint4>(lhs) - as_type<uint4>(as_type<int4>((as_type<uint4>((lhs / v_4)) * as_type<uint4>(v_4))))));
 }
 
 void foo(tint_module_vars_struct tint_module_vars) {

diff --git a/test/tint/statements/compound_assign/vector/plus.wgsl.expected.ir.msl b/test/tint/statements/compound_assign/vector/plus.wgsl.expected.ir.msl
index bc3e299..791b69b 100644
--- a/test/tint/statements/compound_assign/vector/plus.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/vector/plus.wgsl.expected.ir.msl

@@ -10,5 +10,5 @@
 };
 
 void foo(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.v).a = ((*tint_module_vars.v).a + int4(2));
+  (*tint_module_vars.v).a = as_type<int4>((as_type<uint4>((*tint_module_vars.v).a) + as_type<uint4>(int4(2))));
 }

diff --git a/test/tint/statements/compound_assign/vector/times.wgsl.expected.ir.msl b/test/tint/statements/compound_assign/vector/times.wgsl.expected.ir.msl
index b1219fb..43b841a 100644
--- a/test/tint/statements/compound_assign/vector/times.wgsl.expected.ir.msl
+++ b/test/tint/statements/compound_assign/vector/times.wgsl.expected.ir.msl

@@ -10,5 +10,5 @@
 };
 
 void foo(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.v).a = ((*tint_module_vars.v).a * int4(2));
+  (*tint_module_vars.v).a = as_type<int4>((as_type<uint4>((*tint_module_vars.v).a) * as_type<uint4>(int4(2))));
 }

diff --git a/test/tint/statements/decrement/complex.wgsl.expected.ir.msl b/test/tint/statements/decrement/complex.wgsl.expected.ir.msl
index 0c8623d..d932ca4 100644
--- a/test/tint/statements/decrement/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/complex.wgsl.expected.ir.msl

@@ -61,7 +61,7 @@
     int const v_1 = idx1(tint_module_vars);
     device int4* const v_2 = (&(*tint_module_vars.tint_symbol)[v_1].a[idx2(tint_module_vars)]);
     int const v_3 = idx3(tint_module_vars);
-    (*v_2)[v_3] = ((*v_2)[v_3] - 1);
+    (*v_2)[v_3] = as_type<int>((as_type<uint>((*v_2)[v_3]) - as_type<uint>(1)));
     TINT_ISOLATE_UB(tint_volatile_true) while(true) {
       if (((*tint_module_vars.v) < 10u)) {
       } else {
@@ -71,7 +71,7 @@
         int const v_4 = idx4(tint_module_vars);
         device int4* const v_5 = (&(*tint_module_vars.tint_symbol)[v_4].a[idx5(tint_module_vars)]);
         int const v_6 = idx6(tint_module_vars);
-        (*v_5)[v_6] = ((*v_5)[v_6] - 1);
+        (*v_5)[v_6] = as_type<int>((as_type<uint>((*v_5)[v_6]) - as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/statements/decrement/function.wgsl.expected.ir.msl b/test/tint/statements/decrement/function.wgsl.expected.ir.msl
index 6e6144b..de99bff 100644
--- a/test/tint/statements/decrement/function.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/function.wgsl.expected.ir.msl

@@ -3,5 +3,5 @@
 
 void tint_symbol() {
   int i = 0;
-  i = (i - 1);
+  i = as_type<int>((as_type<uint>(i) - as_type<uint>(1)));
 }

diff --git a/test/tint/statements/decrement/private.wgsl.expected.ir.msl b/test/tint/statements/decrement/private.wgsl.expected.ir.msl
index 16d28e2..02bbb7c 100644
--- a/test/tint/statements/decrement/private.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/private.wgsl.expected.ir.msl

@@ -6,5 +6,5 @@
 };
 
 void tint_symbol(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.i) = ((*tint_module_vars.i) - 1);
+  (*tint_module_vars.i) = as_type<int>((as_type<uint>((*tint_module_vars.i)) - as_type<uint>(1)));
 }

diff --git a/test/tint/statements/decrement/split.wgsl.expected.ir.msl b/test/tint/statements/decrement/split.wgsl.expected.ir.msl
index bd6a06a..2a6572d 100644
--- a/test/tint/statements/decrement/split.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/split.wgsl.expected.ir.msl

@@ -3,5 +3,5 @@
 
 void tint_symbol() {
   int b = 2;
-  int c = (b - -(b));
+  int c = as_type<int>((as_type<uint>(b) - as_type<uint>(-(b))));
 }

diff --git a/test/tint/statements/decrement/workgroup.wgsl.expected.ir.msl b/test/tint/statements/decrement/workgroup.wgsl.expected.ir.msl
index 66d1c1d..acd8cb4 100644
--- a/test/tint/statements/decrement/workgroup.wgsl.expected.ir.msl
+++ b/test/tint/statements/decrement/workgroup.wgsl.expected.ir.msl

@@ -6,5 +6,5 @@
 };
 
 void tint_symbol(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.i) = ((*tint_module_vars.i) - 1);
+  (*tint_module_vars.i) = as_type<int>((as_type<uint>((*tint_module_vars.i)) - as_type<uint>(1)));
 }

diff --git a/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl b/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl
index bb00e2a..39ed593 100644
--- a/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl
+++ b/test/tint/statements/discard/atomic_in_for_loop_continuing.wgsl.expected.ir.msl

@@ -37,7 +37,7 @@
       } else {
         break;
       }
-      result = (result + i);
+      result = as_type<int>((as_type<uint>(result) + as_type<uint>(i)));
       {
         int v = 0;
         if ((*tint_module_vars.continue_execution)) {

diff --git a/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl b/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl
index 3de5ec2..7303f09 100644
--- a/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl
+++ b/test/tint/statements/discard/multiple_returns.wgsl.expected.ir.msl

@@ -37,7 +37,7 @@
           return;
         }
         {
-          i = (i + 1);
+          i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
           if ((i == 5)) { break; }
         }
         continue;

diff --git a/test/tint/statements/for/basic.wgsl.expected.ir.msl b/test/tint/statements/for/basic.wgsl.expected.ir.msl
index 569ca8b..fb42656 100644
--- a/test/tint/statements/for/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/basic.wgsl.expected.ir.msl

@@ -18,7 +18,7 @@
       }
       some_loop_body();
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/statements/for/complex.wgsl.expected.ir.msl b/test/tint/statements/for/complex.wgsl.expected.ir.msl
index 15f2f24..16f4e4e 100644
--- a/test/tint/statements/for/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/complex.wgsl.expected.ir.msl

@@ -24,9 +24,9 @@
         break;
       }
       some_loop_body();
-      j = (i * 30);
+      j = as_type<int>((as_type<uint>(i) * as_type<uint>(30)));
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/statements/for/continuing/array_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/continuing/array_ctor.wgsl.expected.ir.msl
index 43b21e1..9fe44a7 100644
--- a/test/tint/statements/for/continuing/array_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/continuing/array_ctor.wgsl.expected.ir.msl

@@ -14,7 +14,7 @@
         break;
       }
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl b/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl
index 43b21e1..9fe44a7 100644
--- a/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/continuing/basic.wgsl.expected.ir.msl

@@ -14,7 +14,7 @@
         break;
       }
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl b/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl
index 3152937..c5a4b75 100644
--- a/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl
+++ b/test/tint/statements/for/continuing/struct_ctor.wgsl.expected.ir.msl

@@ -14,7 +14,7 @@
         break;
       }
       {
-        i = (i + 1);
+        i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/statements/increment/complex.wgsl.expected.ir.msl b/test/tint/statements/increment/complex.wgsl.expected.ir.msl
index 132300a..7f62d2b 100644
--- a/test/tint/statements/increment/complex.wgsl.expected.ir.msl
+++ b/test/tint/statements/increment/complex.wgsl.expected.ir.msl

@@ -61,7 +61,7 @@
     int const v_1 = idx1(tint_module_vars);
     device int4* const v_2 = (&(*tint_module_vars.tint_symbol)[v_1].a[idx2(tint_module_vars)]);
     int const v_3 = idx3(tint_module_vars);
-    (*v_2)[v_3] = ((*v_2)[v_3] + 1);
+    (*v_2)[v_3] = as_type<int>((as_type<uint>((*v_2)[v_3]) + as_type<uint>(1)));
     TINT_ISOLATE_UB(tint_volatile_true) while(true) {
       if (((*tint_module_vars.v) < 10u)) {
       } else {
@@ -71,7 +71,7 @@
         int const v_4 = idx4(tint_module_vars);
         device int4* const v_5 = (&(*tint_module_vars.tint_symbol)[v_4].a[idx5(tint_module_vars)]);
         int const v_6 = idx6(tint_module_vars);
-        (*v_5)[v_6] = ((*v_5)[v_6] + 1);
+        (*v_5)[v_6] = as_type<int>((as_type<uint>((*v_5)[v_6]) + as_type<uint>(1)));
       }
       continue;
     }

diff --git a/test/tint/statements/increment/function.wgsl.expected.ir.msl b/test/tint/statements/increment/function.wgsl.expected.ir.msl
index 960e703..3ba0be2 100644
--- a/test/tint/statements/increment/function.wgsl.expected.ir.msl
+++ b/test/tint/statements/increment/function.wgsl.expected.ir.msl

@@ -3,5 +3,5 @@
 
 void tint_symbol() {
   int i = 0;
-  i = (i + 1);
+  i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
 }

diff --git a/test/tint/statements/increment/private.wgsl.expected.ir.msl b/test/tint/statements/increment/private.wgsl.expected.ir.msl
index 9b8cc059..1d4431c 100644
--- a/test/tint/statements/increment/private.wgsl.expected.ir.msl
+++ b/test/tint/statements/increment/private.wgsl.expected.ir.msl

@@ -6,5 +6,5 @@
 };
 
 void tint_symbol(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.i) = ((*tint_module_vars.i) + 1);
+  (*tint_module_vars.i) = as_type<int>((as_type<uint>((*tint_module_vars.i)) + as_type<uint>(1)));
 }

diff --git a/test/tint/statements/increment/workgroup.wgsl.expected.ir.msl b/test/tint/statements/increment/workgroup.wgsl.expected.ir.msl
index 091538d..decca87 100644
--- a/test/tint/statements/increment/workgroup.wgsl.expected.ir.msl
+++ b/test/tint/statements/increment/workgroup.wgsl.expected.ir.msl

@@ -6,5 +6,5 @@
 };
 
 void tint_symbol(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.i) = ((*tint_module_vars.i) + 1);
+  (*tint_module_vars.i) = as_type<int>((as_type<uint>((*tint_module_vars.i)) + as_type<uint>(1)));
 }

diff --git a/test/tint/struct/type_initializer.wgsl.expected.ir.msl b/test/tint/struct/type_initializer.wgsl.expected.ir.msl
index 94f059d..5b31bfe 100644
--- a/test/tint/struct/type_initializer.wgsl.expected.ir.msl
+++ b/test/tint/struct/type_initializer.wgsl.expected.ir.msl

@@ -39,18 +39,18 @@
   int const x = 42;
   S1 const empty = S1{};
   S1 const nonempty = S1{.a=1, .b=2, .c=3, .d=4};
-  S1 const nonempty_with_expr = S1{.a=1, .b=x, .c=(x + 1), .d=nonempty.d};
+  S1 const nonempty_with_expr = S1{.a=1, .b=x, .c=as_type<int>((as_type<uint>(x) + as_type<uint>(1))), .d=nonempty.d};
   S3 const nested_empty = S3{};
   S3 const nested_nonempty = S3{.g=1, .h=S1{.a=2, .b=3, .c=4, .d=5}, .i=S2{.e=6, .f=S1{.a=7, .b=8, .c=9, .d=10}}};
-  S1 const v = S1{.a=2, .b=x, .c=(x + 1), .d=nested_nonempty.i.f.d};
+  S1 const v = S1{.a=2, .b=x, .c=as_type<int>((as_type<uint>(x) + as_type<uint>(1))), .d=nested_nonempty.i.f.d};
   S3 const nested_nonempty_with_expr = S3{.g=1, .h=v, .i=S2{.e=6, .f=nonempty}};
   int const subexpr_empty = 0;
   int const subexpr_nonempty = 2;
-  int const subexpr_nonempty_with_expr = S1{.a=1, .b=x, .c=(x + 1), .d=nonempty.d}.c;
+  int const subexpr_nonempty_with_expr = S1{.a=1, .b=x, .c=as_type<int>((as_type<uint>(x) + as_type<uint>(1))), .d=nonempty.d}.c;
   S1 const subexpr_nested_empty = S1{};
   S1 const subexpr_nested_nonempty = S1{.a=2, .b=3, .c=4, .d=5};
-  S1 const subexpr_nested_nonempty_with_expr = S2{.e=1, .f=S1{.a=2, .b=x, .c=(x + 1), .d=nested_nonempty.i.f.d}}.f;
+  S1 const subexpr_nested_nonempty_with_expr = S2{.e=1, .f=S1{.a=2, .b=x, .c=as_type<int>((as_type<uint>(x) + as_type<uint>(1))), .d=nested_nonempty.i.f.d}}.f;
   tint_array<T, 2> const aosoa_empty = tint_array<T, 2>{};
   tint_array<T, 2> const aosoa_nonempty = tint_array<T, 2>{T{.a=tint_array<int, 2>{1, 2}}, T{.a=tint_array<int, 2>{3, 4}}};
-  tint_array<T, 2> const aosoa_nonempty_with_expr = tint_array<T, 2>{T{.a=tint_array<int, 2>{1, (aosoa_nonempty[0].a[0] + 1)}}, aosoa_nonempty[1]};
+  tint_array<T, 2> const aosoa_nonempty_with_expr = tint_array<T, 2>{T{.a=tint_array<int, 2>{1, as_type<int>((as_type<uint>(aosoa_nonempty[0].a[0]) + as_type<uint>(1)))}}, aosoa_nonempty[1]};
 }

diff --git a/test/tint/switch/switch.wgsl.expected.ir.msl b/test/tint/switch/switch.wgsl.expected.ir.msl
index bca1c72..ac01dec 100644
--- a/test/tint/switch/switch.wgsl.expected.ir.msl
+++ b/test/tint/switch/switch.wgsl.expected.ir.msl

@@ -14,7 +14,7 @@
     }
     default:
     {
-      a_1 = (a_1 + 2);
+      a_1 = as_type<int>((as_type<uint>(a_1) + as_type<uint>(2)));
       break;
     }
   }

diff --git a/test/tint/var/uses/private.wgsl.expected.ir.msl b/test/tint/var/uses/private.wgsl.expected.ir.msl
index 0f148c8..3f1bfd3 100644
--- a/test/tint/var/uses/private.wgsl.expected.ir.msl
+++ b/test/tint/var/uses/private.wgsl.expected.ir.msl

@@ -8,11 +8,11 @@
 };
 
 void uses_a(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.a) = ((*tint_module_vars.a) + 1);
+  (*tint_module_vars.a) = as_type<int>((as_type<uint>((*tint_module_vars.a)) + as_type<uint>(1)));
 }
 
 void uses_b(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.b) = ((*tint_module_vars.b) * 2);
+  (*tint_module_vars.b) = as_type<int>((as_type<uint>((*tint_module_vars.b)) * as_type<uint>(2)));
 }
 
 void uses_a_and_b(tint_module_vars_struct tint_module_vars) {

diff --git a/test/tint/var/uses/workgroup.wgsl.expected.ir.msl b/test/tint/var/uses/workgroup.wgsl.expected.ir.msl
index 9150b12..6e10372 100644
--- a/test/tint/var/uses/workgroup.wgsl.expected.ir.msl
+++ b/test/tint/var/uses/workgroup.wgsl.expected.ir.msl

@@ -21,11 +21,11 @@
 };
 
 void uses_a(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.a) = ((*tint_module_vars.a) + 1);
+  (*tint_module_vars.a) = as_type<int>((as_type<uint>((*tint_module_vars.a)) + as_type<uint>(1)));
 }
 
 void uses_b(tint_module_vars_struct tint_module_vars) {
-  (*tint_module_vars.b) = ((*tint_module_vars.b) * 2);
+  (*tint_module_vars.b) = as_type<int>((as_type<uint>((*tint_module_vars.b)) * as_type<uint>(2)));
 }
 
 void uses_a_and_b(tint_module_vars_struct tint_module_vars) {