diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index f774315..95b32c0 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -152,6 +152,32 @@
     std::ostream& s;
 };
 
+class ScopedCast {
+  public:
+    ScopedCast(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, cast the scalar to the
+        // vector element type.
+        if (curr_type->is_scalar() && target_vec_type) {
+            target_type = target_vec_type->type();
+        }
+
+        // Cast
+        generator->EmitType(s, target_type, "");
+        s << "(";
+    }
+
+    ~ScopedCast() { s << ")"; }
+
+  private:
+    std::ostream& s;
+};
+
 }  // namespace
 
 SanitizedResult::SanitizedResult() = default;
@@ -521,7 +547,7 @@
             // In case the type is packed, cast to our own type in order to remove the packing.
             // Otherwise, this just casts to itself.
             if (lhs_type->is_signed_integer_vector()) {
-                ScopedBitCast lhs_self_cast(this, out, lhs_type, lhs_type);
+                ScopedCast lhs_self_cast(this, out, lhs_type, lhs_type);
                 if (!EmitExpression(out, expr->lhs)) {
                     return false;
                 }
@@ -540,7 +566,7 @@
             // In case the type is packed, cast to our own type in order to remove the packing.
             // Otherwise, this just casts to itself.
             if (rhs_type->is_signed_integer_vector()) {
-                ScopedBitCast rhs_self_cast(this, out, rhs_type, rhs_type);
+                ScopedCast rhs_self_cast(this, out, rhs_type, rhs_type);
                 if (!EmitExpression(out, expr->rhs)) {
                     return false;
                 }
diff --git a/test/tint/bug/tint/1121.wgsl.expected.msl b/test/tint/bug/tint/1121.wgsl.expected.msl
index 33fb2df..34c70b2 100644
--- a/test/tint/bug/tint/1121.wgsl.expected.msl
+++ b/test/tint/bug/tint/1121.wgsl.expected.msl
@@ -88,7 +88,7 @@
     for(int x_1 = 0; (x_1 < TILE_COUNT_X); x_1 = as_type<int>((as_type<uint>(x_1) + as_type<uint>(1)))) {
       int2 tilePixel0Idx = int2(as_type<int>((as_type<uint>(x_1) * as_type<uint>(TILE_SIZE))), as_type<int>((as_type<uint>(y_1) * as_type<uint>(TILE_SIZE))));
       float2 floorCoord = (((2.0f * float2(tilePixel0Idx)) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f));
-      float2 ceilCoord = (((2.0f * float2(as_type<int2>((as_type<uint2>(as_type<int2>(tilePixel0Idx)) + as_type<uint2>(as_type<int2>(int2(TILE_SIZE))))))) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f));
+      float2 ceilCoord = (((2.0f * float2(as_type<int2>((as_type<uint2>(int2(tilePixel0Idx)) + as_type<uint2>(int2(int2(TILE_SIZE))))))) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f));
       float2 viewFloorCoord = float2((((-(viewNear) * floorCoord[0]) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * floorCoord[1]) - (M[2][1] * viewNear)) / M[1][1]));
       float2 viewCeilCoord = float2((((-(viewNear) * ceilCoord[0]) - (M[2][0] * viewNear)) / M[0][0]), (((-(viewNear) * ceilCoord[1]) - (M[2][1] * viewNear)) / M[1][1]));
       frustumPlanes[0] = float4(1.0f, 0.0f, (-(viewFloorCoord[0]) / viewNear), 0.0f);
diff --git a/test/tint/bug/tint/1520.spvasm.expected.msl b/test/tint/bug/tint/1520.spvasm.expected.msl
index a409d7d..3b8e4f7 100644
--- a/test/tint/bug/tint/1520.spvasm.expected.msl
+++ b/test/tint/bug/tint/1520.spvasm.expected.msl
@@ -45,13 +45,13 @@
   ok = x_41;
   int4 const x_44 = int4(x_27, x_27, x_27, x_27);
   val = x_44;
-  int4 const x_47 = as_type<int4>((as_type<uint4>(as_type<int4>(x_44)) + as_type<uint4>(as_type<int4>(int4(1)))));
+  int4 const x_47 = as_type<int4>((as_type<uint4>(int4(x_44)) + as_type<uint4>(int4(int4(1)))));
   val = x_47;
-  int4 const x_48 = as_type<int4>((as_type<uint4>(as_type<int4>(x_47)) - as_type<uint4>(as_type<int4>(int4(1)))));
+  int4 const x_48 = as_type<int4>((as_type<uint4>(int4(x_47)) - as_type<uint4>(int4(int4(1)))));
   val = x_48;
-  int4 const x_49 = as_type<int4>((as_type<uint4>(as_type<int4>(x_48)) + as_type<uint4>(as_type<int4>(int4(1)))));
+  int4 const x_49 = as_type<int4>((as_type<uint4>(int4(x_48)) + as_type<uint4>(int4(int4(1)))));
   val = x_49;
-  int4 const x_50 = as_type<int4>((as_type<uint4>(as_type<int4>(x_49)) - as_type<uint4>(as_type<int4>(int4(1)))));
+  int4 const x_50 = as_type<int4>((as_type<uint4>(int4(x_49)) - as_type<uint4>(int4(int4(1)))));
   val = x_50;
   x_55 = false;
   if (x_41) {
@@ -59,11 +59,11 @@
     x_55 = x_54;
   }
   ok = x_55;
-  int4 const x_58 = as_type<int4>((as_type<uint4>(as_type<int4>(x_50)) * as_type<uint4>(as_type<int4>(int4(2)))));
+  int4 const x_58 = as_type<int4>((as_type<uint4>(int4(x_50)) * as_type<uint4>(int4(int4(2)))));
   val = x_58;
   int4 const x_59 = (x_58 / int4(2));
   val = x_59;
-  int4 const x_60 = as_type<int4>((as_type<uint4>(as_type<int4>(x_59)) * as_type<uint4>(as_type<int4>(int4(2)))));
+  int4 const x_60 = as_type<int4>((as_type<uint4>(int4(x_59)) * as_type<uint4>(int4(int4(2)))));
   val = x_60;
   int4 const x_61 = (x_60 / int4(2));
   val = x_61;
diff --git a/test/tint/bug/tint/1677.wgsl.expected.msl b/test/tint/bug/tint/1677.wgsl.expected.msl
index 8c3c019..6ee8bbd 100644
--- a/test/tint/bug/tint/1677.wgsl.expected.msl
+++ b/test/tint/bug/tint/1677.wgsl.expected.msl
@@ -20,7 +20,7 @@
 };
 
 void tint_symbol_inner(uint3 id, const device Input* const tint_symbol_1) {
-  int3 const pos = as_type<int3>((as_type<uint3>(as_type<int3>((*(tint_symbol_1)).position)) - as_type<uint3>(as_type<int3>(int3(0)))));
+  int3 const pos = as_type<int3>((as_type<uint3>(int3((*(tint_symbol_1)).position)) - as_type<uint3>(int3(int3(0)))));
 }
 
 kernel void tint_symbol(const device Input* tint_symbol_2 [[buffer(0)]], uint3 id [[thread_position_in_grid]]) {
diff --git a/test/tint/bug/tint/942.wgsl.expected.msl b/test/tint/bug/tint/942.wgsl.expected.msl
index 868cd1e..d605a4f 100644
--- a/test/tint/bug/tint/942.wgsl.expected.msl
+++ b/test/tint/bug/tint/942.wgsl.expected.msl
@@ -32,10 +32,10 @@
   threadgroup_barrier(mem_flags::mem_threadgroup);
   uint const filterOffset = (((*(tint_symbol_2)).filterDim - 1u) / 2u);
   int2 const dims = int2(tint_symbol_3.get_width(0), tint_symbol_3.get_height(0));
-  int2 const baseIndex = as_type<int2>((as_type<uint2>(as_type<int2>(int2(((uint3(WorkGroupID).xy * uint2((*(tint_symbol_2)).blockDim, 4u)) + (uint3(LocalInvocationID).xy * uint2(4u, 1u)))))) - as_type<uint2>(as_type<int2>(int2(int(filterOffset), 0)))));
+  int2 const baseIndex = as_type<int2>((as_type<uint2>(int2(int2(((uint3(WorkGroupID).xy * uint2((*(tint_symbol_2)).blockDim, 4u)) + (uint3(LocalInvocationID).xy * uint2(4u, 1u)))))) - as_type<uint2>(int2(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>(as_type<int2>(baseIndex)) + as_type<uint2>(as_type<int2>(int2(int(c), int(r))))));
+      int2 loadIndex = as_type<int2>((as_type<uint2>(int2(baseIndex)) + as_type<uint2>(int2(int2(int(c), int(r))))));
       if (((*(tint_symbol_4)).value != 0u)) {
         loadIndex = int2(loadIndex).yx;
       }
@@ -45,7 +45,7 @@
   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>(as_type<int2>(baseIndex)) + as_type<uint2>(as_type<int2>(int2(int(c), int(r))))));
+      int2 writeIndex = as_type<int2>((as_type<uint2>(int2(baseIndex)) + as_type<uint2>(int2(int2(int(c), int(r))))));
       if (((*(tint_symbol_4)).value != 0u)) {
         writeIndex = int2(writeIndex).yx;
       }
diff --git a/test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.msl
index 8ba40b7..8658c03 100644
--- a/test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/add/scalar-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
 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>(as_type<int3>(b))));
