diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 6560b6f..d2c69e1 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -517,6 +517,8 @@
     "transform/multiplanar_external_texture.h",
     "transform/num_workgroups_from_uniform.cc",
     "transform/num_workgroups_from_uniform.h",
+    "transform/packed_vec3.cc",
+    "transform/packed_vec3.h",
     "transform/pad_structs.cc",
     "transform/pad_structs.h",
     "transform/promote_initializers_to_let.cc",
@@ -1225,6 +1227,7 @@
       "transform/module_scope_var_to_entry_point_param_test.cc",
       "transform/multiplanar_external_texture_test.cc",
       "transform/num_workgroups_from_uniform_test.cc",
+      "transform/packed_vec3_test.cc",
       "transform/pad_structs_test.cc",
       "transform/promote_initializers_to_let_test.cc",
       "transform/promote_side_effects_to_decl_test.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index af4384c..47b44a5 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -430,6 +430,8 @@
   transform/multiplanar_external_texture.h
   transform/num_workgroups_from_uniform.cc
   transform/num_workgroups_from_uniform.h
+  transform/packed_vec3.cc
+  transform/packed_vec3.h
   transform/pad_structs.cc
   transform/pad_structs.h
   transform/promote_initializers_to_let.cc
@@ -1141,6 +1143,7 @@
       transform/module_scope_var_to_entry_point_param_test.cc
       transform/multiplanar_external_texture_test.cc
       transform/num_workgroups_from_uniform_test.cc
+      transform/packed_vec3_test.cc
       transform/pad_structs_test.cc
       transform/promote_initializers_to_let_test.cc
       transform/promote_side_effects_to_decl_test.cc
