Restore "MSL writer: make signed int overflow defined behaviour"

This reverts commit e33b0baa08.

Added tests/expressions/literals/intmin.wgsl test.

Bug: tint:124
Change-Id: I3d46f939ff20fa377ddb5fcb52f9afe728b8e430
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60441
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Reviewed-by: David Neto <dneto@google.com>

src/writer/msl/generator_impl.cc

@@ -75,6 +75,32 @@ bool last_is_break_or_fallthrough(const ast::BlockStatement* stmts) {
          stmts->last()->Is<ast::FallthroughStatement>();
 }
 
+class ScopedBitCast {
+ public:
+  ScopedBitCast(GeneratorImpl* generator,
+                std::ostream& stream,
+                const sem::Type* curr_type,
+                const sem::Type* target_type)
+      : s(stream) {
+    auto* target_vec_type = target_type->As<sem::Vector>();
+
+    // If we need to promote from scalar to vector, bitcast the scalar to the
+    // vector element type.
+    if (curr_type->is_scalar() && target_vec_type) {
+      target_type = target_vec_type->type();
+    }
+
+    // Bit cast
+    s << "as_type<";
+    generator->EmitType(s, target_type, "");
+    s << ">(";
+  }
+
+  ~ScopedBitCast() { s << ")"; }
+
+ private:
+  std::ostream& s;
+};
 }  // namespace
 
 GeneratorImpl::GeneratorImpl(const Program* program) : TextGenerator(program) {}