+  int3 const r = as_type<int3>((as_type<uint>(a) + as_type<uint3>(int3(b))));
   return;
 }
 
diff --git a/test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.msl b/test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.msl
index d5f6b0e..de4195f 100644
--- a/test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/add/vec3-scalar/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
 kernel void f() {
   int3 const a = int3(1, 2, 3);
   int const b = 4;
-  int3 const r = as_type<int3>((as_type<uint3>(as_type<int3>(a)) + as_type<uint>(b)));
+  int3 const r = as_type<int3>((as_type<uint3>(int3(a)) + as_type<uint>(b)));
   return;
 }
 
diff --git a/test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.msl
index a60e42a..009e87e 100644
--- a/test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/add/vec3-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
 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>(as_type<int3>(a)) + as_type<uint3>(as_type<int3>(b))));
+  int3 const r = as_type<int3>((as_type<uint3>(int3(a)) + as_type<uint3>(int3(b))));
   return;
 }
 
diff --git a/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
index 6f38d4b..6d2f8bc 100644
--- a/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
 kernel void f() {
   int a = 4;
   int3 b = int3(0, 2, 0);
-  int3 const r = (a / as_type<int3>((as_type<uint3>(as_type<int3>(b)) + as_type<uint3>(as_type<int3>(b)))));
+  int3 const r = (a / as_type<int3>((as_type<uint3>(int3(b)) + as_type<uint3>(int3(b)))));
   return;
 }
 
