[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/src/tint/lang/msl/writer/binary_test.cc b/src/tint/lang/msl/writer/binary_test.cc
index e7f7f5c..9e8f89d 100644
--- a/src/tint/lang/msl/writer/binary_test.cc
+++ b/src/tint/lang/msl/writer/binary_test.cc
@@ -202,13 +202,13 @@
BinaryData{"(left >= right)", core::BinaryOp::kGreaterThanEqual}));
using MslWriterBinaryTest_SignedOverflowDefinedBehaviour = MslWriterTestWithParam<BinaryData>;
-TEST_P(MslWriterBinaryTest_SignedOverflowDefinedBehaviour, DISABLED_Emit) {
+TEST_P(MslWriterBinaryTest_SignedOverflowDefinedBehaviour, Emit_Scalar) {
auto params = GetParam();
auto* func = b.Function("foo", ty.void_());
b.Append(func->Block(), [&] {
- auto* l = b.Let("a", b.Constant(1_i));
- auto* r = b.Let("b", b.Constant(3_i));
+ auto* l = b.Let("left", b.Constant(1_i));
+ auto* r = b.Let("right", b.Constant(3_i));
auto* bin = b.Binary(params.op, ty.i32(), l, r);
b.Let("val", bin);
@@ -220,16 +220,41 @@
void foo() {
int const left = 1;
int const right = 3;
- int const val = )" + params.result +
- R"(;
+ int const val = as_type<int>((as_type<uint>(left) )" +
+ params.result + R"( as_type<uint>(right)));
+}
+)");
+}
+
+TEST_P(MslWriterBinaryTest_SignedOverflowDefinedBehaviour, Emit_Vector) {
+ auto params = GetParam();
+
+ auto* func = b.Function("foo", ty.void_());
+ b.Append(func->Block(), [&] {
+ auto* l = b.Let("left", b.Splat<vec4<i32>>(1_i));
+ auto* r = b.Let("right", b.Splat<vec4<i32>>(3_i));
+
+ auto* bin = b.Binary(params.op, ty.vec4<i32>(), l, r);
+ b.Let("val", bin);
+ b.Return(func);
+ });
+
+ ASSERT_TRUE(Generate()) << err_ << output_.msl;
+ EXPECT_EQ(output_.msl, MetalHeader() + R"(
+void foo() {
+ int4 const left = int4(1);
+ int4 const right = int4(3);
+ int4 const val = as_type<int4>((as_type<uint4>(left) )" +
+ params.result + R"( as_type<uint4>(right)));
}
)");
}
constexpr BinaryData signed_overflow_defined_behaviour_cases[] = {
- {"as_type<int>((as_type<uint>(left) + as_type<uint>(right)))", core::BinaryOp::kAdd},
- {"as_type<int>((as_type<uint>(left) - as_type<uint>(right)))", core::BinaryOp::kSubtract},
- {"as_type<int>((as_type<uint>(left) * as_type<uint>(right)))", core::BinaryOp::kMultiply}};
+ {"+", core::BinaryOp::kAdd},
+ {"-", core::BinaryOp::kSubtract},
+ {"*", core::BinaryOp::kMultiply},
+};
INSTANTIATE_TEST_SUITE_P(MslWriterTest,
MslWriterBinaryTest_SignedOverflowDefinedBehaviour,
testing::ValuesIn(signed_overflow_defined_behaviour_cases));
diff --git a/src/tint/lang/msl/writer/raise/binary_polyfill.cc b/src/tint/lang/msl/writer/raise/binary_polyfill.cc
index 0e579df..e2f8dc6 100644
--- a/src/tint/lang/msl/writer/raise/binary_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/binary_polyfill.cc
@@ -52,15 +52,20 @@
// Find the binary operators that need replacing.
Vector<core::ir::CoreBinary*, 4> fmod_worklist;
Vector<core::ir::CoreBinary*, 4> logical_bool_worklist;
+ Vector<core::ir::CoreBinary*, 4> signed_integer_arithmetic_worklist;
for (auto* inst : ir.Instructions()) {
if (auto* binary = inst->As<core::ir::CoreBinary>()) {
- if (binary->Op() == core::BinaryOp::kModulo &&
- binary->LHS()->Type()->IsFloatScalarOrVector()) {
+ auto op = binary->Op();
+ auto* lhs_type = binary->LHS()->Type();
+ if (op == core::BinaryOp::kModulo && lhs_type->IsFloatScalarOrVector()) {
fmod_worklist.Push(binary);
- } else if ((binary->Op() == core::BinaryOp::kAnd ||
- binary->Op() == core::BinaryOp::kOr) &&
- binary->LHS()->Type()->IsBoolScalarOrVector()) {
+ } else if ((op == core::BinaryOp::kAnd || op == core::BinaryOp::kOr) &&
+ lhs_type->IsBoolScalarOrVector()) {
logical_bool_worklist.Push(binary);
+ } else if ((op == core::BinaryOp::kAdd || op == core::BinaryOp::kMultiply ||
+ op == core::BinaryOp::kSubtract) &&
+ lhs_type->IsSignedIntegerScalarOrVector()) {
+ signed_integer_arithmetic_worklist.Push(binary);
}
}
}
@@ -72,6 +77,9 @@
for (auto* logical_bool : logical_bool_worklist) {
LogicalBool(logical_bool);
}
+ for (auto* signed_arith : signed_integer_arithmetic_worklist) {
+ SignedIntegerArithmetic(signed_arith);
+ }
}
/// Replace a floating point modulo binary instruction with the equivalent MSL intrinsic.
@@ -90,10 +98,7 @@
// integers. Make this explicit in the IR and then convert the result of the binary
// instruction back to a boolean.
auto* result_ty = binary->Result(0)->Type();
- const core::type::Type* int_ty = ty.u32();
- if (auto* vec = result_ty->As<core::type::Vector>()) {
- int_ty = ty.vec(int_ty, vec->Width());
- }
+ auto* int_ty = ty.match_width(ty.u32(), result_ty);
b.InsertBefore(binary, [&] {
auto* int_lhs = b.Convert(int_ty, binary->LHS());
auto* int_rhs = b.Convert(int_ty, binary->RHS());
@@ -102,6 +107,26 @@
});
binary->Destroy();
}
+
+ /// Replace a signed integer arithmetic instruction.
+ /// @param binary the signed integer arithmetic instruction
+ void SignedIntegerArithmetic(core::ir::CoreBinary* binary) {
+ // MSL does not define the behavior of signed integer overflow, so bitcast the operands to
+ // unsigned integers, perform the operation, and then bitcast the result back to a signed
+ // integer.
+ auto* signed_result_ty = binary->Result(0)->Type();
+ auto* unsigned_result_ty = ty.match_width(ty.u32(), signed_result_ty);
+ auto* unsigned_lhs_ty = ty.match_width(ty.u32(), binary->LHS()->Type());
+ auto* unsigned_rhs_ty = ty.match_width(ty.u32(), binary->RHS()->Type());
+ b.InsertBefore(binary, [&] {
+ auto* uint_lhs = b.Bitcast(unsigned_lhs_ty, binary->LHS());
+ auto* uint_rhs = b.Bitcast(unsigned_rhs_ty, binary->RHS());
+ auto* uint_binary = b.Binary(binary->Op(), unsigned_result_ty, uint_lhs, uint_rhs);
+ auto* bitcast = b.Bitcast(signed_result_ty, uint_binary);
+ binary->Result(0)->ReplaceAllUsesWith(bitcast->Result(0));
+ });
+ binary->Destroy();
+ }
};
} // namespace
diff --git a/src/tint/lang/msl/writer/raise/binary_polyfill_test.cc b/src/tint/lang/msl/writer/raise/binary_polyfill_test.cc
index 20ceaaf..8560849 100644
--- a/src/tint/lang/msl/writer/raise/binary_polyfill_test.cc
+++ b/src/tint/lang/msl/writer/raise/binary_polyfill_test.cc
@@ -285,5 +285,227 @@
EXPECT_EQ(expect, str());
}
+TEST_F(MslWriter_BinaryPolyfillTest, IntAdd_Scalar) {
+ auto* lhs = b.FunctionParam<i32>("lhs");
+ auto* rhs = b.FunctionParam<i32>("rhs");
+ auto* func = b.Function("foo", ty.i32());
+ func->SetParams({lhs, rhs});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Add<i32>(lhs, rhs);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%lhs:i32, %rhs:i32):i32 {
+ $B1: {
+ %4:i32 = add %lhs, %rhs
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%lhs:i32, %rhs:i32):i32 {
+ $B1: {
+ %4:u32 = bitcast %lhs
+ %5:u32 = bitcast %rhs
+ %6:u32 = add %4, %5
+ %7:i32 = bitcast %6
+ ret %7
+ }
+}
+)";
+
+ Run(BinaryPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BinaryPolyfillTest, IntMul_Scalar) {
+ auto* lhs = b.FunctionParam<i32>("lhs");
+ auto* rhs = b.FunctionParam<i32>("rhs");
+ auto* func = b.Function("foo", ty.i32());
+ func->SetParams({lhs, rhs});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Multiply<i32>(lhs, rhs);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%lhs:i32, %rhs:i32):i32 {
+ $B1: {
+ %4:i32 = mul %lhs, %rhs
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%lhs:i32, %rhs:i32):i32 {
+ $B1: {
+ %4:u32 = bitcast %lhs
+ %5:u32 = bitcast %rhs
+ %6:u32 = mul %4, %5
+ %7:i32 = bitcast %6
+ ret %7
+ }
+}
+)";
+
+ Run(BinaryPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BinaryPolyfillTest, IntSub_Scalar) {
+ auto* lhs = b.FunctionParam<i32>("lhs");
+ auto* rhs = b.FunctionParam<i32>("rhs");
+ auto* func = b.Function("foo", ty.i32());
+ func->SetParams({lhs, rhs});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Subtract<i32>(lhs, rhs);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%lhs:i32, %rhs:i32):i32 {
+ $B1: {
+ %4:i32 = sub %lhs, %rhs
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%lhs:i32, %rhs:i32):i32 {
+ $B1: {
+ %4:u32 = bitcast %lhs
+ %5:u32 = bitcast %rhs
+ %6:u32 = sub %4, %5
+ %7:i32 = bitcast %6
+ ret %7
+ }
+}
+)";
+
+ Run(BinaryPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BinaryPolyfillTest, IntAdd_Vector) {
+ auto* lhs = b.FunctionParam<vec4<i32>>("lhs");
+ auto* rhs = b.FunctionParam<vec4<i32>>("rhs");
+ auto* func = b.Function("foo", ty.vec4<i32>());
+ func->SetParams({lhs, rhs});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Add<vec4<i32>>(lhs, rhs);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%lhs:vec4<i32>, %rhs:vec4<i32>):vec4<i32> {
+ $B1: {
+ %4:vec4<i32> = add %lhs, %rhs
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%lhs:vec4<i32>, %rhs:vec4<i32>):vec4<i32> {
+ $B1: {
+ %4:vec4<u32> = bitcast %lhs
+ %5:vec4<u32> = bitcast %rhs
+ %6:vec4<u32> = add %4, %5
+ %7:vec4<i32> = bitcast %6
+ ret %7
+ }
+}
+)";
+
+ Run(BinaryPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BinaryPolyfillTest, IntAdd_ScalarVector) {
+ auto* lhs = b.FunctionParam<i32>("lhs");
+ auto* rhs = b.FunctionParam<vec4<i32>>("rhs");
+ auto* func = b.Function("foo", ty.vec4<i32>());
+ func->SetParams({lhs, rhs});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Add<vec4<i32>>(lhs, rhs);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%lhs:i32, %rhs:vec4<i32>):vec4<i32> {
+ $B1: {
+ %4:vec4<i32> = add %lhs, %rhs
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%lhs:i32, %rhs:vec4<i32>):vec4<i32> {
+ $B1: {
+ %4:u32 = bitcast %lhs
+ %5:vec4<u32> = bitcast %rhs
+ %6:vec4<u32> = add %4, %5
+ %7:vec4<i32> = bitcast %6
+ ret %7
+ }
+}
+)";
+
+ Run(BinaryPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
+TEST_F(MslWriter_BinaryPolyfillTest, IntAdd_VectorScalar) {
+ auto* lhs = b.FunctionParam<vec4<i32>>("lhs");
+ auto* rhs = b.FunctionParam<i32>("rhs");
+ auto* func = b.Function("foo", ty.vec4<i32>());
+ func->SetParams({lhs, rhs});
+ b.Append(func->Block(), [&] {
+ auto* result = b.Add<vec4<i32>>(lhs, rhs);
+ b.Return(func, result);
+ });
+
+ auto* src = R"(
+%foo = func(%lhs:vec4<i32>, %rhs:i32):vec4<i32> {
+ $B1: {
+ %4:vec4<i32> = add %lhs, %rhs
+ ret %4
+ }
+}
+)";
+ EXPECT_EQ(src, str());
+
+ auto* expect = R"(
+%foo = func(%lhs:vec4<i32>, %rhs:i32):vec4<i32> {
+ $B1: {
+ %4:vec4<u32> = bitcast %lhs
+ %5:u32 = bitcast %rhs
+ %6:vec4<u32> = add %4, %5
+ %7:vec4<i32> = bitcast %6
+ ret %7
+ }
+}
+)";
+
+ Run(BinaryPolyfill);
+
+ EXPECT_EQ(expect, str());
+}
+
} // namespace
} // namespace tint::msl::writer::raise
diff --git a/src/tint/lang/msl/writer/writer_test.cc b/src/tint/lang/msl/writer/writer_test.cc
index 07dbc21..b0cb64d 100644
--- a/src/tint/lang/msl/writer/writer_test.cc
+++ b/src/tint/lang/msl/writer/writer_test.cc
@@ -75,7 +75,7 @@
(*tint_module_vars.b) = 0;
}
threadgroup_barrier(mem_flags::mem_threadgroup);
- (*tint_module_vars.a) = ((*tint_module_vars.a) + (*tint_module_vars.b));
+ (*tint_module_vars.a) = as_type<int>((as_type<uint>((*tint_module_vars.a)) + as_type<uint>((*tint_module_vars.b))));
}
kernel void bar() {
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) {