@@ -227,8 +253,104 @@ bool GeneratorImpl::EmitAssign(ast::AssignmentStatement* stmt) {
 }
 
 bool GeneratorImpl::EmitBinary(std::ostream& out, ast::BinaryExpression* expr) {
+  auto emit_op = [&] {
+    out << " ";
+
+    switch (expr->op()) {
+      case ast::BinaryOp::kAnd:
+        out << "&";
+        break;
+      case ast::BinaryOp::kOr:
+        out << "|";
+        break;
+      case ast::BinaryOp::kXor:
+        out << "^";
+        break;
+      case ast::BinaryOp::kLogicalAnd:
+        out << "&&";
+        break;
+      case ast::BinaryOp::kLogicalOr:
+        out << "||";
+        break;
+      case ast::BinaryOp::kEqual:
+        out << "==";
+        break;
+      case ast::BinaryOp::kNotEqual:
+        out << "!=";
+        break;
+      case ast::BinaryOp::kLessThan:
+        out << "<";
+        break;
+      case ast::BinaryOp::kGreaterThan:
+        out << ">";
+        break;
+      case ast::BinaryOp::kLessThanEqual:
+        out << "<=";
+        break;
+      case ast::BinaryOp::kGreaterThanEqual:
+        out << ">=";
+        break;
+      case ast::BinaryOp::kShiftLeft:
+        out << "<<";
+        break;
+      case ast::BinaryOp::kShiftRight:
+        // TODO(dsinclair): MSL is based on C++14, and >> in C++14 has
+        // implementation-defined behaviour for negative LHS.  We may have to
+        // generate extra code to implement WGSL-specified behaviour for
+        // negative LHS.
+        out << R"(>>)";
+        break;
+
+      case ast::BinaryOp::kAdd:
+        out << "+";
+        break;
+      case ast::BinaryOp::kSubtract:
+        out << "-";
+        break;
+      case ast::BinaryOp::kMultiply:
+        out << "*";
+        break;
+      case ast::BinaryOp::kDivide:
+        out << "/";
+        break;
+      case ast::BinaryOp::kModulo:
+        out << "%";
+        break;
+      case ast::BinaryOp::kNone:
+        diagnostics_.add_error(diag::System::Writer,
+                               "missing binary operation type");
+        return false;
+    }
+    out << " ";
+    return true;
+  };
+
+  auto signed_type_of = [&](const sem::Type* ty) -> const sem::Type* {
+    if (ty->is_integer_scalar()) {
+      return builder_.create<sem::I32>();
+    } else if (auto* v = ty->As<sem::Vector>()) {
+      return builder_.create<sem::Vector>(builder_.create<sem::I32>(),
+                                          v->Width());
+    }
+    return {};
+  };
+
+  auto unsigned_type_of = [&](const sem::Type* ty) -> const sem::Type* {
+    if (ty->is_integer_scalar()) {
+      return builder_.create<sem::U32>();
+    } else if (auto* v = ty->As<sem::Vector>()) {
+      return builder_.create<sem::Vector>(builder_.create<sem::U32>(),
+                                          v->Width());
+    }
+    return {};
+  };
+
+  auto* lhs_type = TypeOf(expr->lhs())->UnwrapRef();
+  auto* rhs_type = TypeOf(expr->rhs())->UnwrapRef();
+
+  // Handle fmod
   if (expr->op() == ast::BinaryOp::kModulo &&
-      TypeOf(expr)->UnwrapRef()->is_float_scalar_or_vector()) {
+      lhs_type->is_float_scalar_or_vector()) {
     out << "fmod";
     ScopedParen sp(out);
     if (!EmitExpression(out, expr->lhs())) {
@@ -241,80 +363,75 @@ bool GeneratorImpl::EmitBinary(std::ostream& out, ast::BinaryExpression* expr) {
     return true;
   }
 
-  ScopedParen sp(out);
+  // Handle +/-/* of signed values
+  if ((expr->IsAdd() || expr->IsSubtract() || expr->IsMultiply()) &&
+      lhs_type->is_signed_scalar_or_vector() &&
+      rhs_type->is_signed_scalar_or_vector()) {
+    // If lhs or rhs is a vector, use that type (support implicit scalar to
+    // vector promotion)
+    auto* target_type =
+        lhs_type->Is<sem::Vector>()
+            ? lhs_type
+            : (rhs_type->Is<sem::Vector>() ? rhs_type : lhs_type);
 
+    // WGSL defines behaviour for signed overflow, MSL does not. For these
+    // cases, bitcast operands to unsigned, then cast result to signed.
+    ScopedBitCast outer_int_cast(this, out, target_type,
+                                 signed_type_of(target_type));
+    ScopedParen sp(out);
+    {
+      ScopedBitCast lhs_uint_cast(this, out, lhs_type,
+                                  unsigned_type_of(target_type));
+      if (!EmitExpression(out, expr->lhs())) {
+        return false;
+      }
+    }
+    if (!emit_op()) {
+      return false;
+    }
+    {
+      ScopedBitCast rhs_uint_cast(this, out, rhs_type,
+                                  unsigned_type_of(target_type));
+      if (!EmitExpression(out, expr->rhs())) {
+        return false;
+      }
+    }
+    return true;
+  }
+
+  // Handle left bit shifting a signed value
+  // TODO(crbug.com/tint/1077): This may not be necessary. The MSL spec
+  // seems to imply that left shifting a signed value is treated the same as
+  // left shifting an unsigned value, but we need to make sure.
+  if (expr->IsShiftLeft() && lhs_type->is_signed_scalar_or_vector()) {
+    // Shift left: discards top bits, so convert first operand to unsigned
+    // first, then convert result back to signed
+    ScopedBitCast outer_int_cast(this, out, lhs_type, signed_type_of(lhs_type));
+    ScopedParen sp(out);
+    {
+      ScopedBitCast lhs_uint_cast(this, out, lhs_type,
+                                  unsigned_type_of(lhs_type));
+      if (!EmitExpression(out, expr->lhs())) {
+        return false;
+      }
+    }
+    if (!emit_op()) {
+      return false;
+    }
+    if (!EmitExpression(out, expr->rhs())) {
+      return false;
+    }
+    return true;
+  }
+
+  // Emit as usual
+  ScopedParen sp(out);
   if (!EmitExpression(out, expr->lhs())) {
     return false;
   }
-  out << " ";
+  if (!emit_op()) {
-
+    return false;
-  switch (expr->op()) {
-    case ast::BinaryOp::kAnd:
-      out << "&";
-      break;
-    case ast::BinaryOp::kOr:
-      out << "|";
-      break;
-    case ast::BinaryOp::kXor:
-      out << "^";
-      break;
-    case ast::BinaryOp::kLogicalAnd:
-      out << "&&";
-      break;
-    case ast::BinaryOp::kLogicalOr:
-      out << "||";
-      break;
-    case ast::BinaryOp::kEqual:
-      out << "==";
-      break;
-    case ast::BinaryOp::kNotEqual:
-      out << "!=";
-      break;
-    case ast::BinaryOp::kLessThan:
-      out << "<";
-      break;
-    case ast::BinaryOp::kGreaterThan:
-      out << ">";
-      break;
-    case ast::BinaryOp::kLessThanEqual:
-      out << "<=";
-      break;
-    case ast::BinaryOp::kGreaterThanEqual:
-      out << ">=";
-      break;
-    case ast::BinaryOp::kShiftLeft:
-      out << "<<";
-      break;
-    case ast::BinaryOp::kShiftRight:
-      // TODO(dsinclair): MSL is based on C++14, and >> in C++14 has
-      // implementation-defined behaviour for negative LHS.  We may have to
-      // generate extra code to implement WGSL-specified behaviour for negative
-      // LHS.
-      out << R"(>>)";
-      break;
-
-    case ast::BinaryOp::kAdd:
-      out << "+";
-      break;
-    case ast::BinaryOp::kSubtract:
-      out << "-";
-      break;
-    case ast::BinaryOp::kMultiply:
-      out << "*";
-      break;
-    case ast::BinaryOp::kDivide:
-      out << "/";
-      break;
-    case ast::BinaryOp::kModulo:
-      out << "%";
-      break;
-    case ast::BinaryOp::kNone:
-      diagnostics_.add_error(diag::System::Writer,
-                             "missing binary operation type");
-      return false;
   }
-  out << " ";
-
   if (!EmitExpression(out, expr->rhs())) {
     return false;
   }
@@ -2338,6 +2455,53 @@ bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
 
 bool GeneratorImpl::EmitUnaryOp(std::ostream& out,
                                 ast::UnaryOpExpression* expr) {
+  // Handle `-e` when `e` is signed, so that we ensure that if `e` is the
+  // largest negative value, it returns `e`.
+  auto* expr_type = TypeOf(expr->expr())->UnwrapRef();
+  if (expr->op() == ast::UnaryOp::kNegation &&
+      expr_type->is_signed_scalar_or_vector()) {
+    auto fn =
+        utils::GetOrCreate(unary_minus_funcs_, expr_type, [&]() -> std::string {
+          // e.g.:
+          // int tint_unary_minus(const int v) {
+          //     return (v == -2147483648) ? v : -v;
+          // }
+          TextBuffer b;
+          TINT_DEFER(helpers_.Append(b));
+
+          auto fn_name = UniqueIdentifier("tint_unary_minus");
+          {
+            auto decl = line(&b);
+            if (!EmitTypeAndName(decl, expr_type, fn_name)) {
+              return "";
+            }
+            decl << "(const ";
+            if (!EmitType(decl, expr_type, "")) {
+              return "";
+            }
+            decl << " v) {";
+          }
+
+          {
+            ScopedIndent si(&b);
+            const auto largest_negative_value =
+                std::to_string(std::numeric_limits<int32_t>::min());
+            line(&b) << "return select(-v, v, v == " << largest_negative_value
+                     << ");";
+          }
+          line(&b) << "}";
+          line(&b);
+          return fn_name;
+        });
+
+    out << fn << "(";
+    if (!EmitExpression(out, expr->expr())) {
+      return false;
+    }
+    out << ")";
+    return true;
+  }
+
   switch (expr->op()) {
     case ast::UnaryOp::kAddressOf:
       out << "&";

src/writer/msl/generator_impl.h

@@ -356,6 +356,7 @@ class GeneratorImpl : public TextGenerator {
   bool has_invariant_ = false;
 
   std::unordered_map<const sem::Intrinsic*, std::string> intrinsics_;
+  std::unordered_map<const sem::Type*, std::string> unary_minus_funcs_;
 };
 
 }  // namespace msl

src/writer/msl/generator_impl_binary_test.cc

@@ -74,6 +74,85 @@ INSTANTIATE_TEST_SUITE_P(
         BinaryData{"(left / right)", ast::BinaryOp::kDivide},
         BinaryData{"(left % right)", ast::BinaryOp::kModulo}));
 
+using MslBinaryTest_SignedOverflowDefinedBehaviour =
+    TestParamHelper<BinaryData>;
+TEST_P(MslBinaryTest_SignedOverflowDefinedBehaviour, Emit) {
+  auto params = GetParam();
+
+  auto* a_type = ty.i32();
+  auto* b_type = (params.op == ast::BinaryOp::kShiftLeft ||
+                  params.op == ast::BinaryOp::kShiftRight)
+                     ? static_cast<ast::Type*>(ty.u32())
+                     : ty.i32();
+
+  auto* a = Var("a", a_type);
+  auto* b = Var("b", b_type);
+
+  auto* expr = create<ast::BinaryExpression>(params.op, Expr(a), Expr(b));
+  WrapInFunction(a, b, expr);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, expr)) << gen.error();
+  EXPECT_EQ(out.str(), params.result);
+}
+using Op = ast::BinaryOp;
+constexpr BinaryData signed_overflow_defined_behaviour_cases[] = {
+    {"as_type<int>((as_type<uint>(a) << b))", Op::kShiftLeft},
+    {"(a >> b)", Op::kShiftRight},
+    {"as_type<int>((as_type<uint>(a) + as_type<uint>(b)))", Op::kAdd},
+    {"as_type<int>((as_type<uint>(a) - as_type<uint>(b)))", Op::kSubtract},
+    {"as_type<int>((as_type<uint>(a) * as_type<uint>(b)))", Op::kMultiply}};
+INSTANTIATE_TEST_SUITE_P(
+    MslGeneratorImplTest,
+    MslBinaryTest_SignedOverflowDefinedBehaviour,
+    testing::ValuesIn(signed_overflow_defined_behaviour_cases));
+
+using MslBinaryTest_SignedOverflowDefinedBehaviour_Chained =
+    TestParamHelper<BinaryData>;
+TEST_P(MslBinaryTest_SignedOverflowDefinedBehaviour_Chained, Emit) {
+  auto params = GetParam();
+
+  auto* a_type = ty.i32();
+  auto* b_type = (params.op == ast::BinaryOp::kShiftLeft ||
+                  params.op == ast::BinaryOp::kShiftRight)
+                     ? static_cast<ast::Type*>(ty.u32())
+                     : ty.i32();
+
+  auto* a = Var("a", a_type);
+  auto* b = Var("b", b_type);
+
+  auto* expr1 = create<ast::BinaryExpression>(params.op, Expr(a), Expr(b));
+  auto* expr2 = create<ast::BinaryExpression>(params.op, expr1, Expr(b));
+  WrapInFunction(a, b, expr2);
+
+  GeneratorImpl& gen = Build();
+
+  std::stringstream out;
+  ASSERT_TRUE(gen.EmitExpression(out, expr2)) << gen.error();
+  EXPECT_EQ(out.str(), params.result);
+}
+using Op = ast::BinaryOp;
+constexpr BinaryData signed_overflow_defined_behaviour_chained_cases[] = {
+    {"as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) << b))) << "
+     "b))",
+     Op::kShiftLeft},
+    {"((a >> b) >> b)", Op::kShiftRight},
+    {"as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) + "
+     "as_type<uint>(b)))) + as_type<uint>(b)))",
+     Op::kAdd},
+    {"as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) - "
+     "as_type<uint>(b)))) - as_type<uint>(b)))",
+     Op::kSubtract},
+    {"as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) * "
+     "as_type<uint>(b)))) * as_type<uint>(b)))",
+     Op::kMultiply}};
+INSTANTIATE_TEST_SUITE_P(
+    MslGeneratorImplTest,
+    MslBinaryTest_SignedOverflowDefinedBehaviour_Chained,
+    testing::ValuesIn(signed_overflow_defined_behaviour_chained_cases));
+
 TEST_F(MslBinaryTest, ModF32) {
   auto* left = Var("left", ty.f32());
   auto* right = Var("right", ty.f32());

src/writer/msl/generator_impl_intrinsic_test.cc

@@ -360,7 +360,7 @@ TEST_F(MslGeneratorImplTest, Ignore) {
 
 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 func() {

src/writer/msl/generator_impl_loop_test.cc

@@ -259,7 +259,9 @@ TEST_F(MslGeneratorImplTest, Emit_ForLoopWithSimpleCont) {
   gen.increment_indent();
 
   ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(  for(; ; i = (i + 1)) {
+  EXPECT_EQ(
+      gen.result(),
+      R"(  for(; ; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     a_statement();
   }
 )");
@@ -310,7 +312,9 @@ TEST_F(MslGeneratorImplTest, Emit_ForLoopWithSimpleInitCondCont) {
   gen.increment_indent();
 
   ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(  for(int i = 0; true; i = (i + 1)) {
+  EXPECT_EQ(
+      gen.result(),
+      R"(  for(int i = 0; true; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     a_statement();
   }
 )");

src/writer/msl/generator_impl_unary_op_test.cc

@@ -85,7 +85,7 @@ TEST_F(MslUnaryOpTest, Negation) {
 
   std::stringstream out;
   ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
-  EXPECT_EQ(out.str(), "-(expr)");
+  EXPECT_EQ(out.str(), "tint_unary_minus(expr)");
 }
 
 TEST_F(MslUnaryOpTest, NegationOfIntMin) {
@@ -97,7 +97,7 @@ TEST_F(MslUnaryOpTest, NegationOfIntMin) {
 
   std::stringstream out;
   ASSERT_TRUE(gen.EmitExpression(out, op)) << gen.error();
-  EXPECT_EQ(out.str(), "-((-2147483647 - 1))");
+  EXPECT_EQ(out.str(), "tint_unary_minus((-2147483647 - 1))");
 }
 
 }  // namespace

test/array/type_constructor.wgsl.expected.msl

@@ -18,7 +18,7 @@ kernel void tint_symbol() {
   int const x = 42;
   tint_array_wrapper const empty = {.arr={}};
   tint_array_wrapper const nonempty = {.arr={1, 2, 3, 4}};
-  tint_array_wrapper const nonempty_with_expr = {.arr={1, x, (x + 1), nonempty.arr[3]}};
+  tint_array_wrapper const nonempty_with_expr = {.arr={1, x, as_type<int>((as_type<uint>(x) + as_type<uint>(1))), nonempty.arr[3]}};
   tint_array_wrapper_1 const nested_empty = {.arr={}};
   tint_array_wrapper const tint_symbol_1 = {.arr={1, 2, 3, 4}};
   tint_array_wrapper const tint_symbol_2 = {.arr={5, 6, 7, 8}};
@@ -29,15 +29,15 @@ kernel void tint_symbol() {
   tint_array_wrapper const tint_symbol_7 = {.arr={21, 22, 23, 24}};
   tint_array_wrapper_2 const tint_symbol_8 = {.arr={tint_symbol_5, tint_symbol_6, tint_symbol_7}};
   tint_array_wrapper_1 const nested_nonempty = {.arr={tint_symbol_4, tint_symbol_8}};
-  tint_array_wrapper const tint_symbol_9 = {.arr={1, 2, x, (x + 1)}};
+  tint_array_wrapper const tint_symbol_9 = {.arr={1, 2, x, as_type<int>((as_type<uint>(x) + as_type<uint>(1)))}};
-  tint_array_wrapper const tint_symbol_10 = {.arr={5, 6, nonempty.arr[2], (nonempty.arr[3] + 1)}};
+  tint_array_wrapper const tint_symbol_10 = {.arr={5, 6, nonempty.arr[2], as_type<int>((as_type<uint>(nonempty.arr[3]) + as_type<uint>(1)))}};
   tint_array_wrapper_2 const tint_symbol_11 = {.arr={tint_symbol_9, tint_symbol_10, nonempty}};
   tint_array_wrapper_1 const nested_nonempty_with_expr = {.arr={tint_symbol_11, nested_nonempty.arr[1]}};
   tint_array_wrapper const tint_symbol_12 = {.arr={}};
   int const subexpr_empty = tint_symbol_12.arr[1];
   tint_array_wrapper const tint_symbol_13 = {.arr={1, 2, 3, 4}};
   int const subexpr_nonempty = tint_symbol_13.arr[2];
-  tint_array_wrapper const tint_symbol_14 = {.arr={1, x, (x + 1), nonempty.arr[3]}};
+  tint_array_wrapper const tint_symbol_14 = {.arr={1, x, as_type<int>((as_type<uint>(x) + as_type<uint>(1))), nonempty.arr[3]}};
   int const subexpr_nonempty_with_expr = tint_symbol_14.arr[2];
   tint_array_wrapper_3 const tint_symbol_15 = {.arr={}};
   tint_array_wrapper const subexpr_nested_empty = tint_symbol_15.arr[1];
@@ -45,7 +45,7 @@ kernel void tint_symbol() {
   tint_array_wrapper const tint_symbol_17 = {.arr={5, 6, 7, 8}};
   tint_array_wrapper_3 const tint_symbol_18 = {.arr={tint_symbol_16, tint_symbol_17}};
   tint_array_wrapper const subexpr_nested_nonempty = tint_symbol_18.arr[1];
-  tint_array_wrapper const tint_symbol_19 = {.arr={1, x, (x + 1), nonempty.arr[3]}};
+  tint_array_wrapper const tint_symbol_19 = {.arr={1, x, as_type<int>((as_type<uint>(x) + as_type<uint>(1))), nonempty.arr[3]}};
   tint_array_wrapper_3 const tint_symbol_20 = {.arr={tint_symbol_19, nested_nonempty.arr[1].arr[2]}};
   tint_array_wrapper const subexpr_nested_nonempty_with_expr = tint_symbol_20.arr[1];
   return;

test/bug/fxc/vector_assignment_in_loop/loop_call_with_loop.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void foo(thread float2* const tint_symbol_1, thread int3* const tint_symbol_2, thread uint4* const tint_symbol_3, thread bool2* const tint_symbol_4) {
-  for(int i = 0; (i < 2); i = (i + 1)) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     (*(tint_symbol_1))[i] = 1.0f;
     (*(tint_symbol_2))[i] = 1;
     (*(tint_symbol_3))[i] = 1u;
@@ -15,7 +15,7 @@ kernel void tint_symbol() {
   thread int3 tint_symbol_6 = 0;
   thread uint4 tint_symbol_7 = 0u;
   thread bool2 tint_symbol_8 = false;
-  for(int i = 0; (i < 2); i = (i + 1)) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
   }
   return;

test/bug/fxc/vector_assignment_in_loop/loop_call_with_no_loop.wgsl.expected.msl

@@ -14,7 +14,7 @@ kernel void tint_symbol() {
   thread int3 tint_symbol_6 = 0;
   thread uint4 tint_symbol_7 = 0u;
   thread bool2 tint_symbol_8 = false;
-  for(int i = 0; (i < 2); i = (i + 1)) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     foo(&(tint_symbol_5), &(tint_symbol_6), &(tint_symbol_7), &(tint_symbol_8));
   }
   return;

test/bug/fxc/vector_assignment_in_loop/loop_types_all.wgsl.expected.msl

@@ -14,7 +14,7 @@ kernel void tint_symbol() {
   bool2 v2b = false;
   bool3 v3b = false;
   bool4 v4b = false;
-  for(int i = 0; (i < 2); i = (i + 1)) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     v2f[i] = 1.0f;
     v3f[i] = 1.0f;
     v4f[i] = 1.0f;

test/bug/fxc/vector_assignment_in_loop/loop_types_repeated.wgsl.expected.msl

@@ -10,7 +10,7 @@ kernel void tint_symbol() {
   uint4 v4u_2 = 0u;
   bool2 v2b = false;
   bool2 v2b_2 = false;
-  for(int i = 0; (i < 2); i = (i + 1)) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     v2f[i] = 1.0f;
     v3i[i] = 1;
     v4u[i] = 1u;

test/bug/fxc/vector_assignment_in_loop/loop_types_some.wgsl.expected.msl

@@ -14,7 +14,7 @@ kernel void tint_symbol() {
   bool2 v2b = false;
   bool3 v3b = false;
   bool4 v4b = false;
-  for(int i = 0; (i < 2); i = (i + 1)) {
+  for(int i = 0; (i < 2); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     v2f[i] = 1.0f;
     v2i[i] = 1;
     v2u[i] = 1u;

test/bug/tint/534.wgsl.expected.msl

@@ -20,7 +20,7 @@ kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_2 [[texture
   int2 dstTexCoord = int2(GlobalInvocationID.xy);
   int2 srcTexCoord = dstTexCoord;
   if ((uniforms.dstTextureFlipY == 1u)) {
-    srcTexCoord.y = ((size.y - dstTexCoord.y) - 1);
+    srcTexCoord.y = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(size.y) - as_type<uint>(dstTexCoord.y)))) - as_type<uint>(1)));
   }
   float4 srcColor = tint_symbol_2.read(uint2(srcTexCoord), 0);
   float4 dstColor = tint_symbol_3.read(uint2(dstTexCoord), 0);

test/bug/tint/749.spvasm.expected.msl

@@ -152,7 +152,7 @@ int performPartition_i1_i1_(thread int* const l, thread int* const h, thread Qui
   int const x_959 = *(l);
   *(l) = 0;
   *(l) = x_959;
-  i_1 = (x_45 - as_type<int>(1u));
+  i_1 = as_type<int>((as_type<uint>(x_45) - as_type<uint>(as_type<int>(1u))));
   int const x_49 = *(l);
   float3 const x_536 = float3(x_534.x, x_534.z, x_535.x);
   j_1 = 10;
@@ -192,7 +192,7 @@ int performPartition_i1_i1_(thread int* const l, thread int* const h, thread Qui
     int const x_968 = param;
     param = 0;
     param = x_968;
-    if ((x_55 <= (x_56 - as_type<int>(1u)))) {
+    if ((x_55 <= as_type<int>((as_type<uint>(x_56) - as_type<uint>(as_type<int>(1u)))))) {
     } else {
       break;
     }
@@ -242,7 +242,7 @@ int performPartition_i1_i1_(thread int* const l, thread int* const h, thread Qui
       int const x_979 = param;
       param = 0;
       param = x_979;
-      i_1 = (x_67 + as_type<int>(1u));
+      i_1 = as_type<int>((as_type<uint>(x_67) + as_type<uint>(as_type<int>(1u))));
       int const x_980 = *(l);
       *(l) = 0;
       *(l) = x_980;
@@ -290,7 +290,7 @@ int performPartition_i1_i1_(thread int* const l, thread int* const h, thread Qui
       int const x_990 = param;
       param = 0;
       param = x_990;
-      j_1 = (1 + x_74);
+      j_1 = as_type<int>((as_type<uint>(1) + as_type<uint>(x_74)));
       int const x_991 = param_1;
       param_1 = 0;
       param_1 = x_991;
@@ -313,7 +313,7 @@ int performPartition_i1_i1_(thread int* const l, thread int* const h, thread Qui
   int const x_995 = *(h);
   *(h) = 0;
   *(h) = x_995;
-  i_1 = (1 + x_76);
+  i_1 = as_type<int>((as_type<uint>(1) + as_type<uint>(x_76)));
   int const x_996 = param_1;
   param_1 = 0;
   param_1 = x_996;
@@ -392,7 +392,7 @@ void quicksort_(thread QuicksortObject* const tint_symbol_85) {
   int const x_1011 = p;
   p = 0;
   p = x_1011;
-  int const x_94 = (x_93 + as_type<int>(1u));
+  int const x_94 = as_type<int>((as_type<uint>(x_93) + as_type<uint>(as_type<int>(1u))));
   int const x_1012 = top;
   top = 0;
   top = x_1012;
@@ -436,7 +436,7 @@ void quicksort_(thread QuicksortObject* const tint_symbol_85) {
   int const x_1021 = stack.arr[x_96_save];
   stack.arr[x_96_save] = 0;
   stack.arr[x_96_save] = x_1021;
-  int const x_98 = (x_97 + 1);
+  int const x_98 = as_type<int>((as_type<uint>(x_97) + as_type<uint>(1)));
   int const x_1022 = stack.arr[x_96_save];
   stack.arr[x_96_save] = 0;
   stack.arr[x_96_save] = x_1022;
@@ -502,7 +502,7 @@ void quicksort_(thread QuicksortObject* const tint_symbol_85) {
     int const x_1035 = p;
     p = 0;
     p = x_1035;
-    top = (x_108 - as_type<int>(1u));
+    top = as_type<int>((as_type<uint>(x_108) - as_type<uint>(as_type<int>(1u))));
     int const x_1036 = p;
     p = 0;
     p = x_1036;
@@ -536,7 +536,7 @@ void quicksort_(thread QuicksortObject* const tint_symbol_85) {
     stack.arr[x_100_save] = 0;
     stack.arr[x_100_save] = x_1043;
     float2 const x_573 = float2(float3(1.0f, 2.0f, 3.0f).y, float3(1.0f, 2.0f, 3.0f).z);
-    top = (x_112 - 1);
+    top = as_type<int>((as_type<uint>(x_112) - as_type<uint>(1)));
     int const x_1044 = param_5;
     param_5 = 0;
     param_5 = x_1044;
@@ -604,7 +604,7 @@ void quicksort_(thread QuicksortObject* const tint_symbol_85) {
     int const x_1059 = stack.arr[x_100_save];
     stack.arr[x_100_save] = 0;
     stack.arr[x_100_save] = x_1059;
-    if (((x_122 - as_type<int>(1u)) > x_124)) {
+    if ((as_type<int>((as_type<uint>(x_122) - as_type<uint>(as_type<int>(1u)))) > x_124)) {
       int const x_1060 = param_4;
       param_4 = 0;
       param_4 = x_1060;
@@ -627,7 +627,7 @@ void quicksort_(thread QuicksortObject* const tint_symbol_85) {
       int const x_1064 = param_5;
       param_5 = 0;
       param_5 = x_1064;
-      int const x_131_save = (1 + x_128);
+      int const x_131_save = as_type<int>((as_type<uint>(1) + as_type<uint>(x_128)));
       int const x_1065 = stack.arr[x_110_save];
       stack.arr[x_110_save] = 0;
       stack.arr[x_110_save] = x_1065;
@@ -666,7 +666,7 @@ void quicksort_(thread QuicksortObject* const tint_symbol_85) {
       int const x_1073 = stack.arr[x_114_save];
       stack.arr[x_114_save] = 0;
       stack.arr[x_114_save] = x_1073;
-      stack.arr[x_136_save] = (x_134 - as_type<int>(1u));
+      stack.arr[x_136_save] = as_type<int>((as_type<uint>(x_134) - as_type<uint>(as_type<int>(1u))));
       int const x_1074 = stack.arr[x_96_save];
       stack.arr[x_96_save] = 0;
       stack.arr[x_96_save] = x_1074;
@@ -720,7 +720,7 @@ void quicksort_(thread QuicksortObject* const tint_symbol_85) {
       stack.arr[x_114_save] = 0;
       stack.arr[x_114_save] = x_1086;
       float3 const x_597 = float3(x_562.y, x_560.y, x_560.y);
-      int const x_144 = (x_143 + 1);
+      int const x_144 = as_type<int>((as_type<uint>(x_143) + as_type<uint>(1)));
       int const x_1087 = param_5;
       param_5 = 0;
       param_5 = x_1087;
@@ -759,7 +759,7 @@ void quicksort_(thread QuicksortObject* const tint_symbol_85) {
       tint_array_wrapper const tint_symbol_35 = {.arr={0, 0, 0, 0, 0, 0, 0, 0, 0, 0}};
       stack = tint_symbol_35;
       stack = x_1095;
-      int const x_149 = (x_148 + as_type<int>(1u));
+      int const x_149 = as_type<int>((as_type<uint>(x_148) + as_type<uint>(as_type<int>(1u))));
       int const x_1096 = stack.arr[x_147_save];
       stack.arr[x_147_save] = 0;
       stack.arr[x_147_save] = x_1096;

test/bug/tint/942.wgsl.expected.msl

@@ -25,10 +25,10 @@ kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_3 [[texture
   threadgroup_barrier(mem_flags::mem_threadgroup);
   uint const filterOffset = ((params.filterDim - 1u) / 2u);
   int2 const dims = int2(tint_symbol_3.get_width(0), tint_symbol_3.get_height(0));
-  int2 const baseIndex = (int2(((WorkGroupID.xy * uint2(params.blockDim, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u)))) - int2(int(filterOffset), 0));
+  int2 const baseIndex = as_type<int2>((as_type<uint2>(int2(((WorkGroupID.xy * uint2(params.blockDim, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u))))) - as_type<uint2>(int2(int(filterOffset), 0))));
   for(uint r = 0u; (r < 4u); r = (r + 1u)) {
     for(uint c = 0u; (c < 4u); c = (c + 1u)) {
-      int2 loadIndex = (baseIndex + int2(int(c), int(r)));
+      int2 loadIndex = as_type<int2>((as_type<uint2>(baseIndex) + as_type<uint2>(int2(int(c), int(r)))));
       if ((flip.value != 0u)) {
         loadIndex = loadIndex.yx;
       }
@@ -38,7 +38,7 @@ kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_3 [[texture
   threadgroup_barrier(mem_flags::mem_threadgroup);
   for(uint r = 0u; (r < 4u); r = (r + 1u)) {
     for(uint c = 0u; (c < 4u); c = (c + 1u)) {
-      int2 writeIndex = (baseIndex + int2(int(c), int(r)));
+      int2 writeIndex = as_type<int2>((as_type<uint2>(baseIndex) + as_type<uint2>(int2(int(c), int(r)))));
       if ((flip.value != 0u)) {
         writeIndex = writeIndex.yx;
       }

test/bug/tint/943.spvasm.expected.msl

@@ -60,7 +60,7 @@ float mm_readA_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, thread
   float x_430 = 0.0f;
   int const x_417 = x_48.aShape.y;
   int const x_419 = x_48.aShape.z;
-  batchASize = (x_417 * x_419);
+  batchASize = as_type<int>((as_type<uint>(x_417) * as_type<uint>(x_419)));
   int const x_421 = *(row);
   int const x_422 = *(col);
   int const x_424 = *(tint_symbol_3);
@@ -74,7 +74,7 @@ float mm_readA_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, thread
     int const x_441 = *(row);
     int const x_442 = *(tint_symbol_4);
     int const x_445 = *(col);
-    float const x_448 = x_165.A[(((x_438 * x_439) + (x_441 * x_442)) + x_445)];
+    float const x_448 = x_165.A[as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(x_438) * as_type<uint>(x_439)))) + as_type<uint>(as_type<int>((as_type<uint>(x_441) * as_type<uint>(x_442))))))) + as_type<uint>(x_445)))];
     x_430 = x_448;
   } else {
     x_430 = 0.0f;
@@ -90,7 +90,7 @@ float mm_readB_i1_i1_(constant Uniforms& x_48, const device ssbB& x_185, thread
   float x_468 = 0.0f;
   int const x_455 = x_48.bShape.y;
   int const x_457 = x_48.bShape.z;
-  batchBSize = (x_455 * x_457);
+  batchBSize = as_type<int>((as_type<uint>(x_455) * as_type<uint>(x_457)));
   int const x_459 = *(row_1);
   int const x_460 = *(col_1);
   int const x_462 = *(tint_symbol_6);
@@ -104,7 +104,7 @@ float mm_readB_i1_i1_(constant Uniforms& x_48, const device ssbB& x_185, thread
     int const x_478 = *(row_1);
     int const x_479 = *(tint_symbol_7);
     int const x_482 = *(col_1);
-    float const x_485 = x_185.B[(((x_475 * x_476) + (x_478 * x_479)) + x_482)];
+    float const x_485 = x_185.B[as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(x_475) * as_type<uint>(x_476)))) + as_type<uint>(as_type<int>((as_type<uint>(x_478) * as_type<uint>(x_479))))))) + as_type<uint>(x_482)))];
     x_468 = x_485;
   } else {
     x_468 = 0.0f;
@@ -204,15 +204,15 @@ void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, cons
   int param_8 = 0;
   float param_9 = 0.0f;
   uint const x_132 = (*(tint_symbol_10)).y;
-  tileRow = (as_type<int>(x_132) * 1);
+  tileRow = as_type<int>((as_type<uint>(as_type<int>(x_132)) * as_type<uint>(1)));
   uint const x_137 = (*(tint_symbol_10)).x;
-  tileCol = (as_type<int>(x_137) * 1);
+  tileCol = as_type<int>((as_type<uint>(as_type<int>(x_137)) * as_type<uint>(1)));
   uint const x_143 = (*(tint_symbol_11)).y;
-  globalRow = (as_type<int>(x_143) * 1);
+  globalRow = as_type<int>((as_type<uint>(as_type<int>(x_143)) * as_type<uint>(1)));
   uint const x_148 = (*(tint_symbol_11)).x;
-  globalCol = (as_type<int>(x_148) * 1);
+  globalCol = as_type<int>((as_type<uint>(as_type<int>(x_148)) * as_type<uint>(1)));
   int const x_152 = *(dimInner);
-  numTiles = (((x_152 - 1) / 64) + 1);
+  numTiles = as_type<int>((as_type<uint>((as_type<int>((as_type<uint>(x_152) - as_type<uint>(1))) / 64)) + as_type<uint>(1)));
   innerRow = 0;
   while (true) {
     int const x_163 = innerRow;
@@ -232,18 +232,18 @@ void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, cons
       acc.arr[x_177].arr[x_178] = 0.0f;
       {
         int const x_181 = innerCol;
-        innerCol = (x_181 + 1);
+        innerCol = as_type<int>((as_type<uint>(x_181) + as_type<uint>(1)));
       }
     }
     {
       int const x_183 = innerRow;
-      innerRow = (x_183 + 1);
+      innerRow = as_type<int>((as_type<uint>(x_183) + as_type<uint>(1)));
     }
   }
   uint const x_187 = (*(tint_symbol_10)).x;
-  tileColA = (as_type<int>(x_187) * 64);
+  tileColA = as_type<int>((as_type<uint>(as_type<int>(x_187)) * as_type<uint>(64)));
   uint const x_192 = (*(tint_symbol_10)).y;
-  tileRowB = (as_type<int>(x_192) * 1);
+  tileRowB = as_type<int>((as_type<uint>(as_type<int>(x_192)) * as_type<uint>(1)));
   t = 0;
   while (true) {
     int const x_201 = t;
@@ -268,28 +268,28 @@ void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, cons
         }
         int const x_221 = tileRow;
         int const x_222 = innerRow_1;
-        inputRow = (x_221 + x_222);
+        inputRow = as_type<int>((as_type<uint>(x_221) + as_type<uint>(x_222)));
         int const x_225 = tileColA;
         int const x_226 = innerCol_1;
-        inputCol = (x_225 + x_226);
+        inputCol = as_type<int>((as_type<uint>(x_225) + as_type<uint>(x_226)));
         int const x_233 = inputRow;
         int const x_234 = inputCol;
         int const x_235 = globalRow;
         int const x_236 = innerRow_1;
         int const x_238 = t;
         int const x_240 = inputCol;
-        param_3 = (x_235 + x_236);
+        param_3 = as_type<int>((as_type<uint>(x_235) + as_type<uint>(x_236)));
-        param_4 = ((x_238 * 64) + x_240);
+        param_4 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(x_238) * as_type<uint>(64)))) + as_type<uint>(x_240)));
         float const x_244 = mm_readA_i1_i1_(x_48, x_165, &(param_3), &(param_4), tint_symbol_12, tint_symbol_13, tint_symbol_14);
         (*(tint_symbol_15)).arr[x_233].arr[x_234] = x_244;
         {
           int const x_247 = innerCol_1;
-          innerCol_1 = (x_247 + 1);
+          innerCol_1 = as_type<int>((as_type<uint>(x_247) + as_type<uint>(1)));
         }
       }
       {
         int const x_249 = innerRow_1;
-        innerRow_1 = (x_249 + 1);
+        innerRow_1 = as_type<int>((as_type<uint>(x_249) + as_type<uint>(1)));
       }
     }
     innerRow_2 = 0;
@@ -308,28 +308,28 @@ void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, cons
         }
         int const x_268 = tileRowB;
         int const x_269 = innerRow_2;
-        inputRow_1 = (x_268 + x_269);
+        inputRow_1 = as_type<int>((as_type<uint>(x_268) + as_type<uint>(x_269)));
         int const x_272 = tileCol;
         int const x_273 = innerCol_2;
-        inputCol_1 = (x_272 + x_273);
+        inputCol_1 = as_type<int>((as_type<uint>(x_272) + as_type<uint>(x_273)));
         int const x_278 = inputRow_1;
         int const x_279 = inputCol_1;
         int const x_280 = t;
         int const x_282 = inputRow_1;
         int const x_284 = globalCol;
         int const x_285 = innerCol_2;
-        param_5 = ((x_280 * 64) + x_282);
+        param_5 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(x_280) * as_type<uint>(64)))) + as_type<uint>(x_282)));
-        param_6 = (x_284 + x_285);
+        param_6 = as_type<int>((as_type<uint>(x_284) + as_type<uint>(x_285)));
         float const x_289 = mm_readB_i1_i1_(x_48, x_185, &(param_5), &(param_6), tint_symbol_13, tint_symbol_16, tint_symbol_14);
         (*(tint_symbol_17)).arr[x_278].arr[x_279] = x_289;
         {
           int const x_291 = innerCol_2;
-          innerCol_2 = (x_291 + 1);
+          innerCol_2 = as_type<int>((as_type<uint>(x_291) + as_type<uint>(1)));
         }
       }
       {
         int const x_293 = innerRow_2;
-        innerRow_2 = (x_293 + 1);
+        innerRow_2 = as_type<int>((as_type<uint>(x_293) + as_type<uint>(1)));
       }
     }
     threadgroup_barrier(mem_flags::mem_threadgroup);
@@ -351,11 +351,11 @@ void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, cons
         int const x_315 = k;
         int const x_316 = tileCol;
         int const x_317 = inner;
-        float const x_320 = (*(tint_symbol_17)).arr[x_315].arr[(x_316 + x_317)];
+        float const x_320 = (*(tint_symbol_17)).arr[x_315].arr[as_type<int>((as_type<uint>(x_316) + as_type<uint>(x_317)))];
         BCached.arr[x_314] = x_320;
         {
           int const x_322 = inner;
-          inner = (x_322 + 1);
+          inner = as_type<int>((as_type<uint>(x_322) + as_type<uint>(1)));
         }
       }
       innerRow_3 = 0;
@@ -368,7 +368,7 @@ void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, cons
         int const x_333 = tileRow;
         int const x_334 = innerRow_3;
         int const x_336 = k;
-        float const x_338 = (*(tint_symbol_15)).arr[(x_333 + x_334)].arr[x_336];
+        float const x_338 = (*(tint_symbol_15)).arr[as_type<int>((as_type<uint>(x_333) + as_type<uint>(x_334)))].arr[x_336];
         ACached = x_338;
         innerCol_3 = 0;
         while (true) {
@@ -386,23 +386,23 @@ void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, cons
           acc.arr[x_347].arr[x_348] = (x_355 + (x_349 * x_352));
           {
             int const x_358 = innerCol_3;
-            innerCol_3 = (x_358 + 1);
+            innerCol_3 = as_type<int>((as_type<uint>(x_358) + as_type<uint>(1)));
           }
         }
         {
           int const x_360 = innerRow_3;
-          innerRow_3 = (x_360 + 1);
+          innerRow_3 = as_type<int>((as_type<uint>(x_360) + as_type<uint>(1)));
         }
       }
       {
         int const x_362 = k;
-        k = (x_362 + 1);
+        k = as_type<int>((as_type<uint>(x_362) + as_type<uint>(1)));
       }
     }
     threadgroup_barrier(mem_flags::mem_threadgroup);
     {
       int const x_364 = t;
-      t = (x_364 + 1);
+      t = as_type<int>((as_type<uint>(x_364) + as_type<uint>(1)));
     }
   }
   innerRow_4 = 0;
@@ -424,13 +424,13 @@ void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, cons
       int const x_382 = globalCol;
       int const x_383 = innerCol_4;
       int const x_385 = *(dimBOuter);
-      bool const x_386 = ((x_382 + x_383) < x_385);
+      bool const x_386 = (as_type<int>((as_type<uint>(x_382) + as_type<uint>(x_383))) < x_385);
       x_394_phi = x_386;
       if (x_386) {
         int const x_389 = globalRow;
         int const x_390 = innerRow_4;
         int const x_392 = *(dimAOuter);
-        x_393 = ((x_389 + x_390) < x_392);
+        x_393 = (as_type<int>((as_type<uint>(x_389) + as_type<uint>(x_390))) < x_392);
         x_394_phi = x_393;
       }
       bool const x_394 = x_394_phi;
@@ -441,20 +441,20 @@ void mm_matMul_i1_i1_i1_(constant Uniforms& x_48, const device ssbA& x_165, cons
         int const x_401 = innerCol_4;
         int const x_403 = innerRow_4;
         int const x_404 = innerCol_4;
-        param_7 = (x_397 + x_398);
+        param_7 = as_type<int>((as_type<uint>(x_397) + as_type<uint>(x_398)));
-        param_8 = (x_400 + x_401);
+        param_8 = as_type<int>((as_type<uint>(x_400) + as_type<uint>(x_401)));
         float const x_409 = acc.arr[x_403].arr[x_404];
         param_9 = x_409;
         mm_write_i1_i1_f1_(x_48, x_54, &(param_7), &(param_8), &(param_9), tint_symbol_14);
       }
       {
         int const x_411 = innerCol_4;
-        innerCol_4 = (x_411 + 1);
+        innerCol_4 = as_type<int>((as_type<uint>(x_411) + as_type<uint>(1)));
       }
     }
     {
       int const x_413 = innerRow_4;
-      innerRow_4 = (x_413 + 1);
+      innerRow_4 = as_type<int>((as_type<uint>(x_413) + as_type<uint>(1)));
     }
   }
   return;

test/bug/tint/948.wgsl.expected.msl

@@ -180,7 +180,7 @@ void main_1(constant LeftOver& x_20, thread float2* const tint_symbol_8, texture
     }
     {
       int const x_304 = i;
-      i = (x_304 + 1);
+      i = as_type<int>((as_type<uint>(x_304) + as_type<uint>(1)));
     }
   }
   float3 const x_310 = x_20.colorMul;

test/bug/tint/949.wgsl.expected.msl

@@ -339,7 +339,7 @@ void main_1(constant LeftOver& x_269, constant Light0& light0, thread float* con
     }
     {
       int const x_441 = i;
-      i = (x_441 + 1);
+      i = as_type<int>((as_type<uint>(x_441) + as_type<uint>(1)));
     }
   }
   float2 const x_444 = vCurrOffset;

test/expressions/binary/add/scalar-scalar/i32.wgsl.expected.msl

@@ -4,7 +4,7 @@ using namespace metal;
 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)));
   return;
 }
 