diff --git a/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
index 68a2b33..c90c00e 100644
--- a/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/div_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
 kernel void f() {
   int3 a = int3(1, 2, 3);
   int3 b = int3(0, 5, 0);
-  int3 const r = (a / as_type<int3>((as_type<uint3>(as_type<int3>(b)) + as_type<uint3>(as_type<int3>(b)))));
+  int3 const r = (a / as_type<int3>((as_type<uint3>(int3(b)) + as_type<uint3>(int3(b)))));
   return;
 }
 
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
index 2545d86..22f6994 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/scalar-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
 kernel void f() {
   int a = 4;
   int3 b = int3(0, 2, 0);
-  int3 const r = (a % as_type<int3>((as_type<uint3>(as_type<int3>(b)) + as_type<uint3>(as_type<int3>(b)))));
+  int3 const r = (a % as_type<int3>((as_type<uint3>(int3(b)) + as_type<uint3>(int3(b)))));
   return;
 }
 
diff --git a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
index 4ed1ec2..2d83cd3 100644
--- a/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/mod_by_zero/by_expression/vec3-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
 kernel void f() {
   int3 a = int3(1, 2, 3);
   int3 b = int3(0, 5, 0);
-  int3 const r = (a % as_type<int3>((as_type<uint3>(as_type<int3>(b)) + as_type<uint3>(as_type<int3>(b)))));
+  int3 const r = (a % as_type<int3>((as_type<uint3>(int3(b)) + as_type<uint3>(int3(b)))));
   return;
 }
 
diff --git a/test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.msl
index 71e20e6..0a492f4 100644
--- a/test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/mul/scalar-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
 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>(as_type<int3>(b))));
+  int3 const r = as_type<int3>((as_type<uint>(a) * as_type<uint3>(int3(b))));
   return;
 }
 
diff --git a/test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.msl b/test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.msl
index 8ddccd9..134ba04 100644
--- a/test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/mul/vec3-scalar/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
 kernel void f() {
   int3 const a = int3(1, 2, 3);
   int const b = 4;
-  int3 const r = as_type<int3>((as_type<uint3>(as_type<int3>(a)) * as_type<uint>(b)));
+  int3 const r = as_type<int3>((as_type<uint3>(int3(a)) * as_type<uint>(b)));
   return;
 }
 
