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

This reverts commit e5dbe24e94.

Reason for revert: Makes the Tint-Dawn roll fails because of
MSL compilation errors on as_type<uint>(-2147483648):
as_type cast from 'long' to 'uint' (aka 'unsigned int') is not allowed
   
as_type<uint>(-2147483647) compiles fine, so this is most
likely because the MSL compiler types the literal as a long
(since without the - it is larger than the max int32).

Original change's description:
> MSL writer: make signed int overflow defined behaviour
>
> Bug: tint:124
> Change-Id: Icf545b633d6390ceb7f639e80111390005e311a1
> Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60100
> Kokoro: Kokoro <noreply+kokoro@google.com>
> Commit-Queue: Antonio Maiorano <amaiorano@google.com>
> Reviewed-by: David Neto <dneto@google.com>

TBR=dneto@google.com,bclayton@google.com,jrprice@google.com,amaiorano@google.com,noreply+kokoro@google.com,tint-scoped@luci-project-accounts.iam.gserviceaccount.com

Change-Id: I3e3384a9185013bb141a1b7b9b22bad8571bbc50
No-Presubmit: true
No-Tree-Checks: true
No-Try: true
Bug: tint:124
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/60345
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Reviewed-by: Ben Clayton <bclayton@google.com>
Auto-Submit: Corentin Wallez <cwallez@chromium.org>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>

src/writer/msl/generator_impl.cc

@@ -17,7 +17,6 @@
 #include <algorithm>
 #include <cmath>
 #include <iomanip>
-#include <limits>
 #include <utility>
 #include <vector>
 