test/expressions/binary/add/scalar-vec3/i32.wgsl.expected.msl

@@ -4,7 +4,7 @@ using namespace metal;
 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)));
   return;
 }
 

test/expressions/binary/add/vec3-scalar/i32.wgsl.expected.msl

@@ -4,7 +4,7 @@ using namespace metal;
 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)));
   return;
 }
 

test/expressions/binary/add/vec3-vec3/i32.wgsl.expected.msl

@@ -4,7 +4,7 @@ using namespace metal;
 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)));
   return;
 }
 

test/expressions/binary/left-shift/scalar-scalar/i32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = 1;
+    let b = 2u;
+    let r : i32 = a << b;
+}

test/expressions/binary/left-shift/scalar-scalar/i32.wgsl.expected.hlsl

@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const int r = (1 << 2u);
+  return;
+}

test/expressions/binary/left-shift/scalar-scalar/i32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  int const a = 1;
+  uint const b = 2u;
+  int const r = as_type<int>((as_type<uint>(a) << b));
+  return;
+}
+

test/expressions/binary/left-shift/scalar-scalar/i32.wgsl.expected.spvasm

@@ -0,0 +1,21 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 10
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %f "f"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+       %uint = OpTypeInt 32 0
+     %uint_2 = OpConstant %uint 2
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %9 = OpShiftLeftLogical %int %int_1 %uint_2
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/left-shift/scalar-scalar/i32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  let a = 1;
+  let b = 2u;
+  let r : i32 = (a << b);
+}