diff --git a/src/tint/transform/packed_vec3.cc b/src/tint/transform/packed_vec3.cc
new file mode 100644
index 0000000..dde5aca
--- /dev/null
+++ b/src/tint/transform/packed_vec3.cc
@@ -0,0 +1,194 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/tint/transform/packed_vec3.h"
+
+#include <algorithm>
+#include <string>
+#include <utility>
+
+#include "src/tint/program_builder.h"
+#include "src/tint/sem/index_accessor_expression.h"
+#include "src/tint/sem/member_accessor_expression.h"
+#include "src/tint/sem/statement.h"
+#include "src/tint/sem/variable.h"
+#include "src/tint/utils/hashmap.h"
+#include "src/tint/utils/hashset.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::transform::PackedVec3);
+TINT_INSTANTIATE_TYPEINFO(tint::transform::PackedVec3::Attribute);
+
+using namespace tint::number_suffixes;  // NOLINT
+
+namespace tint::transform {
+
+/// The PIMPL state for the PackedVec3 transform
+struct PackedVec3::State {
+    /// Constructor
+    /// @param c the CloneContext
+    explicit State(CloneContext& c) : ctx(c) {}
+
+    /// Runs the transform
+    void Run() {
+        // Packed vec3<T> struct members
+        utils::Hashset<const sem::StructMember*, 8> members;
+
+        // Find all the packed vector struct members, and apply the @internal(packed_vector)
+        // attribute.
+        for (auto* decl : ctx.src->AST().GlobalDeclarations()) {
+            if (auto* str = sem.Get<sem::Struct>(decl)) {
+                if (str->IsHostShareable()) {
+                    for (auto* member : str->Members()) {
+                        if (auto* vec = member->Type()->As<sem::Vector>()) {
+                            if (vec->Width() == 3) {
+                                members.Add(member);
+
+                                // Apply the PackedVec3::Attribute to the member
+                                auto* member_decl = member->Declaration();
+                                auto name = ctx.Clone(member_decl->symbol);
+                                auto* type = ctx.Clone(member_decl->type);
+                                utils::Vector<const ast::Attribute*, 4> attrs{
+                                    b.ASTNodes().Create<Attribute>(b.ID(), b.AllocateNodeID()),
+                                };
+                                for (auto* attr : member_decl->attributes) {
+                                    attrs.Push(ctx.Clone(attr));
+                                }
+                                ctx.Replace(member_decl, b.Member(name, type, std::move(attrs)));
+                            }
+                        }
+                    }
+                }
+            }
+        }
+
+        // Walk the nodes, starting with the most deeply nested, finding all the AST expressions
+        // that load a whole packed vector (not a scalar / swizzle of the vector).
+        utils::Hashset<const sem::Expression*, 16> refs;
+        for (auto* node : ctx.src->ASTNodes().Objects()) {
+            Switch(
+                sem.Get(node),  //
+                [&](const sem::StructMemberAccess* access) {
+                    if (members.Contains(access->Member())) {
+                        // Access to a packed vector member. Seed the expression tracking.
+                        refs.Add(access);
+                    }
+                },
+                [&](const sem::IndexAccessorExpression* access) {
+                    // Not loading a whole packed vector. Ignore.
+                    refs.Remove(access->Object());
+                },
+                [&](const sem::Swizzle* access) {
+                    // Not loading a whole packed vector. Ignore.
+                    refs.Remove(access->Object());
+                },
+                [&](const sem::VariableUser* user) {
+                    auto* v = user->Variable();
+                    if (v->Declaration()->Is<ast::Let>() &&  // if variable is let...
+                        v->Type()->Is<sem::Pointer>() &&     // and let is a pointer...
+                        refs.Contains(v->Initializer())) {   // and pointer is to a packed vector...
+                        refs.Add(user);  // then propagate tracking to pointer usage
+                    }
+                },
+                [&](const sem::Expression* expr) {
+                    if (auto* unary = expr->Declaration()->As<ast::UnaryOpExpression>()) {
+                        if (unary->op == ast::UnaryOp::kAddressOf ||
+                            unary->op == ast::UnaryOp::kIndirection) {
+                            // Memory access on the packed vector. Track these.
+                            auto* inner = sem.Get(unary->expr);
+                            if (refs.Remove(inner)) {
+                                refs.Add(expr);
+                            }
+                        }
+                        // Note: non-memory ops (e.g. '-') are ignored, leaving any tracked
+                        // reference at the inner expression, so we'd cast, then apply the unary op.
+                    }
+                },
+                [&](const sem::Statement* e) {
+                    if (auto* assign = e->Declaration()->As<ast::AssignmentStatement>()) {
+                        // We don't want to cast packed_vectors if they're being assigned to.
+                        refs.Remove(sem.Get(assign->lhs));
+                    }
+                });
+        }
+
+        // Wrap the load expressions with a cast to the unpacked type.
+        utils::Hashmap<const sem::Vector*, Symbol, 3> unpack_fns;
+        for (auto* ref : refs) {
+            // ref is either a packed vec3 that needs casting, or a pointer to a vec3 which we just
+            // leave alone.
+            if (auto* vec_ty = ref->Type()->UnwrapRef()->As<sem::Vector>()) {
+                auto* expr = ref->Declaration();
+                ctx.Replace(expr, [this, vec_ty, expr] {  //
+                    auto* packed = ctx.CloneWithoutTransform(expr);
+                    return b.Construct(CreateASTTypeFor(ctx, vec_ty), packed);
+                });
+            }
+        }
+
+        ctx.Clone();
+    }
+
+    /// @returns true if this transform should be run for the given program
+    /// @param program the program to inspect
+    static bool ShouldRun(const Program* program) {
+        for (auto* decl : program->AST().GlobalDeclarations()) {
+            if (auto* str = program->Sem().Get<sem::Struct>(decl)) {
+                if (str->IsHostShareable()) {
+                    for (auto* member : str->Members()) {
+                        if (auto* vec = member->Type()->As<sem::Vector>()) {
+                            if (vec->Width() == 3) {
+                                return true;
+                            }
+                        }
+                    }
+                }
+            }
+        }
+        return false;
+    }
+
+  private:
+    /// The clone context
+    CloneContext& ctx;
+    /// Alias to the semantic info in ctx.src
+    const sem::Info& sem = ctx.src->Sem();
+    /// Alias to the symbols in ctx.src
+    const SymbolTable& sym = ctx.src->Symbols();
+    /// Alias to the ctx.dst program builder
+    ProgramBuilder& b = *ctx.dst;
+};
+
+PackedVec3::Attribute::Attribute(ProgramID pid, ast::NodeID nid) : Base(pid, nid) {}
+PackedVec3::Attribute::~Attribute() = default;
+
+const PackedVec3::Attribute* PackedVec3::Attribute::Clone(CloneContext* ctx) const {
+    return ctx->dst->ASTNodes().Create<Attribute>(ctx->dst->ID(), ctx->dst->AllocateNodeID());
+}
+
+std::string PackedVec3::Attribute::InternalName() const {
+    return "packed_vector";
+}
+
+PackedVec3::PackedVec3() = default;
+PackedVec3::~PackedVec3() = default;
+
+bool PackedVec3::ShouldRun(const Program* program, const DataMap&) const {
+    return State::ShouldRun(program);
+}
+
+void PackedVec3::Run(CloneContext& ctx, const DataMap&, DataMap&) const {
+    State(ctx).Run();
+}
+
+}  // namespace tint::transform
diff --git a/src/tint/transform/packed_vec3.h b/src/tint/transform/packed_vec3.h
new file mode 100644
index 0000000..9d899cb
--- /dev/null
+++ b/src/tint/transform/packed_vec3.h
@@ -0,0 +1,78 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_TINT_TRANSFORM_PACKED_VEC3_H_
+#define SRC_TINT_TRANSFORM_PACKED_VEC3_H_
+
+#include <string>
+
+#include "src/tint/ast/internal_attribute.h"
+#include "src/tint/transform/transform.h"
+
+namespace tint::transform {
+
+/// A transform to be used by the MSL backend which will:
+/// * Apply the `@internal('packed_vector')` attribute (PackedVec3::Attribute) to all host-sharable
+///   structure members that have a vec3<T> type.
+/// * Cast all direct (not sub-accessed) loads of these packed vectors to the 'unpacked' vec3<T>
+///   type before usage.
+///
+/// This transform papers over overload holes in the MSL standard library where an MSL
+/// `packed_vector` type cannot be interchangable used as a regular `vec` type.
+class PackedVec3 final : public Castable<PackedVec3, Transform> {
+  public:
+    /// Attribute is the attribute applied to padded vector structure members.
+    class Attribute final : public Castable<Attribute, ast::InternalAttribute> {
+      public:
+        /// Constructor
+        /// @param pid the identifier of the program that owns this node
+        /// @param nid the unique node identifier
+        Attribute(ProgramID pid, ast::NodeID nid);
+        /// Destructor
+        ~Attribute() override;
+
+        /// @returns "packed_vector".
+        std::string InternalName() const override;
+
+        /// Performs a deep clone of this object using the CloneContext `ctx`.
+        /// @param ctx the clone context
+        /// @return the newly cloned object
+        const Attribute* Clone(CloneContext* ctx) const override;
+    };
+
+    /// Constructor
+    PackedVec3();
+    /// Destructor
+    ~PackedVec3() override;
+
+    /// @param program the program to inspect
+    /// @param data optional extra transform-specific input data
+    /// @returns true if this transform should be run for the given program
+    bool ShouldRun(const Program* program, const DataMap& data = {}) const override;
+
+  private:
+    struct State;
+
+    /// Runs the transform using the CloneContext built for transforming a
+    /// program. Run() is responsible for calling Clone() on the CloneContext.
+    /// @param ctx the CloneContext primed with the input program and
+    /// ProgramBuilder
+    /// @param inputs optional extra transform-specific input data
+    /// @param outputs optional extra transform-specific output data
+    void Run(CloneContext& ctx, const DataMap& inputs, DataMap& outputs) const override;
+};
+
+}  // namespace tint::transform
+
+#endif  // SRC_TINT_TRANSFORM_PACKED_VEC3_H_
diff --git a/src/tint/transform/packed_vec3_test.cc b/src/tint/transform/packed_vec3_test.cc
new file mode 100644
index 0000000..0f5c92e
--- /dev/null
+++ b/src/tint/transform/packed_vec3_test.cc
@@ -0,0 +1,662 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//     http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/tint/transform/packed_vec3.h"
+
+#include <string>
+#include <utility>
+#include <vector>
+
+#include "src/tint/transform/test_helper.h"
+#include "src/tint/utils/string.h"
+
+namespace tint::transform {
+namespace {
+
+using PackedVec3Test = TransformTest;
+
+TEST_F(PackedVec3Test, ShouldRun_EmptyModule) {
+    auto* src = R"()";
+
+    EXPECT_FALSE(ShouldRun<PackedVec3>(src));
+}
+
+TEST_F(PackedVec3Test, ShouldRun_NonHostSharableStruct) {
+    auto* src = R"(
+struct S {
+  v : vec3<f32>,
+}
+
+fn f() {
+  var v : S; // function address-space - not host sharable
+}
+)";
+
+    EXPECT_FALSE(ShouldRun<PackedVec3>(src));
+}
+
+TEST_F(PackedVec3Test, ShouldRun_Vec4Vec2) {
+    auto* src = R"(
+struct S {
+  v4 : vec4<f32>,
+  v2 : vec2<f32>,
+}
+
+@group(0) @binding(0) var<uniform> P : S; // Host sharable
+)";
+
+    EXPECT_FALSE(ShouldRun<PackedVec3>(src));
+}
+
+TEST_F(PackedVec3Test, ShouldRun_HostSharableStruct) {
+    auto* src = R"(
+struct S {
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<uniform> P : S; // Host sharable
+)";
+
+    EXPECT_TRUE(ShouldRun<PackedVec3>(src));
+}
+
+TEST_F(PackedVec3Test, UniformAddressSpace) {
+    auto* src = R"(
+struct S {
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<uniform> P : S;
+
+fn f() {
+  let x = P.v;
+}
+)";
+
+    auto* expect = R"(
+struct S {
+  @internal(packed_vector)
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<uniform> P : S;
+
+fn f() {
+  let x = vec3<f32>(P.v);
+}
+)";
+
+    DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, StorageAddressSpace) {
+    auto* src = R"(
+struct S {
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = P.v;
+}
+)";
+
+    auto* expect = R"(
+struct S {
+  @internal(packed_vector)
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = vec3<f32>(P.v);
+}
+)";
+
+    DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ExistingMemberAttributes) {
+    auto* src = R"(
+struct S {
+  @align(32) @size(64) v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = P.v;
+}
+)";
+
+    auto* expect = R"(
+struct S {
+  @internal(packed_vector) @align(32) @size(64)
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = vec3<f32>(P.v);
+}
+)";
+
+    DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, MultipleVectors) {
+    auto* src = R"(
+struct S {
+  v2_a : vec2<f32>,
+  v3_a : vec3<f32>,
+  v4_a : vec4<f32>,
+  v2_b : vec2<f32>,
+  v3_b : vec3<f32>,
+  v4_b : vec4<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let v2_a = P.v2_a;
+  let v3_a = P.v3_a;
+  let v4_a = P.v4_a;
+  let v2_b = P.v2_b;
+  let v3_b = P.v3_b;
+  let v4_b = P.v4_b;
+}
+)";
+
+    auto* expect = R"(
+struct S {
+  v2_a : vec2<f32>,
+  @internal(packed_vector)
+  v3_a : vec3<f32>,
+  v4_a : vec4<f32>,
+  v2_b : vec2<f32>,
+  @internal(packed_vector)
+  v3_b : vec3<f32>,
+  v4_b : vec4<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let v2_a = P.v2_a;
+  let v3_a = vec3<f32>(P.v3_a);
+  let v4_a = P.v4_a;
+  let v2_b = P.v2_b;
+  let v3_b = vec3<f32>(P.v3_b);
+  let v4_b = P.v4_b;
+}
+)";
+
+    DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, MixedAddressSpace) {
+    auto* src = R"(
+struct S {
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  var f : S;
+  let x = f.v;
+}
+)";
+
+    auto* expect = R"(
+struct S {
+  @internal(packed_vector)
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  var f : S;
+  let x = vec3<f32>(f.v);
+}
+)";
+
+    DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadMemberAccessChain) {
+    auto* src = R"(
+struct S {
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = P.v.yz.x;
+}
+)";
+
+    auto* expect = R"(
+struct S {
+  @internal(packed_vector)
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = P.v.yz.x;
+}
+)";
+
+    DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadVector) {
+    auto* src = R"(
+struct S {
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = P.v;
+}
+)";
+
+    auto* expect = R"(
+struct S {
+  @internal(packed_vector)
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = vec3<f32>(P.v);
+}
+)";
+
+    DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadIndexAccessor) {
+    auto* src = R"(
+struct S {
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = P.v[1];
+}
+)";
+
+    auto* expect = R"(
+struct S {
+  @internal(packed_vector)
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = P.v[1];
+}
+)";
+
+    DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadViaStructPtrDirect) {
+    auto* src = R"(
+struct S {
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = (*(&(*(&P)))).v;
+}
+)";
+
+    auto* expect = R"(
+struct S {
+  @internal(packed_vector)
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = vec3<f32>((*(&(*(&(P))))).v);
+}
+)";
+
+    DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadViaVectorPtrDirect) {
+    auto* src = R"(
+struct S {
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = *(&(*(&(P.v))));
+}
+)";
+
+    auto* expect = R"(
+struct S {
+  @internal(packed_vector)
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = vec3<f32>(*(&(*(&(P.v)))));
+}
+)";
+
+    DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadViaStructPtrViaLet) {
+    auto* src = R"(
+struct S {
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let p0 = &P;
+  let p1 = &(*(p0));
+  let a = (*p1).v;
+  let p2 = &(*(p1));
+  let b = (*p2).v;
+  let c = (*p2).v;
+}
+)";
+
+    auto* expect = R"(
+struct S {
+  @internal(packed_vector)
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let p0 = &(P);
+  let p1 = &(*(p0));
+  let a = vec3<f32>((*(p1)).v);
+  let p2 = &(*(p1));
+  let b = vec3<f32>((*(p2)).v);
+  let c = vec3<f32>((*(p2)).v);
+}
+)";
+
+    DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadViaVectorPtrViaLet) {
+    auto* src = R"(
+struct S {
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let p0 = &(P.v);
+  let p1 = &(*(p0));
+  let a = *p1;
+  let p2 = &(*(p1));
+  let b = *p2;
+  let c = *p2;
+}
+)";
+
+    auto* expect = R"(
+struct S {
+  @internal(packed_vector)
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let p0 = &(P.v);
+  let p1 = &(*(p0));
+  let a = vec3<f32>(*(p1));
+  let p2 = &(*(p1));
+  let b = vec3<f32>(*(p2));
+  let c = vec3<f32>(*(p2));
+}
+)";
+
+    DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadUnaryOp) {
+    auto* src = R"(
+struct S {
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = -P.v;
+}
+)";
+
+    auto* expect = R"(
+struct S {
+  @internal(packed_vector)
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = -(vec3<f32>(P.v));
+}
+)";
+
+    DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ReadBinaryOp) {
+    auto* src = R"(
+struct S {
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = P.v + P.v;
+}
+)";
+
+    auto* expect = R"(
+struct S {
+  @internal(packed_vector)
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = (vec3<f32>(P.v) + vec3<f32>(P.v));
+}
+)";
+
+    DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, WriteVector) {
+    auto* src = R"(
+struct S {
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S;
+
+fn f() {
+  P.v = vec3(1.23);
+}
+)";
+
+    auto* expect = R"(
+struct S {
+  @internal(packed_vector)
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S;
+
+fn f() {
+  P.v = vec3(1.23);
+}
+)";
+
+    DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, WriteMemberAccess) {
+    auto* src = R"(
+struct S {
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S;
+
+fn f() {
+  P.v.y = 1.23;
+}
+)";
+
+    auto* expect = R"(
+struct S {
+  @internal(packed_vector)
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S;
+
+fn f() {
+  P.v.y = 1.23;
+}
+)";
+
+    DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, WriteIndexAccessor) {
+    auto* src = R"(
+struct S {
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S;
+
+fn f() {
+  P.v[1] = 1.23;
+}
+)";
+
+    auto* expect = R"(
+struct S {
+  @internal(packed_vector)
+  v : vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S;
+
+fn f() {
+  P.v[1] = 1.23;
+}
+)";
+
+    DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+}  // namespace
+}  // namespace tint::transform
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index 0bc8c39..cf67e10 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -66,6 +66,7 @@
 #include "src/tint/transform/expand_compound_assignment.h"
 #include "src/tint/transform/manager.h"
 #include "src/tint/transform/module_scope_var_to_entry_point_param.h"