diff --git a/test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.msl
index 554bc14..dbbcac5 100644
--- a/test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/mul/vec3-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
 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>(as_type<int3>(a)) * as_type<uint3>(as_type<int3>(b))));
+  int3 const r = as_type<int3>((as_type<uint3>(int3(a)) * as_type<uint3>(int3(b))));
   return;
 }
 
diff --git a/test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.msl
index 2021a7a..5403240 100644
--- a/test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/sub/scalar-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
 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>(as_type<int3>(b))));
+  int3 const r = as_type<int3>((as_type<uint>(a) - as_type<uint3>(int3(b))));
   return;
 }
 
diff --git a/test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.msl b/test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.msl
index efe23d7..9e59fcd 100644
--- a/test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/sub/vec3-scalar/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
 kernel void f() {
   int3 const a = int3(1, 2, 3);
   int const b = 4;
-  int3 const r = as_type<int3>((as_type<uint3>(as_type<int3>(a)) - as_type<uint>(b)));
+  int3 const r = as_type<int3>((as_type<uint3>(int3(a)) - as_type<uint>(b)));
   return;
 }
 
diff --git a/test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.msl b/test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.msl
index 2cbef7a..a42ec6c 100644
--- a/test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/sub/vec3-vec3/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
 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>(as_type<int3>(a)) - as_type<uint3>(as_type<int3>(b))));
+  int3 const r = as_type<int3>((as_type<uint3>(int3(a)) - as_type<uint3>(int3(b))));
   return;
 }
 
diff --git a/test/tint/statements/compound_assign/vector/minus.wgsl.expected.msl b/test/tint/statements/compound_assign/vector/minus.wgsl.expected.msl
index 7984259..d6ddc12 100644
--- a/test/tint/statements/compound_assign/vector/minus.wgsl.expected.msl
+++ b/test/tint/statements/compound_assign/vector/minus.wgsl.expected.msl
@@ -6,6 +6,6 @@
 };
 
 void foo(device S* const tint_symbol) {
-  (*(tint_symbol)).a = as_type<int4>((as_type<uint4>(as_type<int4>((*(tint_symbol)).a)) - as_type<uint4>(as_type<int4>(int4(2)))));
+  (*(tint_symbol)).a = as_type<int4>((as_type<uint4>(int4((*(tint_symbol)).a)) - as_type<uint4>(int4(int4(2)))));
 }
 
diff --git a/test/tint/statements/compound_assign/vector/plus.wgsl.expected.msl b/test/tint/statements/compound_assign/vector/plus.wgsl.expected.msl
index 9795878..455e61d 100644
--- a/test/tint/statements/compound_assign/vector/plus.wgsl.expected.msl
+++ b/test/tint/statements/compound_assign/vector/plus.wgsl.expected.msl
@@ -6,6 +6,6 @@
 };
 
 void foo(device S* const tint_symbol) {
-  (*(tint_symbol)).a = as_type<int4>((as_type<uint4>(as_type<int4>((*(tint_symbol)).a)) + as_type<uint4>(as_type<int4>(int4(2)))));
+  (*(tint_symbol)).a = as_type<int4>((as_type<uint4>(int4((*(tint_symbol)).a)) + as_type<uint4>(int4(int4(2)))));
 }
 
diff --git a/test/tint/statements/compound_assign/vector/times.wgsl.expected.msl b/test/tint/statements/compound_assign/vector/times.wgsl.expected.msl
index 9fec5f3..3dfa83f 100644
--- a/test/tint/statements/compound_assign/vector/times.wgsl.expected.msl
+++ b/test/tint/statements/compound_assign/vector/times.wgsl.expected.msl
@@ -6,6 +6,6 @@
 };
 
 void foo(device S* const tint_symbol) {
-  (*(tint_symbol)).a = as_type<int4>((as_type<uint4>(as_type<int4>((*(tint_symbol)).a)) * as_type<uint4>(as_type<int4>(int4(2)))));
+  (*(tint_symbol)).a = as_type<int4>((as_type<uint4>(int4((*(tint_symbol)).a)) * as_type<uint4>(int4(int4(2)))));
 }