test/expressions/binary/left-shift/scalar-scalar/u32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = 1u;
+    let b = 2u;
+    let r : u32 = a << b;
+}

test/expressions/binary/left-shift/scalar-scalar/u32.wgsl.expected.hlsl

@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const uint r = (1u << 2u);
+  return;
+}

test/expressions/binary/left-shift/scalar-scalar/u32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  uint const a = 1u;
+  uint const b = 2u;
+  uint const r = (a << b);
+  return;
+}
+

test/expressions/binary/left-shift/scalar-scalar/u32.wgsl.expected.spvasm

@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 9
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %f "f"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %8 = OpShiftLeftLogical %uint %uint_1 %uint_2
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/left-shift/scalar-scalar/u32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  let a = 1u;
+  let b = 2u;
+  let r : u32 = (a << b);
+}

test/expressions/binary/left-shift/vector-vector/i32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = vec3<i32>(1, 2, 3);
+    let b = vec3<u32>(4u, 5u, 6u);
+    let r : vec3<i32> = a << b;
+}

test/expressions/binary/left-shift/vector-vector/i32.wgsl.expected.hlsl

@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const int3 a = int3(1, 2, 3);
+  const uint3 b = uint3(4u, 5u, 6u);
+  const int3 r = (a << b);
+  return;
+}