+#include "src/tint/transform/packed_vec3.h"
 #include "src/tint/transform/promote_initializers_to_let.h"
 #include "src/tint/transform/promote_side_effects_to_decl.h"
 #include "src/tint/transform/remove_phonies.h"
@@ -154,32 +155,6 @@
     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;
@@ -259,6 +234,7 @@
     // it assumes that the form of the array length argument is &var.array.
     manager.Add<transform::ArrayLengthFromUniform>();
     manager.Add<transform::ModuleScopeVarToEntryPointParam>();
+    manager.Add<transform::PackedVec3>();
     data.Add<transform::ArrayLengthFromUniform::Config>(std::move(array_length_from_uniform_cfg));
     data.Add<transform::CanonicalizeEntryPointIO::Config>(std::move(entry_point_io_cfg));
     auto out = manager.Run(in, data);
@@ -554,18 +530,8 @@
         ScopedParen sp(out);
         {
             ScopedBitCast lhs_uint_cast(this, out, lhs_type, unsigned_type_of(target_type));
-
-            // 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()) {
-                ScopedCast lhs_self_cast(this, out, lhs_type, lhs_type);
-                if (!EmitExpression(out, expr->lhs)) {
-                    return false;
-                }
-            } else {
-                if (!EmitExpression(out, expr->lhs)) {
-                    return false;
-                }
+            if (!EmitExpression(out, expr->lhs)) {
+                return false;
             }
         }
         if (!emit_op()) {
@@ -573,18 +539,8 @@
         }
         {
             ScopedBitCast rhs_uint_cast(this, out, rhs_type, unsigned_type_of(target_type));
-
-            // 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()) {
-                ScopedCast rhs_self_cast(this, out, rhs_type, rhs_type);
-                if (!EmitExpression(out, expr->rhs)) {
-                    return false;
-                }
-            } else {
-                if (!EmitExpression(out, expr->rhs)) {
-                    return false;
-                }
+            if (!EmitExpression(out, expr->rhs)) {
+                return false;
             }
         }
         return true;