@@ -75,32 +74,6 @@ 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) {}
@@ -253,7 +226,25 @@ bool GeneratorImpl::EmitAssign(ast::AssignmentStatement* stmt) {
 }
 
 bool GeneratorImpl::EmitBinary(std::ostream& out, ast::BinaryExpression* expr) {
-  auto emit_op = [&] {
+  if (expr->op() == ast::BinaryOp::kModulo &&
+      TypeOf(expr)->UnwrapRef()->is_float_scalar_or_vector()) {
+    out << "fmod";
+    ScopedParen sp(out);
+    if (!EmitExpression(out, expr->lhs())) {
+      return false;
+    }
+    out << ", ";
+    if (!EmitExpression(out, expr->rhs())) {
+      return false;
+    }
+    return true;
+  }
+
+  ScopedParen sp(out);
+
+  if (!EmitExpression(out, expr->lhs())) {
+    return false;
+  }
   out << " ";
 
   switch (expr->op()) {
@@ -296,8 +287,8 @@ bool GeneratorImpl::EmitBinary(std::ostream& out, ast::BinaryExpression* expr) {
     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
+      // generate extra code to implement WGSL-specified behaviour for negative
-        // negative LHS.
+      // LHS.
       out << R"(>>)";
       break;
 
@@ -322,116 +313,7 @@ bool GeneratorImpl::EmitBinary(std::ostream& out, ast::BinaryExpression* expr) {
       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 &&
-      lhs_type->is_float_scalar_or_vector()) {
-    out << "fmod";
-    ScopedParen sp(out);
-    if (!EmitExpression(out, expr->lhs())) {
-      return false;
-    }
-    out << ", ";
-    if (!EmitExpression(out, expr->rhs())) {
-      return false;
-    }
-    return true;
-  }
-
-  // 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;
-  }
-  if (!emit_op()) {
-    return false;
-  }
   if (!EmitExpression(out, expr->rhs())) {
     return false;
   }
@@ -2445,53 +2327,6 @@ 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,7 +356,6 @@ 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,85 +74,6 @@ 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 as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(a) + as_type<uint>(b)))) * as_type<uint>(c)));
+  return ((a + b) * c);
 }
 
 kernel void func() {

src/writer/msl/generator_impl_loop_test.cc

@@ -259,7 +259,7 @@ TEST_F(MslGeneratorImplTest, Emit_ForLoopWithSimpleCont) {
   gen.increment_indent();
 
   ASSERT_TRUE(gen.EmitStatement(f)) << gen.error();
-  EXPECT_EQ(gen.result(), R"(  for(; ; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  EXPECT_EQ(gen.result(), R"(  for(; ; i = (i + 1)) {
     a_statement();
   }
 )");
@@ -310,7 +310,7 @@ 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 = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  EXPECT_EQ(gen.result(), R"(  for(int i = 0; true; i = (i + 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(), "tint_unary_minus(expr)");
+  EXPECT_EQ(out.str(), "-(expr)");
 }
 
 }  // 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, as_type<int>((as_type<uint>(x) + as_type<uint>(1))), nonempty.arr[3]}};
+  tint_array_wrapper const nonempty_with_expr = {.arr={1, x, (x + 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, as_type<int>((as_type<uint>(x) + as_type<uint>(1)))}};
+  tint_array_wrapper const tint_symbol_9 = {.arr={1, 2, x, (x + 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 const tint_symbol_10 = {.arr={5, 6, nonempty.arr[2], (nonempty.arr[3] + 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, as_type<int>((as_type<uint>(x) + as_type<uint>(1))), nonempty.arr[3]}};
+  tint_array_wrapper const tint_symbol_14 = {.arr={1, x, (x + 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, as_type<int>((as_type<uint>(x) + as_type<uint>(1))), nonempty.arr[3]}};
+  tint_array_wrapper const tint_symbol_19 = {.arr={1, x, (x + 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 = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 2); i = (i + 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 = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 2); i = (i + 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 = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 2); i = (i + 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 = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 2); i = (i + 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 = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 2); i = (i + 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 = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  for(int i = 0; (i < 2); i = (i + 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 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(size.y) - as_type<uint>(dstTexCoord.y)))) - as_type<uint>(1)));
+    srcTexCoord.y = ((size.y - dstTexCoord.y) - 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 = as_type<int>((as_type<uint>(x_45) - as_type<uint>(as_type<int>(1u))));
+  i_1 = (x_45 - 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 <= as_type<int>((as_type<uint>(x_56) - as_type<uint>(as_type<int>(1u)))))) {
+    if ((x_55 <= (x_56 - 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 = as_type<int>((as_type<uint>(x_67) + as_type<uint>(as_type<int>(1u))));
+      i_1 = (x_67 + 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 = as_type<int>((as_type<uint>(1) + as_type<uint>(x_74)));
+      j_1 = (1 + 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 = as_type<int>((as_type<uint>(1) + as_type<uint>(x_76)));
+  i_1 = (1 + 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 = as_type<int>((as_type<uint>(x_93) + as_type<uint>(as_type<int>(1u))));
+  int const x_94 = (x_93 + 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 = as_type<int>((as_type<uint>(x_97) + as_type<uint>(1)));
+  int const x_98 = (x_97 + 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 = as_type<int>((as_type<uint>(x_108) - as_type<uint>(as_type<int>(1u))));
+    top = (x_108 - 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 = as_type<int>((as_type<uint>(x_112) - as_type<uint>(1)));
+    top = (x_112 - 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 ((as_type<int>((as_type<uint>(x_122) - as_type<uint>(as_type<int>(1u)))) > x_124)) {
+    if (((x_122 - 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 = as_type<int>((as_type<uint>(1) + as_type<uint>(x_128)));
+      int const x_131_save = (1 + 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] = as_type<int>((as_type<uint>(x_134) - as_type<uint>(as_type<int>(1u))));
+      stack.arr[x_136_save] = (x_134 - 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 = as_type<int>((as_type<uint>(x_143) + as_type<uint>(1)));
+      int const x_144 = (x_143 + 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 = as_type<int>((as_type<uint>(x_148) + as_type<uint>(as_type<int>(1u))));
+      int const x_149 = (x_148 + 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

@@ -24,10 +24,10 @@ kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_4 [[texture
   threadgroup_barrier(mem_flags::mem_threadgroup);
   uint const filterOffset = ((params.filterDim - 1u) / 2u);
   int2 const dims = int2(tint_symbol_4.get_width(0), tint_symbol_4.get_height(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))));
+  int2 const baseIndex = (int2(((WorkGroupID.xy * uint2(params.blockDim, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u)))) - int2(int(filterOffset), 0));
   for(uint r = 0u; (r < 4u); r = (r + 1u)) {
     for(uint c = 0u; (c < 4u); c = (c + 1u)) {
-      int2 loadIndex = as_type<int2>((as_type<uint2>(baseIndex) + as_type<uint2>(int2(int(c), int(r)))));
+      int2 loadIndex = (baseIndex + int2(int(c), int(r)));
       if ((flip.value != 0u)) {
         loadIndex = loadIndex.yx;
       }
@@ -37,7 +37,7 @@ kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_4 [[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 = as_type<int2>((as_type<uint2>(baseIndex) + as_type<uint2>(int2(int(c), int(r)))));
+      int2 writeIndex = (baseIndex + 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 = as_type<int>((as_type<uint>(x_417) * as_type<uint>(x_419)));
+  batchASize = (x_417 * x_419);
   int const x_421 = *(row);
   int const x_422 = *(col);
   int const x_424 = *(tint_symbol_5);
@@ -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_6);
     int const x_445 = *(col);
-    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)))];
+    float const x_448 = x_165.A[(((x_438 * x_439) + (x_441 * x_442)) + 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 = as_type<int>((as_type<uint>(x_455) * as_type<uint>(x_457)));
+  batchBSize = (x_455 * x_457);
   int const x_459 = *(row_1);
   int const x_460 = *(col_1);
   int const x_462 = *(tint_symbol_8);
@@ -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_9);
     int const x_482 = *(col_1);
-    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)))];
+    float const x_485 = x_185.B[(((x_475 * x_476) + (x_478 * x_479)) + 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_12)).y;
-  tileRow = as_type<int>((as_type<uint>(as_type<int>(x_132)) * as_type<uint>(1)));
+  tileRow = (as_type<int>(x_132) * 1);
   uint const x_137 = (*(tint_symbol_12)).x;
-  tileCol = as_type<int>((as_type<uint>(as_type<int>(x_137)) * as_type<uint>(1)));
+  tileCol = (as_type<int>(x_137) * 1);
   uint const x_143 = (*(tint_symbol_13)).y;
-  globalRow = as_type<int>((as_type<uint>(as_type<int>(x_143)) * as_type<uint>(1)));
+  globalRow = (as_type<int>(x_143) * 1);
   uint const x_148 = (*(tint_symbol_13)).x;
-  globalCol = as_type<int>((as_type<uint>(as_type<int>(x_148)) * as_type<uint>(1)));
+  globalCol = (as_type<int>(x_148) * 1);
   int const x_152 = *(dimInner);
-  numTiles = as_type<int>((as_type<uint>((as_type<int>((as_type<uint>(x_152) - as_type<uint>(1))) / 64)) + as_type<uint>(1)));
+  numTiles = (((x_152 - 1) / 64) + 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 = as_type<int>((as_type<uint>(x_181) + as_type<uint>(1)));
+        innerCol = (x_181 + 1);
       }
     }
     {
       int const x_183 = innerRow;
-      innerRow = as_type<int>((as_type<uint>(x_183) + as_type<uint>(1)));
+      innerRow = (x_183 + 1);
     }
   }
   uint const x_187 = (*(tint_symbol_12)).x;
-  tileColA = as_type<int>((as_type<uint>(as_type<int>(x_187)) * as_type<uint>(64)));
+  tileColA = (as_type<int>(x_187) * 64);
   uint const x_192 = (*(tint_symbol_12)).y;
-  tileRowB = as_type<int>((as_type<uint>(as_type<int>(x_192)) * as_type<uint>(1)));
+  tileRowB = (as_type<int>(x_192) * 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 = as_type<int>((as_type<uint>(x_221) + as_type<uint>(x_222)));
+        inputRow = (x_221 + x_222);
         int const x_225 = tileColA;
         int const x_226 = innerCol_1;
-        inputCol = as_type<int>((as_type<uint>(x_225) + as_type<uint>(x_226)));
+        inputCol = (x_225 + 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 = as_type<int>((as_type<uint>(x_235) + as_type<uint>(x_236)));
+        param_3 = (x_235 + x_236);
-        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)));
+        param_4 = ((x_238 * 64) + x_240);
         float const x_244 = mm_readA_i1_i1_(x_48, x_165, &(param_3), &(param_4), tint_symbol_14, tint_symbol_15, tint_symbol_16);
         (*(tint_symbol_17)).arr[x_233].arr[x_234] = x_244;
         {
           int const x_247 = innerCol_1;
-          innerCol_1 = as_type<int>((as_type<uint>(x_247) + as_type<uint>(1)));
+          innerCol_1 = (x_247 + 1);
         }
       }
       {
         int const x_249 = innerRow_1;
-        innerRow_1 = as_type<int>((as_type<uint>(x_249) + as_type<uint>(1)));
+        innerRow_1 = (x_249 + 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 = as_type<int>((as_type<uint>(x_268) + as_type<uint>(x_269)));
+        inputRow_1 = (x_268 + x_269);
         int const x_272 = tileCol;
         int const x_273 = innerCol_2;
-        inputCol_1 = as_type<int>((as_type<uint>(x_272) + as_type<uint>(x_273)));
+        inputCol_1 = (x_272 + 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 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(x_280) * as_type<uint>(64)))) + as_type<uint>(x_282)));
+        param_5 = ((x_280 * 64) + x_282);
-        param_6 = as_type<int>((as_type<uint>(x_284) + as_type<uint>(x_285)));
+        param_6 = (x_284 + x_285);
         float const x_289 = mm_readB_i1_i1_(x_48, x_185, &(param_5), &(param_6), tint_symbol_15, tint_symbol_18, tint_symbol_16);
         (*(tint_symbol_19)).arr[x_278].arr[x_279] = x_289;
         {
           int const x_291 = innerCol_2;
-          innerCol_2 = as_type<int>((as_type<uint>(x_291) + as_type<uint>(1)));
+          innerCol_2 = (x_291 + 1);
         }
       }
       {
         int const x_293 = innerRow_2;
-        innerRow_2 = as_type<int>((as_type<uint>(x_293) + as_type<uint>(1)));
+        innerRow_2 = (x_293 + 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_19)).arr[x_315].arr[as_type<int>((as_type<uint>(x_316) + as_type<uint>(x_317)))];
+        float const x_320 = (*(tint_symbol_19)).arr[x_315].arr[(x_316 + x_317)];
         BCached.arr[x_314] = x_320;
         {
           int const x_322 = inner;
-          inner = as_type<int>((as_type<uint>(x_322) + as_type<uint>(1)));
+          inner = (x_322 + 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_17)).arr[as_type<int>((as_type<uint>(x_333) + as_type<uint>(x_334)))].arr[x_336];
+        float const x_338 = (*(tint_symbol_17)).arr[(x_333 + 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 = as_type<int>((as_type<uint>(x_358) + as_type<uint>(1)));
+            innerCol_3 = (x_358 + 1);
           }
         }
         {
           int const x_360 = innerRow_3;
-          innerRow_3 = as_type<int>((as_type<uint>(x_360) + as_type<uint>(1)));
+          innerRow_3 = (x_360 + 1);
         }
       }
       {
         int const x_362 = k;
-        k = as_type<int>((as_type<uint>(x_362) + as_type<uint>(1)));
+        k = (x_362 + 1);
       }
     }
     threadgroup_barrier(mem_flags::mem_threadgroup);
     {
       int const x_364 = t;
-      t = as_type<int>((as_type<uint>(x_364) + as_type<uint>(1)));
+      t = (x_364 + 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 = (as_type<int>((as_type<uint>(x_382) + as_type<uint>(x_383))) < x_385);
+      bool const x_386 = ((x_382 + 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 = (as_type<int>((as_type<uint>(x_389) + as_type<uint>(x_390))) < x_392);
+        x_393 = ((x_389 + 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 = as_type<int>((as_type<uint>(x_397) + as_type<uint>(x_398)));
+        param_7 = (x_397 + x_398);
-        param_8 = as_type<int>((as_type<uint>(x_400) + as_type<uint>(x_401)));
+        param_8 = (x_400 + 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_16);
       }
       {
         int const x_411 = innerCol_4;
-        innerCol_4 = as_type<int>((as_type<uint>(x_411) + as_type<uint>(1)));
+        innerCol_4 = (x_411 + 1);
       }
     }
     {
       int const x_413 = innerRow_4;
-      innerRow_4 = as_type<int>((as_type<uint>(x_413) + as_type<uint>(1)));
+      innerRow_4 = (x_413 + 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 = as_type<int>((as_type<uint>(x_304) + as_type<uint>(1)));
+      i = (x_304 + 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 = as_type<int>((as_type<uint>(x_441) + as_type<uint>(1)));
+      i = (x_441 + 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 = as_type<int>((as_type<uint>(a) + as_type<uint>(b)));
+  int const r = (a + 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 = as_type<int3>((as_type<uint>(a) + as_type<uint3>(b)));
+  int3 const r = (a + 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 = as_type<int3>((as_type<uint3>(a) + as_type<uint>(b)));
+  int3 const r = (a + 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 = as_type<int3>((as_type<uint3>(a) + as_type<uint3>(b)));
+  int3 const r = (a + b);
   return;
 }
 

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

@@ -1,6 +0,0 @@
-[[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

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

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

@@ -1,10 +0,0 @@
-#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

@@ -1,21 +0,0 @@
-; 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

@@ -1,6 +0,0 @@
-[[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

@@ -1,6 +0,0 @@
-[[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

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

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

@@ -1,10 +0,0 @@
-#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

@@ -1,20 +0,0 @@
-; 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

@@ -1,6 +0,0 @@
-[[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

@@ -1,6 +0,0 @@
-[[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

@@ -1,7 +0,0 @@
-[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

@@ -1,10 +0,0 @@
-#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

@@ -1,29 +0,0 @@
-; 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

@@ -1,6 +0,0 @@
-[[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

@@ -1,6 +0,0 @@
-[[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

@@ -1,7 +0,0 @@
-[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

@@ -1,10 +0,0 @@
-#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

@@ -1,27 +0,0 @@
-; 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

@@ -1,6 +0,0 @@
-[[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 = as_type<int>((as_type<uint>(a) * as_type<uint>(b)));
+  int const r = (a * 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 = as_type<int3>((as_type<uint>(a) * as_type<uint3>(b)));
+  int3 const r = (a * 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 = as_type<int3>((as_type<uint3>(a) * as_type<uint>(b)));
+  int3 const r = (a * 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 = as_type<int3>((as_type<uint3>(a) * as_type<uint3>(b)));
+  int3 const r = (a * b);
   return;
 }
 

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

@@ -1,6 +0,0 @@
-[[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

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

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

@@ -1,10 +0,0 @@
-#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

@@ -1,21 +0,0 @@
-; 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

@@ -1,6 +0,0 @@
-[[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

@@ -1,6 +0,0 @@
-[[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

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

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

@@ -1,10 +0,0 @@
-#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

@@ -1,20 +0,0 @@
-; 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

@@ -1,6 +0,0 @@
-[[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

@@ -1,6 +0,0 @@
-[[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

@@ -1,7 +0,0 @@
-[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

@@ -1,10 +0,0 @@
-#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

@@ -1,29 +0,0 @@
-; 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

@@ -1,6 +0,0 @@
-[[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

@@ -1,6 +0,0 @@
-[[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

@@ -1,7 +0,0 @@
-[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

@@ -1,10 +0,0 @@
-#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

@@ -1,27 +0,0 @@
-; 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

@@ -1,6 +0,0 @@
-[[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 = as_type<int>((as_type<uint>(a) - as_type<uint>(b)));
+  int const r = (a - 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 = as_type<int3>((as_type<uint>(a) - as_type<uint3>(b)));
+  int3 const r = (a - 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 = as_type<int3>((as_type<uint3>(a) - as_type<uint>(b)));
+  int3 const r = (a - 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 = as_type<int3>((as_type<uint3>(a) - as_type<uint3>(b)));
+  int3 const r = (a - b);
   return;
 }
 

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

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

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

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

test/expressions/unary/negate/negate.wgsl

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

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

@@ -1,12 +0,0 @@
-[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

@@ -1,20 +0,0 @@
-#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

@@ -1,36 +0,0 @@
-; 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

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

test/intrinsics/ignore.wgsl.expected.msl

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

test/loops/loop.wgsl.expected.msl

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

test/loops/loop_with_continuing.wgsl.expected.msl

@@ -8,7 +8,7 @@ int f() {
       return i;
     }
     {
-      i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
+      i = (i + 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 = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
+    i = (i + 1);
     if ((i > 4)) {
       return 1;
     }
     while (true) {
-      j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)));
+      j = (j + 1);
       if ((j > 4)) {
         return 2;
       }

test/loops/nested_loops_with_continuing.wgsl.expected.msl

@@ -13,11 +13,11 @@ int f() {
         return 2;
       }
       {
-        j = as_type<int>((as_type<uint>(j) + as_type<uint>(1)));
+        j = (j + 1);
       }
     }
     {
-      i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
+      i = (i + 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 = as_type<int>((as_type<uint>(x_9) + as_type<uint>(1)));
+  int const x_11 = (x_9 + 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 = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
+  int const use = (i + 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 = as_type<int>((as_type<uint>(x_10) + as_type<uint>(1)));
+  int const x_12 = (x_10 + 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 = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
+  int const use = (i + 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 = as_type<int>((as_type<uint>(i) + as_type<uint>(1)));
+  int const use = (i + 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 = as_type<int>((as_type<uint>(tint_symbol_1) + as_type<uint>(1)));
+  int const use = (tint_symbol_1 + 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 = as_type<int>((as_type<uint>(v.a) + as_type<uint>(1)));
+  int const use = (v.a + 1);
   return;
 }
 

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

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

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

@@ -8,7 +8,7 @@ kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgrou
   }
   threadgroup_barrier(mem_flags::mem_threadgroup);
   tint_symbol_2 = 123;
-  int const use = as_type<int>((as_type<uint>(tint_symbol_2) + as_type<uint>(1)));
+  int const use = (tint_symbol_2 + 1);
   return;
 }
 

test/ptr_ref/load/param/ptr.spvasm.expected.msl

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

test/ptr_ref/load/param/ptr.wgsl.expected.msl

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

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

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

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

@@ -4,7 +4,7 @@ using namespace metal;
 kernel void tint_symbol() {
   thread int tint_symbol_1 = 0;
   tint_symbol_1 = 123;
-  tint_symbol_1 = as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(100) + as_type<uint>(20)))) + as_type<uint>(3)));
+  tint_symbol_1 = ((100 + 20) + 3);
   return;
 }