test/expressions/binary/left-shift/vector-vector/i32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  int3 const a = int3(1, 2, 3);
+  uint3 const b = uint3(4u, 5u, 6u);
+  int3 const r = as_type<int3>((as_type<uint3>(a) << b));
+  return;
+}
+

test/expressions/binary/left-shift/vector-vector/i32.wgsl.expected.spvasm

@@ -0,0 +1,29 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %f "f"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+      %int_1 = OpConstant %int 1
+      %int_2 = OpConstant %int 2
+      %int_3 = OpConstant %int 3
+         %10 = OpConstantComposite %v3int %int_1 %int_2 %int_3
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+     %uint_4 = OpConstant %uint 4
+     %uint_5 = OpConstant %uint 5
+     %uint_6 = OpConstant %uint 6
+         %16 = OpConstantComposite %v3uint %uint_4 %uint_5 %uint_6
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+         %17 = OpShiftLeftLogical %v3int %10 %16
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/left-shift/vector-vector/i32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  let a = vec3<i32>(1, 2, 3);
+  let b = vec3<u32>(4u, 5u, 6u);
+  let r : vec3<i32> = (a << b);
+}

test/expressions/binary/left-shift/vector-vector/u32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = vec3<u32>(1u, 2u, 3u);
+    let b = vec3<u32>(4u, 5u, 6u);
+    let r : vec3<u32> = a << b;
+}