@@ -601,18 +557,8 @@
         ScopedParen sp(out);
         {
             ScopedBitCast lhs_uint_cast(this, out, lhs_type, unsigned_type_of(lhs_type));
-
-            // 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()) {
-                ScopedCast lhs_self_cast(this, out, lhs_type, lhs_type);
-                if (!EmitExpression(out, expr->lhs)) {
-                    return false;
-                }
-            } else {
-                if (!EmitExpression(out, expr->lhs)) {
-                    return false;
-                }
+            if (!EmitExpression(out, expr->lhs)) {
+                return false;
             }
         }
         if (!emit_op()) {
@@ -2780,41 +2726,6 @@
     return false;
 }
 
-bool GeneratorImpl::EmitPackedType(std::ostream& out,
-                                   const sem::Type* type,
-                                   const std::string& name) {
-    auto* vec = type->As<sem::Vector>();
-    if (vec && vec->Width() == 3) {
-        out << "packed_";
-        if (!EmitType(out, vec, "")) {
-            return false;
-        }
-
-        if (vec->is_float_vector() && !matrix_packed_vector_overloads_) {
-            // Overload operators for matrix-vector arithmetic where the vector
-            // operand is packed, as these overloads to not exist in the metal
-            // namespace.
-            TextBuffer b;
-            TINT_DEFER(helpers_.Append(b));
-            line(&b) << R"(template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-)";
-            matrix_packed_vector_overloads_ = true;
-        }
-
-        return true;
-    }
-
-    return EmitType(out, type, name);
-}
-
 bool GeneratorImpl::EmitStructType(TextBuffer* b, const sem::Struct* str) {
     line(b) << "struct " << StructName(str) << " {";
 
@@ -2861,14 +2772,15 @@
             }
 
             add_byte_offset_comment(out, msl_offset);
+        }
 
-            if (!EmitPackedType(out, mem->Type(), mem_name)) {
-                return false;
+        if (auto* decl = mem->Declaration()) {
+            if (ast::HasAttribute<transform::PackedVec3::Attribute>(decl->attributes)) {
+                out << "packed_";
             }
-        } else {
-            if (!EmitType(out, mem->Type(), mem_name)) {
-                return false;
-            }
+        }
+        if (!EmitType(out, mem->Type(), mem_name)) {
+            return false;
         }
 
         auto* ty = mem->Type();
@@ -2934,6 +2846,7 @@
                     [&](const ast::StructMemberOffsetAttribute*) { return true; },
                     [&](const ast::StructMemberAlignAttribute*) { return true; },
                     [&](const ast::StructMemberSizeAttribute*) { return true; },
+                    [&](const transform::PackedVec3::Attribute*) { return true; },
                     [&](Default) {
                         TINT_ICE(Writer, diagnostics_)
                             << "unhandled struct member attribute: " << attr->Name();
diff --git a/src/tint/writer/msl/generator_impl.h b/src/tint/writer/msl/generator_impl.h
index 188bfea..c3e1ac0 100644
--- a/src/tint/writer/msl/generator_impl.h
+++ b/src/tint/writer/msl/generator_impl.h
@@ -328,14 +328,6 @@
     /// @param sc the address space to generate
     /// @returns true if the address space is emitted
     bool EmitAddressSpace(std::ostream& out, ast::AddressSpace sc);
-    /// Handles generating an MSL-packed storage type.
-    /// If the type does not have a packed form, the standard non-packed form is
-    /// emitted.
-    /// @param out the output of the type stream
-    /// @param type the type to generate
-    /// @param name the name of the variable, only used for array emission
-    /// @returns true if the type is emitted
-    bool EmitPackedType(std::ostream& out, const sem::Type* type, const std::string& name);
     /// Handles generating a struct declaration
     /// @param buffer the text buffer that the type declaration will be written to
     /// @param str the struct to generate
@@ -431,8 +423,8 @@
     /// Non-empty only if an invariant attribute has been generated.
     std::string invariant_define_name_;
 
-    /// True if matrix-packed_vector operator overloads have been generated.
-    bool matrix_packed_vector_overloads_ = false;
+    /// The generated name for the packed vec3 type.
+    std::string packed_vec3_ty_;
 
     /// Unique name of the tint_array<T, N> template.
     /// Non-empty only if the template has been generated.
diff --git a/src/tint/writer/msl/generator_impl_binary_test.cc b/src/tint/writer/msl/generator_impl_binary_test.cc
index e08c420..fc03507 100644
--- a/src/tint/writer/msl/generator_impl_binary_test.cc
+++ b/src/tint/writer/msl/generator_impl_binary_test.cc
@@ -128,18 +128,14 @@
 }
 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))",