test/expressions/binary/left-shift/vector-vector/u32.wgsl.expected.hlsl

@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const uint3 a = uint3(1u, 2u, 3u);
+  const uint3 b = uint3(4u, 5u, 6u);
+  const uint3 r = (a << b);
+  return;
+}

test/expressions/binary/left-shift/vector-vector/u32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  uint3 const a = uint3(1u, 2u, 3u);
+  uint3 const b = uint3(4u, 5u, 6u);
+  uint3 const r = (a << b);
+  return;
+}
+

test/expressions/binary/left-shift/vector-vector/u32.wgsl.expected.spvasm

@@ -0,0 +1,27 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %f "f"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+     %uint_3 = OpConstant %uint 3
+         %10 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
+     %uint_4 = OpConstant %uint 4
+     %uint_5 = OpConstant %uint 5
+     %uint_6 = OpConstant %uint 6
+         %14 = OpConstantComposite %v3uint %uint_4 %uint_5 %uint_6
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+         %15 = OpShiftLeftLogical %v3uint %10 %14
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/left-shift/vector-vector/u32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  let a = vec3<u32>(1u, 2u, 3u);
+  let b = vec3<u32>(4u, 5u, 6u);
+  let r : vec3<u32> = (a << b);
+}

test/expressions/binary/mul/scalar-scalar/i32.wgsl.expected.msl