+    {R"(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)))",
+    {R"(((a >> b) >> b))", Op::kShiftRight},
+    {R"(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)))",
+    {R"(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)))",
+    {R"(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,
diff --git a/src/tint/writer/msl/generator_impl_type_test.cc b/src/tint/writer/msl/generator_impl_type_test.cc
index bc9767a..27de589 100644
--- a/src/tint/writer/msl/generator_impl_type_test.cc
+++ b/src/tint/writer/msl/generator_impl_type_test.cc
@@ -60,7 +60,7 @@
 // Size and alignments taken from the MSL spec:
 // https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
 DECLARE_TYPE(float2, 8, 8);
-DECLARE_TYPE(packed_float3, 12, 4);
+DECLARE_TYPE(float3, 12, 4);
 DECLARE_TYPE(float4, 16, 16);
 DECLARE_TYPE(float2x2, 16, 8);
 DECLARE_TYPE(float2x3, 32, 16);
@@ -301,7 +301,7 @@
     FIELD(0x0100, float2, 0, c)            \
     FIELD(0x0108, uint, 0, d)              \
     FIELD(0x010c, int8_t, 4, tint_pad_2)   \
-    FIELD(0x0110, packed_float3, 0, e)     \
+    FIELD(0x0110, float3, 0, e)            \
     FIELD(0x011c, uint, 0, f)              \
     FIELD(0x0120, float4, 0, g)            \
     FIELD(0x0130, uint, 0, h)              \
@@ -641,7 +641,7 @@
   /* 0x0100 */ float2 tint_pad_33;
   /* 0x0108 */ uint tint_pad_1;
   /* 0x010c */ tint_array<int8_t, 4> tint_pad_12;
-  /* 0x0110 */ packed_float3 tint_pad_3;
+  /* 0x0110 */ float3 tint_pad_3;
   /* 0x011c */ uint tint_pad_7;
   /* 0x0120 */ float4 tint_pad_25;
   /* 0x0130 */ uint tint_pad_5;
diff --git a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl
index 549e70f..52e5e7d 100644
--- a/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl
+++ b/test/tint/buffer/storage/dynamic_index/read.wgsl.expected.msl
@@ -2,16 +2,6 @@
 
 using namespace metal;
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 template<typename T, size_t N>
 struct tint_array {
     const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -42,11 +32,11 @@
 };
 
 void tint_symbol_inner(uint idx, const device S* const tint_symbol_1) {
-  int3 const a = (*(tint_symbol_1)).arr[idx].a;
+  int3 const a = int3((*(tint_symbol_1)).arr[idx].a);
   int const b = (*(tint_symbol_1)).arr[idx].b;
-  uint3 const c = (*(tint_symbol_1)).arr[idx].c;
+  uint3 const c = uint3((*(tint_symbol_1)).arr[idx].c);
   uint const d = (*(tint_symbol_1)).arr[idx].d;
-  float3 const e = (*(tint_symbol_1)).arr[idx].e;
+  float3 const e = float3((*(tint_symbol_1)).arr[idx].e);
   float const f = (*(tint_symbol_1)).arr[idx].f;
   float2x3 const g = (*(tint_symbol_1)).arr[idx].g;
   float3x2 const h = (*(tint_symbol_1)).arr[idx].h;
diff --git a/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl b/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl
index e2430df..bb82069 100644
--- a/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl
+++ b/test/tint/buffer/storage/dynamic_index/write.wgsl.expected.msl
@@ -2,16 +2,6 @@
 
 using namespace metal;
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 template<typename T, size_t N>
 struct tint_array {
     const constant T& operator[](size_t i) const constant { return elements[i]; }
diff --git a/test/tint/buffer/storage/static_index/read.wgsl.expected.msl b/test/tint/buffer/storage/static_index/read.wgsl.expected.msl
index 8f891a2..db51912 100644
--- a/test/tint/buffer/storage/static_index/read.wgsl.expected.msl
+++ b/test/tint/buffer/storage/static_index/read.wgsl.expected.msl
@@ -2,16 +2,6 @@
 
 using namespace metal;
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 template<typename T, size_t N>
 struct tint_array {
     const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -43,11 +33,11 @@
 };
 
 kernel void tint_symbol(const device S* tint_symbol_1 [[buffer(0)]]) {
-  int3 const a = (*(tint_symbol_1)).a;
+  int3 const a = int3((*(tint_symbol_1)).a);
   int const b = (*(tint_symbol_1)).b;
-  uint3 const c = (*(tint_symbol_1)).c;
+  uint3 const c = uint3((*(tint_symbol_1)).c);
   uint const d = (*(tint_symbol_1)).d;
-  float3 const e = (*(tint_symbol_1)).e;
+  float3 const e = float3((*(tint_symbol_1)).e);
   float const f = (*(tint_symbol_1)).f;
   float2x3 const g = (*(tint_symbol_1)).g;
   float3x2 const h = (*(tint_symbol_1)).h;
diff --git a/test/tint/buffer/storage/static_index/write.wgsl.expected.msl b/test/tint/buffer/storage/static_index/write.wgsl.expected.msl
index 251e06f..1b21710 100644
--- a/test/tint/buffer/storage/static_index/write.wgsl.expected.msl
+++ b/test/tint/buffer/storage/static_index/write.wgsl.expected.msl
@@ -2,16 +2,6 @@
 
 using namespace metal;
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 template<typename T, size_t N>
 struct tint_array {
     const constant T& operator[](size_t i) const constant { return elements[i]; }
diff --git a/test/tint/buffer/uniform/dynamic_index/read.wgsl.expected.msl b/test/tint/buffer/uniform/dynamic_index/read.wgsl.expected.msl
index e361022..37817c3 100644
--- a/test/tint/buffer/uniform/dynamic_index/read.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/dynamic_index/read.wgsl.expected.msl
@@ -2,16 +2,6 @@
 
 using namespace metal;
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 template<typename T, size_t N>
 struct tint_array {
     const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -44,11 +34,11 @@
 };
 
 void tint_symbol_inner(uint idx, const constant S* const tint_symbol_1) {
-  int3 const a = (*(tint_symbol_1)).arr[idx].a;
+  int3 const a = int3((*(tint_symbol_1)).arr[idx].a);
   int const b = (*(tint_symbol_1)).arr[idx].b;
-  uint3 const c = (*(tint_symbol_1)).arr[idx].c;
+  uint3 const c = uint3((*(tint_symbol_1)).arr[idx].c);
   uint const d = (*(tint_symbol_1)).arr[idx].d;
-  float3 const e = (*(tint_symbol_1)).arr[idx].e;
+  float3 const e = float3((*(tint_symbol_1)).arr[idx].e);
   float const f = (*(tint_symbol_1)).arr[idx].f;
   int2 const g = (*(tint_symbol_1)).arr[idx].g;
   int2 const h = (*(tint_symbol_1)).arr[idx].h;
diff --git a/test/tint/buffer/uniform/static_index/read.wgsl.expected.msl b/test/tint/buffer/uniform/static_index/read.wgsl.expected.msl
index 26497bd..02d091d 100644
--- a/test/tint/buffer/uniform/static_index/read.wgsl.expected.msl
+++ b/test/tint/buffer/uniform/static_index/read.wgsl.expected.msl
@@ -14,16 +14,6 @@
     T elements[N];
 };
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 struct Inner {
   /* 0x0000 */ int x;
   /* 0x0004 */ tint_array<int8_t, 12> tint_pad;
@@ -46,11 +36,11 @@
 };
 
 kernel void tint_symbol(const constant S* tint_symbol_1 [[buffer(0)]]) {
-  int3 const a = (*(tint_symbol_1)).a;
+  int3 const a = int3((*(tint_symbol_1)).a);
   int const b = (*(tint_symbol_1)).b;
-  uint3 const c = (*(tint_symbol_1)).c;
+  uint3 const c = uint3((*(tint_symbol_1)).c);
   uint const d = (*(tint_symbol_1)).d;
-  float3 const e = (*(tint_symbol_1)).e;
+  float3 const e = float3((*(tint_symbol_1)).e);
   float const f = (*(tint_symbol_1)).f;
   int2 const g = (*(tint_symbol_1)).g;
   int2 const h = (*(tint_symbol_1)).h;
diff --git a/test/tint/bug/chromium/1273230.wgsl.expected.msl b/test/tint/bug/chromium/1273230.wgsl.expected.msl
index 08f7854..1e627d7 100644
--- a/test/tint/bug/chromium/1273230.wgsl.expected.msl
+++ b/test/tint/bug/chromium/1273230.wgsl.expected.msl
@@ -2,16 +2,6 @@
 
 using namespace metal;
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 template<typename T, size_t N>
 struct tint_array {
     const constant T& operator[](size_t i) const constant { return elements[i]; }
diff --git a/test/tint/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.msl b/test/tint/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.msl
index 6507f58..c63e9f1 100644
--- a/test/tint/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.msl
+++ b/test/tint/bug/fxc/indexed_assign_to_array_in_struct/1206.wgsl.expected.msl
@@ -14,16 +14,6 @@
     T elements[N];
 };
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 struct Simulation {
   /* 0x0000 */ uint i;
 };
diff --git a/test/tint/bug/tint/1113.wgsl.expected.msl b/test/tint/bug/tint/1113.wgsl.expected.msl
index b62d4c5..bf546c1 100644
--- a/test/tint/bug/tint/1113.wgsl.expected.msl
+++ b/test/tint/bug/tint/1113.wgsl.expected.msl
@@ -2,16 +2,6 @@
 
 using namespace metal;
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 template<typename T, size_t N>
 struct tint_array {
     const constant T& operator[](size_t i) const constant { return elements[i]; }
diff --git a/test/tint/bug/tint/1118.wgsl.expected.msl b/test/tint/bug/tint/1118.wgsl.expected.msl
index 40e92e8..5539c67 100644
--- a/test/tint/bug/tint/1118.wgsl.expected.msl
+++ b/test/tint/bug/tint/1118.wgsl.expected.msl
@@ -1,17 +1,6 @@
 #include <metal_stdlib>
 
 using namespace metal;
-
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 struct Scene {
   /* 0x0000 */ float4 vEyePosition;
 };
@@ -77,12 +66,12 @@
   shadow = 1.0f;
   refractionColor = float4(0.0f, 0.0f, 0.0f, 1.0f);
   reflectionColor = float4(0.0f, 0.0f, 0.0f, 1.0f);
-  float3 const x_94 = (*(tint_symbol_11)).vEmissiveColor;
+  float3 const x_94 = float3((*(tint_symbol_11)).vEmissiveColor);
   emissiveColor = x_94;
   float3 const x_96 = diffuseBase;
   float3 const x_97 = diffuseColor;
   float3 const x_99 = emissiveColor;
-  float3 const x_103 = (*(tint_symbol_11)).vAmbientColor;
+  float3 const x_103 = float3((*(tint_symbol_11)).vAmbientColor);
   float4 const x_108 = baseColor;
   finalDiffuse = (clamp((((x_96 * x_97) + x_99) + x_103), float3(0.0f), float3(1.0f)) * float3(x_108[0], x_108[1], x_108[2]));
   finalSpecular = float3(0.0f);
diff --git a/test/tint/bug/tint/1121.wgsl.expected.msl b/test/tint/bug/tint/1121.wgsl.expected.msl
index 34c70b2..22cc6dd 100644
--- a/test/tint/bug/tint/1121.wgsl.expected.msl
+++ b/test/tint/bug/tint/1121.wgsl.expected.msl
@@ -2,16 +2,6 @@
 
 using namespace metal;
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 template<typename T, size_t N>
 struct tint_array {
     const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -88,7 +78,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>(int2(tilePixel0Idx)) + as_type<uint2>(int2(int2(TILE_SIZE))))))) / float4((*(tint_symbol_3)).fullScreenSize).xy) - float2(1.0f));
+      float2 ceilCoord = (((2.0f * float2(as_type<int2>((as_type<uint2>(tilePixel0Idx) + as_type<uint2>(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 3b8e4f7..d9ff17c 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>(int4(x_44)) + as_type<uint4>(int4(int4(1)))));
+  int4 const x_47 = as_type<int4>((as_type<uint4>(x_44) + as_type<uint4>(int4(1))));
   val = x_47;
-  int4 const x_48 = as_type<int4>((as_type<uint4>(int4(x_47)) - as_type<uint4>(int4(int4(1)))));
+  int4 const x_48 = as_type<int4>((as_type<uint4>(x_47) - as_type<uint4>(int4(1))));
   val = x_48;
-  int4 const x_49 = as_type<int4>((as_type<uint4>(int4(x_48)) + as_type<uint4>(int4(int4(1)))));
+  int4 const x_49 = as_type<int4>((as_type<uint4>(x_48) + as_type<uint4>(int4(1))));
   val = x_49;
-  int4 const x_50 = as_type<int4>((as_type<uint4>(int4(x_49)) - as_type<uint4>(int4(int4(1)))));
+  int4 const x_50 = as_type<int4>((as_type<uint4>(x_49) - as_type<uint4>(int4(1))));
   val = x_50;
   x_55 = false;
   if (x_41) {
@@ -59,11 +59,11 @@
     x_55 = x_54;
   }
   ok = x_55;
-  int4 const x_58 = as_type<int4>((as_type<uint4>(int4(x_50)) * as_type<uint4>(int4(int4(2)))));
+  int4 const x_58 = as_type<int4>((as_type<uint4>(x_50) * as_type<uint4>(int4(2))));
   val = x_58;
   int4 const x_59 = (x_58 / int4(2));
   val = x_59;
-  int4 const x_60 = as_type<int4>((as_type<uint4>(int4(x_59)) * as_type<uint4>(int4(int4(2)))));
+  int4 const x_60 = as_type<int4>((as_type<uint4>(x_59) * as_type<uint4>(int4(2))));
   val = x_60;
   int4 const x_61 = (x_60 / int4(2));
   val = x_61;
diff --git a/test/tint/bug/tint/1534.wgsl b/test/tint/bug/tint/1534.wgsl
new file mode 100644
index 0000000..1b0c2a6
--- /dev/null
+++ b/test/tint/bug/tint/1534.wgsl
@@ -0,0 +1,17 @@
+
+struct g {
+  a : vec3<u32>,
+}
+
+struct h {
+  a : u32,
+}
+
+@group(0) @binding(0) var<uniform> i : g;
+
+@group(0) @binding(1) var<storage, read_write> j : h;
+
+@compute @workgroup_size(1) fn main() {
+  let l = dot(i.a, i.a);
+  j.a = i.a.x;
+}
diff --git a/test/tint/bug/tint/1534.wgsl.expected.msl b/test/tint/bug/tint/1534.wgsl.expected.msl
new file mode 100644
index 0000000..6f8a3de
--- /dev/null
+++ b/test/tint/bug/tint/1534.wgsl.expected.msl
@@ -0,0 +1,35 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+    const constant T& operator[](size_t i) const constant { return elements[i]; }
+    device T& operator[](size_t i) device { return elements[i]; }
+    const device T& operator[](size_t i) const device { return elements[i]; }
+    thread T& operator[](size_t i) thread { return elements[i]; }
+    const thread T& operator[](size_t i) const thread { return elements[i]; }
+    threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+    const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+    T elements[N];
+};
+
+template<typename T>
+T tint_dot3(vec<T,3> a, vec<T,3> b) {
+  return a[0]*b[0] + a[1]*b[1] + a[2]*b[2];
+}
+struct g {
+  /* 0x0000 */ packed_uint3 a;
+  /* 0x000c */ tint_array<int8_t, 4> tint_pad;
+};
+
+struct h {
+  /* 0x0000 */ uint a;
+};
+
+kernel void tint_symbol(const constant g* tint_symbol_1 [[buffer(0)]], device h* tint_symbol_2 [[buffer(1)]]) {
+  uint const l = tint_dot3(uint3((*(tint_symbol_1)).a), uint3((*(tint_symbol_1)).a));
+  (*(tint_symbol_2)).a = (*(tint_symbol_1)).a[0];
+  return;
+}
+
diff --git a/test/tint/bug/tint/1677.wgsl.expected.msl b/test/tint/bug/tint/1677.wgsl.expected.msl
index 6ee8bbd..580eba4 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>(int3((*(tint_symbol_1)).position)) - as_type<uint3>(int3(int3(0)))));
+  int3 const pos = as_type<int3>((as_type<uint3>(int3((*(tint_symbol_1)).position)) - as_type<uint3>(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/948.wgsl.expected.msl b/test/tint/bug/tint/948.wgsl.expected.msl
index 6113640..e77abf2 100644
--- a/test/tint/bug/tint/948.wgsl.expected.msl
+++ b/test/tint/bug/tint/948.wgsl.expected.msl
@@ -14,16 +14,6 @@
     T elements[N];
 };
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 struct LeftOver {
   /* 0x0000 */ float time;
   /* 0x0004 */ uint padding;
@@ -193,7 +183,7 @@
       i = as_type<int>((as_type<uint>(x_304) + as_type<uint>(1)));
     }
   }
-  float3 const x_310 = (*(tint_symbol_9)).colorMul;
+  float3 const x_310 = float3((*(tint_symbol_9)).colorMul);
   float4 const x_311 = color;
   float3 const x_313 = (float3(x_311[0], x_311[1], x_311[2]) * x_310);
   float4 const x_314 = color;
diff --git a/test/tint/bug/tint/949.wgsl.expected.msl b/test/tint/bug/tint/949.wgsl.expected.msl
index c7bdad3..44f5320 100644
--- a/test/tint/bug/tint/949.wgsl.expected.msl
+++ b/test/tint/bug/tint/949.wgsl.expected.msl
@@ -14,16 +14,6 @@
     T elements[N];
 };
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 struct lightingInfo {
   float3 diffuse;
   float3 specular;
@@ -248,7 +238,7 @@
   float4 const x_264 = tempTextureRead;
   float const x_273 = (*(tint_symbol_10)).textureInfoName;
   rgb = (float3(x_264[0], x_264[1], x_264[2]) * x_273);
-  float3 const x_279 = (*(tint_symbol_10)).u_cameraPosition;
+  float3 const x_279 = float3((*(tint_symbol_10)).u_cameraPosition);
   float4 const x_282 = *(tint_symbol_11);
   output5 = normalize((x_279 - float3(x_282[0], x_282[1], x_282[2])));
   output4 = float4(0.0f);
@@ -380,7 +370,7 @@
   tempTextureRead1 = x_475;
   float4 const x_477 = tempTextureRead1;
   rgb1 = float3(x_477[0], x_477[1], x_477[2]);
-  float3 const x_481 = (*(tint_symbol_10)).u_cameraPosition;
+  float3 const x_481 = float3((*(tint_symbol_10)).u_cameraPosition);
   float4 const x_482 = *(tint_symbol_11);
   viewDirectionW_1 = normalize((x_481 - float3(x_482[0], x_482[1], x_482[2])));
   shadow = 1.0f;
@@ -400,7 +390,7 @@
   param_14 = float3(x_510[0], x_510[1], x_510[2]);
   float4 const x_514 = (*(tint_symbol_17)).vLightSpecular;
   param_15 = float3(x_514[0], x_514[1], x_514[2]);
-  float3 const x_518 = (*(tint_symbol_17)).vLightGround;
+  float3 const x_518 = float3((*(tint_symbol_17)).vLightGround);
   param_16 = x_518;
   float const x_520 = glossiness_1;
   param_17 = x_520;
diff --git a/test/tint/bug/tint/980.wgsl.expected.msl b/test/tint/bug/tint/980.wgsl.expected.msl
index 2e23ef1..36cce9d 100644
--- a/test/tint/bug/tint/980.wgsl.expected.msl
+++ b/test/tint/bug/tint/980.wgsl.expected.msl
@@ -1,17 +1,6 @@
 #include <metal_stdlib>
 
 using namespace metal;
-
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 float3 Bad(uint index, float3 rd) {
   float3 normal = float3(0.0f);
   normal[index] = -(sign(rd[index]));
@@ -24,7 +13,7 @@
 };
 
 void tint_symbol_inner(uint idx, device S* const tint_symbol_2) {
-  float3 const tint_symbol_1 = Bad((*(tint_symbol_2)).i, (*(tint_symbol_2)).v);
+  float3 const tint_symbol_1 = Bad((*(tint_symbol_2)).i, float3((*(tint_symbol_2)).v));
   (*(tint_symbol_2)).v = tint_symbol_1;
 }
 
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 8658c03..46f1909 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>(int3(b))));
+  int3 const r = as_type<int3>((as_type<uint>(a) + as_type<uint3>(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 de4195f..89306cc 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>(int3(a)) + as_type<uint>(b)));
+  int3 const r = as_type<int3>((as_type<uint3>(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 009e87e..54d88c7 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>(int3(a)) + as_type<uint3>(int3(b))));
+  int3 const r = as_type<int3>((as_type<uint3>(a) + as_type<uint3>(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 6d2f8bc..72cb90c 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>(int3(b)) + as_type<uint3>(int3(b)))));
+  int3 const r = (a / as_type<int3>((as_type<uint3>(b) + as_type<uint3>(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 c90c00e..e2a2a76 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>(int3(b)) + as_type<uint3>(int3(b)))));
+  int3 const r = (a / as_type<int3>((as_type<uint3>(b) + as_type<uint3>(b))));
   return;
 }
 
diff --git a/test/tint/expressions/binary/left-shift/vector-vector/i32.wgsl.expected.msl b/test/tint/expressions/binary/left-shift/vector-vector/i32.wgsl.expected.msl
index 1e4e146..f97db25 100644
--- a/test/tint/expressions/binary/left-shift/vector-vector/i32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/left-shift/vector-vector/i32.wgsl.expected.msl
@@ -4,7 +4,7 @@
 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>(int3(a)) << b));
+  int3 const r = as_type<int3>((as_type<uint3>(a) << 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 22f6994..12f1a72 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>(int3(b)) + as_type<uint3>(int3(b)))));
+  int3 const r = (a % as_type<int3>((as_type<uint3>(b) + as_type<uint3>(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 2d83cd3..61356dc 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>(int3(b)) + as_type<uint3>(int3(b)))));
+  int3 const r = (a % as_type<int3>((as_type<uint3>(b) + as_type<uint3>(b))));
   return;
 }
 
diff --git a/test/tint/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.msl b/test/tint/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.msl
index e210c3c..a36de59 100644
--- a/test/tint/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/mul/mat3x2-vec3/f32.wgsl.expected.msl
@@ -14,16 +14,6 @@
     T elements[N];
 };
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 struct S {
   /* 0x0000 */ float3x2 tint_symbol;
   /* 0x0018 */ tint_array<int8_t, 8> tint_pad;
@@ -32,7 +22,7 @@
 };
 
 fragment void tint_symbol_1(const constant S* tint_symbol_2 [[buffer(0)]]) {
-  float2 const x = ((*(tint_symbol_2)).tint_symbol * (*(tint_symbol_2)).vector);
+  float2 const x = ((*(tint_symbol_2)).tint_symbol * float3((*(tint_symbol_2)).vector));
   return;
 }
 
diff --git a/test/tint/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.msl b/test/tint/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.msl
index 8f4ff55..79c3e09 100644
--- a/test/tint/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/mul/mat3x3-vec3/f32.wgsl.expected.msl
@@ -2,16 +2,6 @@
 
 using namespace metal;
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 template<typename T, size_t N>
 struct tint_array {
     const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -31,7 +21,7 @@
 };
 
 fragment void tint_symbol_1(const constant S* tint_symbol_2 [[buffer(0)]]) {
-  float3 const x = ((*(tint_symbol_2)).tint_symbol * (*(tint_symbol_2)).vector);
+  float3 const x = ((*(tint_symbol_2)).tint_symbol * float3((*(tint_symbol_2)).vector));
   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 0a492f4..216de39 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>(int3(b))));
+  int3 const r = as_type<int3>((as_type<uint>(a) * as_type<uint3>(b)));
   return;
 }
 
diff --git a/test/tint/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.msl b/test/tint/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.msl
index d9fc80b..321a379 100644
--- a/test/tint/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/mul/vec3-mat3x3/f32.wgsl.expected.msl
@@ -2,16 +2,6 @@
 
 using namespace metal;
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 template<typename T, size_t N>
 struct tint_array {
     const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -31,7 +21,7 @@
 };
 
 fragment void tint_symbol_1(const constant S* tint_symbol_2 [[buffer(0)]]) {
-  float3 const x = ((*(tint_symbol_2)).vector * (*(tint_symbol_2)).tint_symbol);
+  float3 const x = (float3((*(tint_symbol_2)).vector) * (*(tint_symbol_2)).tint_symbol);
   return;
 }
 
diff --git a/test/tint/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.msl b/test/tint/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.msl
index 3f66659..b1f6165 100644
--- a/test/tint/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.msl
+++ b/test/tint/expressions/binary/mul/vec3-mat4x3/f32.wgsl.expected.msl
@@ -2,16 +2,6 @@
 
 using namespace metal;
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 template<typename T, size_t N>
 struct tint_array {
     const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -31,7 +21,7 @@
 };
 
 fragment void tint_symbol_1(const constant S* tint_symbol_2 [[buffer(0)]]) {
-  float4 const x = ((*(tint_symbol_2)).vector * (*(tint_symbol_2)).tint_symbol);
+  float4 const x = (float3((*(tint_symbol_2)).vector) * (*(tint_symbol_2)).tint_symbol);
   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 134ba04..f320bd7 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>(int3(a)) * as_type<uint>(b)));
+  int3 const r = as_type<int3>((as_type<uint3>(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 dbbcac5..f4511b2 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>(int3(a)) * as_type<uint3>(int3(b))));
+  int3 const r = as_type<int3>((as_type<uint3>(a) * as_type<uint3>(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 5403240..8837be0 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>(int3(b))));
+  int3 const r = as_type<int3>((as_type<uint>(a) - as_type<uint3>(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 9e59fcd..6149a0c 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>(int3(a)) - as_type<uint>(b)));
+  int3 const r = as_type<int3>((as_type<uint3>(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 a42ec6c..142bb72 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>(int3(a)) - as_type<uint3>(int3(b))));
+  int3 const r = as_type<int3>((as_type<uint3>(a) - as_type<uint3>(b)));
   return;
 }
 
diff --git a/test/tint/expressions/swizzle/read/packed_vec3/f32.wgsl.expected.msl b/test/tint/expressions/swizzle/read/packed_vec3/f32.wgsl.expected.msl
index e2bbca3..1974ac0 100644
--- a/test/tint/expressions/swizzle/read/packed_vec3/f32.wgsl.expected.msl
+++ b/test/tint/expressions/swizzle/read/packed_vec3/f32.wgsl.expected.msl
@@ -2,16 +2,6 @@
 
 using namespace metal;
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 template<typename T, size_t N>
 struct tint_array {
     const constant T& operator[](size_t i) const constant { return elements[i]; }
@@ -30,7 +20,7 @@
 };
 
 void f(const constant S* const tint_symbol) {
-  float3 v = (*(tint_symbol)).v;
+  float3 v = float3((*(tint_symbol)).v);
   float x = (*(tint_symbol)).v[0];
   float y = (*(tint_symbol)).v[1];
   float z = (*(tint_symbol)).v[2];
diff --git a/test/tint/expressions/swizzle/read/packed_vec3/i32.wgsl.expected.msl b/test/tint/expressions/swizzle/read/packed_vec3/i32.wgsl.expected.msl
index 7a93a59..cfabec5 100644
--- a/test/tint/expressions/swizzle/read/packed_vec3/i32.wgsl.expected.msl
+++ b/test/tint/expressions/swizzle/read/packed_vec3/i32.wgsl.expected.msl
@@ -20,7 +20,7 @@
 };
 
 void f(const constant S* const tint_symbol) {
-  int3 v = (*(tint_symbol)).v;
+  int3 v = int3((*(tint_symbol)).v);
   int x = (*(tint_symbol)).v[0];
   int y = (*(tint_symbol)).v[1];
   int z = (*(tint_symbol)).v[2];
diff --git a/test/tint/expressions/swizzle/read/packed_vec3/u32.wgsl.expected.msl b/test/tint/expressions/swizzle/read/packed_vec3/u32.wgsl.expected.msl
index 7444f75..3f1899c 100644
--- a/test/tint/expressions/swizzle/read/packed_vec3/u32.wgsl.expected.msl
+++ b/test/tint/expressions/swizzle/read/packed_vec3/u32.wgsl.expected.msl
@@ -20,7 +20,7 @@
 };
 
 void f(const constant S* const tint_symbol) {
-  uint3 v = (*(tint_symbol)).v;
+  uint3 v = uint3((*(tint_symbol)).v);
   uint x = (*(tint_symbol)).v[0];
   uint y = (*(tint_symbol)).v[1];
   uint z = (*(tint_symbol)).v[2];
diff --git a/test/tint/expressions/swizzle/write/packed_vec3/f32.wgsl.expected.msl b/test/tint/expressions/swizzle/write/packed_vec3/f32.wgsl.expected.msl
index 7599cbc..8640cf6 100644
--- a/test/tint/expressions/swizzle/write/packed_vec3/f32.wgsl.expected.msl
+++ b/test/tint/expressions/swizzle/write/packed_vec3/f32.wgsl.expected.msl
@@ -2,16 +2,6 @@
 
 using namespace metal;
 
-template<typename T, int N, int M>
-inline vec<T, M> operator*(matrix<T, N, M> lhs, packed_vec<T, N> rhs) {
-  return lhs * vec<T, N>(rhs);
-}
-
-template<typename T, int N, int M>
-inline vec<T, N> operator*(packed_vec<T, M> lhs, matrix<T, N, M> rhs) {
-  return vec<T, M>(lhs) * rhs;
-}
-
 template<typename T, size_t N>
 struct tint_array {
     const constant T& operator[](size_t i) const constant { return elements[i]; }
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 d6ddc12..1a225d3 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>(int4((*(tint_symbol)).a)) - as_type<uint4>(int4(int4(2)))));
+  (*(tint_symbol)).a = as_type<int4>((as_type<uint4>((*(tint_symbol)).a) - as_type<uint4>(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 455e61d..7ea366b 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>(int4((*(tint_symbol)).a)) + as_type<uint4>(int4(int4(2)))));
+  (*(tint_symbol)).a = as_type<int4>((as_type<uint4>((*(tint_symbol)).a) + as_type<uint4>(int4(2))));
 }
 
diff --git a/test/tint/statements/compound_assign/vector/shift_left.wgsl.expected.msl b/test/tint/statements/compound_assign/vector/shift_left.wgsl.expected.msl
index 84a60be..2dd819c 100644
--- a/test/tint/statements/compound_assign/vector/shift_left.wgsl.expected.msl
+++ b/test/tint/statements/compound_assign/vector/shift_left.wgsl.expected.msl
@@ -6,6 +6,6 @@
 };
 
 void foo(device S* const tint_symbol) {
-  (*(tint_symbol)).a = as_type<int4>((as_type<uint4>(int4((*(tint_symbol)).a)) << uint4(2u)));
+  (*(tint_symbol)).a = as_type<int4>((as_type<uint4>((*(tint_symbol)).a) << uint4(2u)));
 }
 
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 3dfa83f..ef7fd06 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>(int4((*(tint_symbol)).a)) * as_type<uint4>(int4(int4(2)))));
+  (*(tint_symbol)).a = as_type<int4>((as_type<uint4>((*(tint_symbol)).a) * as_type<uint4>(int4(2))));
 }