@@ -4,7 +4,7 @@ using namespace metal;
 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)));
   return;
 }
 

test/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.msl

@@ -4,7 +4,7 @@ using namespace metal;
 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)));
   return;
 }
 

test/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.msl

@@ -4,7 +4,7 @@ using namespace metal;
 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)));
   return;
 }
 

test/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.msl

@@ -4,7 +4,7 @@ using namespace metal;
 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)));
   return;
 }
 

test/expressions/binary/right-shift/scalar-scalar/i32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = 1;
+    let b = 2u;
+    let r : i32 = a >> b;
+}

test/expressions/binary/right-shift/scalar-scalar/i32.wgsl.expected.hlsl

@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const int r = (1 >> 2u);
+  return;
+}

test/expressions/binary/right-shift/scalar-scalar/i32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  int const a = 1;
+  uint const b = 2u;
+  int const r = (a >> b);
+  return;
+}
+

test/expressions/binary/right-shift/scalar-scalar/i32.wgsl.expected.spvasm

@@ -0,0 +1,21 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 10
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %f "f"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %int_1 = OpConstant %int 1
+       %uint = OpTypeInt 32 0
+     %uint_2 = OpConstant %uint 2
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %9 = OpShiftRightArithmetic %int %int_1 %uint_2
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/right-shift/scalar-scalar/i32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  let a = 1;
+  let b = 2u;
+  let r : i32 = (a >> b);
+}

test/expressions/binary/right-shift/scalar-scalar/u32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = 1u;
+    let b = 2u;
+    let r : u32 = a >> b;
+}

test/expressions/binary/right-shift/scalar-scalar/u32.wgsl.expected.hlsl

@@ -0,0 +1,5 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const uint r = (1u >> 2u);
+  return;
+}

test/expressions/binary/right-shift/scalar-scalar/u32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  uint const a = 1u;
+  uint const b = 2u;
+  uint const r = (a >> b);
+  return;
+}
+

test/expressions/binary/right-shift/scalar-scalar/u32.wgsl.expected.spvasm

@@ -0,0 +1,20 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 9
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %f "f"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+          %8 = OpShiftRightLogical %uint %uint_1 %uint_2
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/right-shift/scalar-scalar/u32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  let a = 1u;
+  let b = 2u;
+  let r : u32 = (a >> b);
+}

test/expressions/binary/right-shift/vector-vector/i32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = vec3<i32>(1, 2, 3);
+    let b = vec3<u32>(4u, 5u, 6u);
+    let r : vec3<i32> = a >> b;
+}

test/expressions/binary/right-shift/vector-vector/i32.wgsl.expected.hlsl

@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const int3 a = int3(1, 2, 3);
+  const uint3 b = uint3(4u, 5u, 6u);
+  const int3 r = (a >> b);
+  return;
+}

test/expressions/binary/right-shift/vector-vector/i32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  int3 const a = int3(1, 2, 3);
+  uint3 const b = uint3(4u, 5u, 6u);
+  int3 const r = (a >> b);
+  return;
+}
+

test/expressions/binary/right-shift/vector-vector/i32.wgsl.expected.spvasm

@@ -0,0 +1,29 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %f "f"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+      %int_1 = OpConstant %int 1
+      %int_2 = OpConstant %int 2
+      %int_3 = OpConstant %int 3
+         %10 = OpConstantComposite %v3int %int_1 %int_2 %int_3
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+     %uint_4 = OpConstant %uint 4
+     %uint_5 = OpConstant %uint 5
+     %uint_6 = OpConstant %uint 6
+         %16 = OpConstantComposite %v3uint %uint_4 %uint_5 %uint_6
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+         %17 = OpShiftRightArithmetic %v3int %10 %16
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/right-shift/vector-vector/i32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  let a = vec3<i32>(1, 2, 3);
+  let b = vec3<u32>(4u, 5u, 6u);
+  let r : vec3<i32> = (a >> b);
+}

test/expressions/binary/right-shift/vector-vector/u32.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+    let a = vec3<u32>(1u, 2u, 3u);
+    let b = vec3<u32>(4u, 5u, 6u);
+    let r : vec3<u32> = a >> b;
+}

test/expressions/binary/right-shift/vector-vector/u32.wgsl.expected.hlsl

@@ -0,0 +1,7 @@
+[numthreads(1, 1, 1)]
+void f() {
+  const uint3 a = uint3(1u, 2u, 3u);
+  const uint3 b = uint3(4u, 5u, 6u);
+  const uint3 r = (a >> b);
+  return;
+}

test/expressions/binary/right-shift/vector-vector/u32.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void f() {
+  uint3 const a = uint3(1u, 2u, 3u);
+  uint3 const b = uint3(4u, 5u, 6u);
+  uint3 const r = (a >> b);
+  return;
+}
+

test/expressions/binary/right-shift/vector-vector/u32.wgsl.expected.spvasm

@@ -0,0 +1,27 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %f "f"
+               OpExecutionMode %f LocalSize 1 1 1
+               OpName %f "f"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+     %uint_1 = OpConstant %uint 1
+     %uint_2 = OpConstant %uint 2
+     %uint_3 = OpConstant %uint 3
+         %10 = OpConstantComposite %v3uint %uint_1 %uint_2 %uint_3
+     %uint_4 = OpConstant %uint 4
+     %uint_5 = OpConstant %uint 5
+     %uint_6 = OpConstant %uint 6
+         %14 = OpConstantComposite %v3uint %uint_4 %uint_5 %uint_6
+          %f = OpFunction %void None %1
+          %4 = OpLabel
+         %15 = OpShiftRightLogical %v3uint %10 %14
+               OpReturn
+               OpFunctionEnd

test/expressions/binary/right-shift/vector-vector/u32.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  let a = vec3<u32>(1u, 2u, 3u);
+  let b = vec3<u32>(4u, 5u, 6u);
+  let r : vec3<u32> = (a >> b);
+}

test/expressions/binary/sub/scalar-scalar/i32.wgsl.expected.msl

@@ -4,7 +4,7 @@ using namespace metal;
 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)));
   return;
 }
 

test/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.msl

@@ -4,7 +4,7 @@ using namespace metal;
 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)));
   return;
 }
 

test/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.msl

@@ -4,7 +4,7 @@ using namespace metal;
 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)));
   return;
 }
 

test/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.msl

@@ -4,7 +4,7 @@ using namespace metal;
 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)));
   return;
 }
 

test/expressions/literals/intmin.wgsl

@@ -0,0 +1,7 @@
+fn add_int_min_explicit() -> i32 {
+    var a = -2147483648;
+    var b = a + 1;
+
+    var c = -2147483648 + 1;
+    return c;
+}

test/expressions/literals/intmin.wgsl.expected.hlsl

@@ -0,0 +1,11 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+int add_int_min_explicit() {
+  int a = -2147483648;
+  int b = (a + 1);
+  int c = (-2147483648 + 1);
+  return c;
+}

test/expressions/literals/intmin.wgsl.expected.msl

@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+int add_int_min_explicit() {
+  int a = (-2147483647 - 1);
+  int b = as_type<int>((as_type<uint>(a) + as_type<uint>(1)));
+  int c = as_type<int>((as_type<uint>((-2147483647 - 1)) + as_type<uint>(1)));
+  return c;
+}
+

test/expressions/literals/intmin.wgsl.expected.spvasm

@@ -0,0 +1,40 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 20
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %add_int_min_explicit "add_int_min_explicit"
+               OpName %a "a"
+               OpName %b "b"
+               OpName %c "c"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+          %5 = OpTypeFunction %int
+%int_n2147483648 = OpConstant %int -2147483648
+%_ptr_Function_int = OpTypePointer Function %int
+         %12 = OpConstantNull %int
+      %int_1 = OpConstant %int 1
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+%add_int_min_explicit = OpFunction %int None %5
+          %8 = OpLabel
+          %a = OpVariable %_ptr_Function_int Function %12
+          %b = OpVariable %_ptr_Function_int Function %12
+          %c = OpVariable %_ptr_Function_int Function %12
+               OpStore %a %int_n2147483648
+         %13 = OpLoad %int %a
+         %15 = OpIAdd %int %13 %int_1
+               OpStore %b %15
+         %17 = OpIAdd %int %int_n2147483648 %int_1
+               OpStore %c %17
+         %19 = OpLoad %int %c
+               OpReturnValue %19
+               OpFunctionEnd

test/expressions/literals/intmin.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+fn add_int_min_explicit() -> i32 {
+  var a = -2147483648;
+  var b = (a + 1);
+  var c = (-2147483648 + 1);
+  return c;
+}

test/expressions/splat/expression/i32.wgsl.expected.msl

@@ -2,8 +2,8 @@
 
 using namespace metal;
 void f() {
-  int2 v2 = int2((1 + 2));
+  int2 v2 = int2(as_type<int>((as_type<uint>(1) + as_type<uint>(2))));
-  int3 v3 = int3((1 + 2));
+  int3 v3 = int3(as_type<int>((as_type<uint>(1) + as_type<uint>(2))));
-  int4 v4 = int4((1 + 2));
+  int4 v4 = int4(as_type<int>((as_type<uint>(1) + as_type<uint>(2))));
 }
 

test/expressions/splat/var/i32.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 void f() {
-  int v = (1 + 2);
+  int v = as_type<int>((as_type<uint>(1) + as_type<uint>(2)));
   int2 v2 = int2(v);
   int3 v3 = int3(v);
   int4 v4 = int4(v);

test/expressions/unary/negate/negate.wgsl

@@ -0,0 +1,7 @@
+fn i(x : i32) -> i32 {
+  return -x;
+}
+
+fn vi(x : vec4<i32>) -> vec4<i32> {
+  return -x;
+}

test/expressions/unary/negate/negate.wgsl.expected.hlsl

@@ -0,0 +1,12 @@
+[numthreads(1, 1, 1)]
+void unused_entry_point() {
+  return;
+}
+
+int i(int x) {
+  return -(x);
+}
+
+int4 vi(int4 x) {
+  return -(x);
+}

test/expressions/unary/negate/negate.wgsl.expected.msl

@@ -0,0 +1,20 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+int tint_unary_minus(const int v) {
+  return select(-v, v, v == -2147483648);
+}
+
+int4 tint_unary_minus_1(const int4 v) {
+  return select(-v, v, v == -2147483648);
+}
+
+int i(int x) {
+  return tint_unary_minus(x);
+}
+
+int4 vi(int4 x) {
+  return tint_unary_minus_1(x);
+}
+

test/expressions/unary/negate/negate.wgsl.expected.spvasm

@@ -0,0 +1,36 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %unused_entry_point "unused_entry_point"
+               OpExecutionMode %unused_entry_point LocalSize 1 1 1
+               OpName %unused_entry_point "unused_entry_point"
+               OpName %i "i"
+               OpName %x "x"
+               OpName %vi "vi"
+               OpName %x_0 "x"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+          %5 = OpTypeFunction %int %int
+      %v4int = OpTypeVector %int 4
+         %11 = OpTypeFunction %v4int %v4int
+%unused_entry_point = OpFunction %void None %1
+          %4 = OpLabel
+               OpReturn
+               OpFunctionEnd
+          %i = OpFunction %int None %5
+          %x = OpFunctionParameter %int
+          %9 = OpLabel
+         %10 = OpSNegate %int %x
+               OpReturnValue %10
+               OpFunctionEnd
+         %vi = OpFunction %v4int None %11
+        %x_0 = OpFunctionParameter %v4int
+         %15 = OpLabel
+         %16 = OpSNegate %v4int %x_0
+               OpReturnValue %16
+               OpFunctionEnd

test/expressions/unary/negate/negate.wgsl.expected.wgsl

@@ -0,0 +1,7 @@
+fn i(x : i32) -> i32 {
+  return -(x);
+}
+
+fn vi(x : vec4<i32>) -> vec4<i32> {
+  return -(x);
+}

test/intrinsics/ignore/call.wgsl.expected.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() {

test/loops/continue_in_switch.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 kernel void f() {
-  for(int i = 0; (i < 4); i = (i + 1)) {
+  for(int i = 0; (i < 4); i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
     switch(i) {
       case 0: {
         continue;

test/loops/loop.wgsl.expected.msl

@@ -4,7 +4,7 @@ using namespace metal;
 int f() {
   int i = 0;
   while (true) {
-    i = (i + 1);
+    i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
     if ((i > 4)) {
       return i;
     }

test/loops/loop_with_continuing.wgsl.expected.msl

@@ -8,7 +8,7 @@ int f() {
       return i;
     }
     {
-      i = (i + 1);
+      i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
     }
   }
   return 0;

test/loops/nested_loops.wgsl.expected.msl

@@ -5,12 +5,12 @@ int f() {
   int i = 0;
   int j = 0;
   while (true) {
-    i = (i + 1);
+    i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
     if ((i > 4)) {
       return 1;
     }
     while (true) {
-      j = (j + 1);
+      j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)));
       if ((j > 4)) {
         return 2;
       }

test/loops/nested_loops_with_continuing.wgsl.expected.msl

@@ -13,11 +13,11 @@ int f() {
         return 2;
       }
       {
-        j = (j + 1);
+        j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)));
       }
     }
     {
-      i = (i + 1);
+      i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
     }
   }
   return 0;

test/ptr_ref/load/global/i32.spvasm.expected.msl

@@ -3,7 +3,7 @@
 using namespace metal;
 void main_1(thread int* const tint_symbol_1) {
   int const x_9 = *(tint_symbol_1);
-  int const x_11 = (x_9 + 1);
+  int const x_11 = as_type<int>((as_type<uint>(x_9) + as_type<uint>(1)));
   return;
 }
 

test/ptr_ref/load/global/i32.wgsl.expected.msl

@@ -4,7 +4,7 @@ using namespace metal;
 kernel void tint_symbol() {
   thread int tint_symbol_1 = 0;
   int const i = tint_symbol_1;
-  int const use = (i + 1);
+  int const use = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
   return;
 }
 

test/ptr_ref/load/local/i32.spvasm.expected.msl

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

test/ptr_ref/load/local/i32.wgsl.expected.msl

@@ -3,7 +3,7 @@
 using namespace metal;
 kernel void tint_symbol() {
   int i = 123;
-  int const use = (i + 1);
+  int const use = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
   return;
 }
 

test/ptr_ref/load/local/ptr_function.wgsl.expected.msl

@@ -3,7 +3,7 @@
 using namespace metal;
 kernel void tint_symbol() {
   int i = 123;
-  int const use = (i + 1);
+  int const use = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
   return;
 }
 

test/ptr_ref/load/local/ptr_private.wgsl.expected.msl

@@ -3,7 +3,7 @@
 using namespace metal;
 kernel void tint_symbol() {
   thread int tint_symbol_1 = 123;
-  int const use = (tint_symbol_1 + 1);
+  int const use = as_type<int>((as_type<uint>(tint_symbol_1) + as_type<uint>(1)));
   return;
 }
 

test/ptr_ref/load/local/ptr_storage.wgsl.expected.msl

@@ -6,7 +6,7 @@ struct S {
 };
 
 kernel void tint_symbol(device S& v [[buffer(0)]]) {
-  int const use = (v.a + 1);
+  int const use = as_type<int>((as_type<uint>(v.a) + as_type<uint>(1)));
   return;
 }