diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 0405576..d09432f 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -553,6 +553,8 @@
     elseif(${DEPENDENCY} STREQUAL "gtest")
       target_include_directories(${TARGET} PRIVATE ${gmock_SOURCE_DIR}/include)
       target_link_libraries(${TARGET} PRIVATE gmock)
+    elseif(${DEPENDENCY} STREQUAL "libprotobuf-mutator")
+      target_link_libraries(${TARGET} PRIVATE libprotobuf-mutator)
     elseif(${DEPENDENCY} STREQUAL "metal")
       find_library(CoreGraphicsFramework CoreGraphics REQUIRED)
       find_library(FoundationFramework Foundation REQUIRED)
@@ -673,7 +675,8 @@
 foreach(PROTO_TARGET ${TINT_PROTO_TARGETS})
   generate_protos(
     TARGET ${PROTO_TARGET}
-    PROTOC_OUT_DIR "${DAWN_BUILD_GEN_DIR}/src/tint/")
+    IMPORT_DIRS "${TINT_ROOT_SOURCE_DIR}"
+    PROTOC_OUT_DIR "${DAWN_BUILD_GEN_DIR}")
   target_include_directories(${PROTO_TARGET} PRIVATE "${DAWN_BUILD_GEN_DIR}/src/tint/")
   target_include_directories(${PROTO_TARGET} PUBLIC "${DAWN_BUILD_GEN_DIR}")
   target_link_libraries(${PROTO_TARGET} libprotobuf)
diff --git a/src/tint/OWNERS b/src/tint/OWNERS
index 933a8fe..07dae49 100644
--- a/src/tint/OWNERS
+++ b/src/tint/OWNERS
@@ -1,5 +1,6 @@
 amaiorano@google.com
 bclayton@google.com
+chouinard@google.com
 dneto@google.com
 dsinclair@chromium.org
 jrprice@google.com
diff --git a/src/tint/cmd/test/BUILD.bazel b/src/tint/cmd/test/BUILD.bazel
index b6d308b..bc1d56f 100644
--- a/src/tint/cmd/test/BUILD.bazel
+++ b/src/tint/cmd/test/BUILD.bazel
@@ -118,8 +118,8 @@
     ":tint_build_msl_writer": [
       "//src/tint/lang/msl/writer/ast_printer:test",
       "//src/tint/lang/msl/writer/common:test",
-      "//src/tint/lang/msl/writer/printer:test",
       "//src/tint/lang/msl/writer/raise:test",
+      "//src/tint/lang/msl/writer:test",
     ],
     "//conditions:default": [],
   }) + select({
diff --git a/src/tint/cmd/test/BUILD.cmake b/src/tint/cmd/test/BUILD.cmake
index 94c5d90..098d937 100644
--- a/src/tint/cmd/test/BUILD.cmake
+++ b/src/tint/cmd/test/BUILD.cmake
@@ -129,8 +129,8 @@
   tint_target_add_dependencies(tint_cmd_test_test_cmd test_cmd
     tint_lang_msl_writer_ast_printer_test
     tint_lang_msl_writer_common_test
-    tint_lang_msl_writer_printer_test
     tint_lang_msl_writer_raise_test
+    tint_lang_msl_writer_test
   )
 endif(TINT_BUILD_MSL_WRITER)
 
diff --git a/src/tint/cmd/test/BUILD.gn b/src/tint/cmd/test/BUILD.gn
index b8647d4..6e8b377 100644
--- a/src/tint/cmd/test/BUILD.gn
+++ b/src/tint/cmd/test/BUILD.gn
@@ -122,9 +122,9 @@
 
     if (tint_build_msl_writer) {
       deps += [
+        "${tint_src_dir}/lang/msl/writer:unittests",
         "${tint_src_dir}/lang/msl/writer/ast_printer:unittests",
         "${tint_src_dir}/lang/msl/writer/common:unittests",
-        "${tint_src_dir}/lang/msl/writer/printer:unittests",
         "${tint_src_dir}/lang/msl/writer/raise:unittests",
       ]
     }
diff --git a/src/tint/lang/core/ir/access.h b/src/tint/lang/core/ir/access.h
index 6ee0eb6..55a332c 100644
--- a/src/tint/lang/core/ir/access.h
+++ b/src/tint/lang/core/ir/access.h
@@ -58,10 +58,10 @@
     Access* Clone(CloneContext& ctx) override;
 
     /// @returns the object used for the access
-    Value* Object() { return operands_[kObjectOperandOffset]; }
+    Value* Object() { return Operand(kObjectOperandOffset); }
 
     /// @returns the object used for the access
-    const Value* Object() const { return operands_[kObjectOperandOffset]; }
+    const Value* Object() const { return Operand(kObjectOperandOffset); }
 
     /// Adds the given index to the end of the access chain
     /// @param idx the index to add
diff --git a/src/tint/lang/core/ir/binary.h b/src/tint/lang/core/ir/binary.h
index 378dae4..d4d83c2 100644
--- a/src/tint/lang/core/ir/binary.h
+++ b/src/tint/lang/core/ir/binary.h
@@ -67,16 +67,16 @@
     void SetOp(BinaryOp op) { op_ = op; }
 
     /// @returns the left-hand-side value for the instruction
-    Value* LHS() { return operands_[kLhsOperandOffset]; }
+    Value* LHS() { return Operand(kLhsOperandOffset); }
 
     /// @returns the left-hand-side value for the instruction
-    const Value* LHS() const { return operands_[kLhsOperandOffset]; }
+    const Value* LHS() const { return Operand(kLhsOperandOffset); }
 
     /// @returns the right-hand-side value for the instruction
-    Value* RHS() { return operands_[kRhsOperandOffset]; }
+    Value* RHS() { return Operand(kRhsOperandOffset); }
 
     /// @returns the right-hand-side value for the instruction
-    const Value* RHS() const { return operands_[kRhsOperandOffset]; }
+    const Value* RHS() const { return Operand(kRhsOperandOffset); }
 
     /// @returns the friendly name for the instruction
     std::string FriendlyName() const override { return "binary"; }
diff --git a/src/tint/lang/core/ir/bitcast.h b/src/tint/lang/core/ir/bitcast.h
index ad30e92..37bc7c0 100644
--- a/src/tint/lang/core/ir/bitcast.h
+++ b/src/tint/lang/core/ir/bitcast.h
@@ -54,10 +54,10 @@
     Bitcast* Clone(CloneContext& ctx) override;
 
     /// @returns the operand value
-    Value* Val() { return operands_[kValueOperandOffset]; }
+    Value* Val() { return Operand(kValueOperandOffset); }
 
     /// @returns the operand value
-    const Value* Val() const { return operands_[kValueOperandOffset]; }
+    const Value* Val() const { return Operand(kValueOperandOffset); }
 
     /// @returns the friendly name for the instruction
     std::string FriendlyName() const override { return "bitcast"; }
diff --git a/src/tint/lang/core/ir/break_if.h b/src/tint/lang/core/ir/break_if.h
index 090f22a..38bf9d8 100644
--- a/src/tint/lang/core/ir/break_if.h
+++ b/src/tint/lang/core/ir/break_if.h
@@ -74,10 +74,10 @@
     size_t ArgsOperandOffset() const override { return kArgsOperandOffset; }
 
     /// @returns the break condition
-    Value* Condition() { return operands_[kConditionOperandOffset]; }
+    Value* Condition() { return Operand(kConditionOperandOffset); }
 
     /// @returns the break condition
-    const Value* Condition() const { return operands_[kConditionOperandOffset]; }
+    const Value* Condition() const { return Operand(kConditionOperandOffset); }
 
     /// @returns the loop containing the break-if
     ir::Loop* Loop() { return loop_; }
diff --git a/src/tint/lang/core/ir/disassembly.cc b/src/tint/lang/core/ir/disassembly.cc
index 6909d45..bd48191 100644
--- a/src/tint/lang/core/ir/disassembly.cc
+++ b/src/tint/lang/core/ir/disassembly.cc
@@ -386,12 +386,6 @@
                 [&](const core::constant::Value* c) {
                     tint::Switch(
                         c,
-                        [&](const core::constant::Scalar<AFloat>* scalar) {
-                            out_ << StyleLiteral(scalar->ValueAs<AFloat>().value);
-                        },
-                        [&](const core::constant::Scalar<AInt>* scalar) {
-                            out_ << StyleLiteral(scalar->ValueAs<AInt>().value);
-                        },
                         [&](const core::constant::Scalar<i32>* scalar) {
                             out_ << StyleLiteral(scalar->ValueAs<i32>().value, "i");
                         },
@@ -423,7 +417,8 @@
                                 need_comma = true;
                             }
                             out_ << ")";
-                        });
+                        },
+                        TINT_ICE_ON_NO_MATCH);
                 };
             emit(constant->Value());
         },
@@ -567,7 +562,7 @@
 
 void Disassembly::EmitOperand(const Instruction* inst, size_t index) {
     SourceMarker marker(this);
-    EmitValue(inst->Operands()[index]);
+    EmitValue(inst->Operand(index));
     marker.Store(IndexedValue{inst, static_cast<uint32_t>(index)});
 }
 
diff --git a/src/tint/lang/core/ir/if.h b/src/tint/lang/core/ir/if.h
index 3a6e0d7..328168b 100644
--- a/src/tint/lang/core/ir/if.h
+++ b/src/tint/lang/core/ir/if.h
@@ -81,10 +81,10 @@
     void ForeachBlock(const std::function<void(const ir::Block*)>& cb) const override;
 
     /// @returns the if condition
-    Value* Condition() { return operands_[kConditionOperandOffset]; }
+    Value* Condition() { return Operand(kConditionOperandOffset); }
 
     /// @returns the if condition
-    const Value* Condition() const { return operands_[kConditionOperandOffset]; }
+    const Value* Condition() const { return Operand(kConditionOperandOffset); }
 
     /// @returns the true block
     ir::Block* True() { return true_; }
diff --git a/src/tint/lang/core/ir/let.h b/src/tint/lang/core/ir/let.h
index 96a318d..b02de3a 100644
--- a/src/tint/lang/core/ir/let.h
+++ b/src/tint/lang/core/ir/let.h
@@ -57,10 +57,10 @@
     void SetValue(ir::Value* value) { SetOperand(kValueOperandOffset, value); }
 
     /// @returns the value
-    ir::Value* Value() { return operands_[kValueOperandOffset]; }
+    ir::Value* Value() { return Operand(kValueOperandOffset); }
 
     /// @returns the value
-    const ir::Value* Value() const { return operands_[kValueOperandOffset]; }
+    const ir::Value* Value() const { return Operand(kValueOperandOffset); }
 
     /// @returns the friendly name for the instruction
     std::string FriendlyName() const override { return "let"; }
diff --git a/src/tint/lang/core/ir/load.h b/src/tint/lang/core/ir/load.h
index 6b65a1d..5a2af5d 100644
--- a/src/tint/lang/core/ir/load.h
+++ b/src/tint/lang/core/ir/load.h
@@ -55,10 +55,10 @@
     Load* Clone(CloneContext& ctx) override;
 
     /// @returns the value being loaded from
-    Value* From() { return operands_[kFromOperandOffset]; }
+    Value* From() { return Operand(kFromOperandOffset); }
 
     /// @returns the value being loaded from
-    const Value* From() const { return operands_[kFromOperandOffset]; }
+    const Value* From() const { return Operand(kFromOperandOffset); }
 
     /// @returns the friendly name for the instruction
     std::string FriendlyName() const override { return "load"; }
diff --git a/src/tint/lang/core/ir/load_vector_element.h b/src/tint/lang/core/ir/load_vector_element.h
index f94ddbf..1c6920a 100644
--- a/src/tint/lang/core/ir/load_vector_element.h
+++ b/src/tint/lang/core/ir/load_vector_element.h
@@ -58,16 +58,16 @@
     LoadVectorElement* Clone(CloneContext& ctx) override;
 
     /// @returns the vector pointer value
-    ir::Value* From() { return operands_[kFromOperandOffset]; }
+    ir::Value* From() { return Operand(kFromOperandOffset); }
 
     /// @returns the vector pointer value
-    const ir::Value* From() const { return operands_[kFromOperandOffset]; }
+    const ir::Value* From() const { return Operand(kFromOperandOffset); }
 
     /// @returns the new vector element index
-    ir::Value* Index() { return operands_[kIndexOperandOffset]; }
+    ir::Value* Index() { return Operand(kIndexOperandOffset); }
 
     /// @returns the new vector element index
-    const ir::Value* Index() const { return operands_[kIndexOperandOffset]; }
+    const ir::Value* Index() const { return Operand(kIndexOperandOffset); }
 
     /// @returns the friendly name for the instruction
     std::string FriendlyName() const override { return "load_vector_element"; }
diff --git a/src/tint/lang/core/ir/return.cc b/src/tint/lang/core/ir/return.cc
index d236eae..4f7608b 100644
--- a/src/tint/lang/core/ir/return.cc
+++ b/src/tint/lang/core/ir/return.cc
@@ -59,11 +59,11 @@
 }
 
 Function* Return::Func() {
-    return tint::As<Function>(operands_[kFunctionOperandOffset]);
+    return tint::As<Function>(Operand(kFunctionOperandOffset));
 }
 
 const Function* Return::Func() const {
-    return tint::As<Function>(operands_[kFunctionOperandOffset]);
+    return tint::As<Function>(Operand(kFunctionOperandOffset));
 }
 
 }  // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/store.h b/src/tint/lang/core/ir/store.h
index bc647c2..b14cbb1 100644
--- a/src/tint/lang/core/ir/store.h
+++ b/src/tint/lang/core/ir/store.h
@@ -57,19 +57,19 @@
     Store* Clone(CloneContext& ctx) override;
 
     /// @returns the value being stored too
-    Value* To() { return operands_[kToOperandOffset]; }
+    Value* To() { return Operand(kToOperandOffset); }
 
     /// @returns the value being stored too
-    const Value* To() const { return operands_[kToOperandOffset]; }
+    const Value* To() const { return Operand(kToOperandOffset); }
 
     /// @param to the value being stored too
     void SetTo(Value* to) { SetOperand(kToOperandOffset, to); }
 
     /// @returns the value being stored
-    Value* From() { return operands_[kFromOperandOffset]; }
+    Value* From() { return Operand(kFromOperandOffset); }
 
     /// @returns the value being stored
-    const Value* From() const { return operands_[kFromOperandOffset]; }
+    const Value* From() const { return Operand(kFromOperandOffset); }
 
     /// @returns the friendly name for the instruction
     std::string FriendlyName() const override { return "store"; }
diff --git a/src/tint/lang/core/ir/store_vector_element.h b/src/tint/lang/core/ir/store_vector_element.h
index 91e8697..9230bf0 100644
--- a/src/tint/lang/core/ir/store_vector_element.h
+++ b/src/tint/lang/core/ir/store_vector_element.h
@@ -61,22 +61,22 @@
     StoreVectorElement* Clone(CloneContext& ctx) override;
 
     /// @returns the vector pointer value
-    ir::Value* To() { return operands_[kToOperandOffset]; }
+    ir::Value* To() { return Operand(kToOperandOffset); }
 
     /// @returns the vector pointer value
-    const ir::Value* To() const { return operands_[kToOperandOffset]; }
+    const ir::Value* To() const { return Operand(kToOperandOffset); }
 
     /// @returns the new vector element index
-    ir::Value* Index() { return operands_[kIndexOperandOffset]; }
+    ir::Value* Index() { return Operand(kIndexOperandOffset); }
 
     /// @returns the new vector element index
-    const ir::Value* Index() const { return operands_[kIndexOperandOffset]; }
+    const ir::Value* Index() const { return Operand(kIndexOperandOffset); }
 
     /// @returns the new vector element value
-    ir::Value* Value() { return operands_[kValueOperandOffset]; }
+    ir::Value* Value() { return Operand(kValueOperandOffset); }
 
     /// @returns the new vector element value
-    const ir::Value* Value() const { return operands_[kValueOperandOffset]; }
+    const ir::Value* Value() const { return Operand(kValueOperandOffset); }
 
     /// @returns the friendly name for the instruction
     std::string FriendlyName() const override { return "store_vector_element"; }
diff --git a/src/tint/lang/core/ir/switch.h b/src/tint/lang/core/ir/switch.h
index c1b5f3a..151ecf3 100644
--- a/src/tint/lang/core/ir/switch.h
+++ b/src/tint/lang/core/ir/switch.h
@@ -105,10 +105,10 @@
     VectorRef<Case> Cases() const { return cases_; }
 
     /// @returns the condition
-    Value* Condition() { return operands_[kConditionOperandOffset]; }
+    Value* Condition() { return Operand(kConditionOperandOffset); }
 
     /// @returns the condition
-    const Value* Condition() const { return operands_[kConditionOperandOffset]; }
+    const Value* Condition() const { return Operand(kConditionOperandOffset); }
 
     /// @returns the friendly name for the instruction
     std::string FriendlyName() const override { return "switch"; }
diff --git a/src/tint/lang/core/ir/swizzle.h b/src/tint/lang/core/ir/swizzle.h
index 88740d0..efe85d0 100644
--- a/src/tint/lang/core/ir/swizzle.h
+++ b/src/tint/lang/core/ir/swizzle.h
@@ -56,10 +56,10 @@
     Swizzle* Clone(CloneContext& ctx) override;
 
     /// @returns the object used for the access
-    Value* Object() { return operands_[kObjectOperandOffset]; }
+    Value* Object() { return Operand(kObjectOperandOffset); }
 
     /// @returns the object used for the access
-    const Value* Object() const { return operands_[kObjectOperandOffset]; }
+    const Value* Object() const { return Operand(kObjectOperandOffset); }
 
     /// @returns the swizzle indices
     VectorRef<uint32_t> Indices() const { return indices_; }
diff --git a/src/tint/lang/core/ir/transform/binary_polyfill.cc b/src/tint/lang/core/ir/transform/binary_polyfill.cc
index e74a828..d07ffb6 100644
--- a/src/tint/lang/core/ir/transform/binary_polyfill.cc
+++ b/src/tint/lang/core/ir/transform/binary_polyfill.cc
@@ -125,7 +125,7 @@
     /// @param match the type to match the component count of
     /// @returns a value with the same number of vector components as @p match
     ir::Constant* MatchWidth(ir::Constant* element, const core::type::Type* match) {
-        if (auto* vec = match->As<core::type::Vector>()) {
+        if (match->Is<core::type::Vector>()) {
             return b.Splat(MatchWidth(element->Type(), match), element);
         }
         return element;
diff --git a/src/tint/lang/core/ir/transform/builtin_polyfill.cc b/src/tint/lang/core/ir/transform/builtin_polyfill.cc
index f62c7c6..6547d36 100644
--- a/src/tint/lang/core/ir/transform/builtin_polyfill.cc
+++ b/src/tint/lang/core/ir/transform/builtin_polyfill.cc
@@ -224,7 +224,7 @@
     /// @param match the type to match the component count of
     /// @returns a value with the same number of vector components as @p match
     ir::Constant* MatchWidth(ir::Constant* element, const core::type::Type* match) {
-        if (auto* vec = match->As<core::type::Vector>()) {
+        if (match->Is<core::type::Vector>()) {
             return b.Splat(MatchWidth(element->Type(), match), element);
         }
         return element;
diff --git a/src/tint/lang/core/ir/transform/conversion_polyfill.cc b/src/tint/lang/core/ir/transform/conversion_polyfill.cc
index 36ea764..e59ef30 100644
--- a/src/tint/lang/core/ir/transform/conversion_polyfill.cc
+++ b/src/tint/lang/core/ir/transform/conversion_polyfill.cc
@@ -206,7 +206,7 @@
     /// @param match the type to match the component count of
     /// @returns a value with the same number of vector components as @p match
     ir::Constant* MatchWidth(ir::Constant* element, const core::type::Type* match) {
-        if (auto* vec = match->As<core::type::Vector>()) {
+        if (match->Is<core::type::Vector>()) {
             return b.Splat(MatchWidth(element->Type(), match), element);
         }
         return element;
diff --git a/src/tint/lang/core/ir/transform/direct_variable_access_wgsl_test.cc b/src/tint/lang/core/ir/transform/direct_variable_access_wgsl_test.cc
index 374d980..dd5da1d 100644
--- a/src/tint/lang/core/ir/transform/direct_variable_access_wgsl_test.cc
+++ b/src/tint/lang/core/ir/transform/direct_variable_access_wgsl_test.cc
@@ -278,7 +278,7 @@
 
     auto* expect =
         R"(
-@group(0) @binding(0) var<uniform> U : array<array<array<vec4<i32>, 8u>, 8u>, 8u>;
+@group(0u) @binding(0u) var<uniform> U : array<array<array<vec4<i32>, 8u>, 8u>, 8u>;
 
 fn a(pre : i32, p_indices : array<u32, 3u>, post : i32) -> vec4<i32> {
   return U[p_indices[0u]][p_indices[1u]][p_indices[2u]];
@@ -344,7 +344,7 @@
 )";
 
     auto* expect = R"(
-@group(0) @binding(0) var<uniform> U : array<array<array<vec4<i32>, 8u>, 8u>, 8u>;
+@group(0u) @binding(0u) var<uniform> U : array<array<array<vec4<i32>, 8u>, 8u>, 8u>;
 
 var<private> i : i32;
 
@@ -425,7 +425,7 @@
 )";
 
     auto* expect = R"(
-@group(0) @binding(0) var<uniform> U : array<array<vec4<i32>, 8u>, 8u>;
+@group(0u) @binding(0u) var<uniform> U : array<array<vec4<i32>, 8u>, 8u>;
 
 var<private> i : i32;
 
@@ -503,7 +503,7 @@
 )";
 
     auto* expect = R"(
-@group(0) @binding(0) var<uniform> U : array<array<vec4<i32>, 8u>, 8u>;
+@group(0u) @binding(0u) var<uniform> U : array<array<vec4<i32>, 8u>, 8u>;
 
 var<private> i : i32;
 
@@ -585,7 +585,7 @@
 )";
 
     auto* expect = R"(
-@group(0) @binding(0) var<uniform> U : array<array<vec4<i32>, 8u>, 8u>;
+@group(0u) @binding(0u) var<uniform> U : array<array<vec4<i32>, 8u>, 8u>;
 
 var<private> i : i32;
 
@@ -667,7 +667,7 @@
 )";
 
     auto* expect = R"(
-@group(0) @binding(0) var<uniform> U : array<array<vec4<i32>, 8u>, 8u>;
+@group(0u) @binding(0u) var<uniform> U : array<array<vec4<i32>, 8u>, 8u>;
 
 var<private> i : i32;
 
@@ -733,7 +733,7 @@
 )";
 
     auto* expect = R"(
-@group(0) @binding(0) var<uniform> U : i32;
+@group(0u) @binding(0u) var<uniform> U : i32;
 
 fn a(pre : i32, post : i32) -> i32 {
   return U;
@@ -764,7 +764,7 @@
 )";
 
     auto* expect = R"(
-@group(0) @binding(0) var<uniform> U : array<vec4<i32>, 8u>;
+@group(0u) @binding(0u) var<uniform> U : array<vec4<i32>, 8u>;
 
 fn a(pre : i32, p_indices : array<u32, 1u>, post : i32) -> vec4<i32> {
   return U[p_indices[0u]];
@@ -852,7 +852,7 @@
   mat : mat3x4<f32>,
 }
 
-@group(0) @binding(0) var<uniform> U : Outer;
+@group(0u) @binding(0u) var<uniform> U : Outer;
 
 fn f0(p_indices : array<u32, 1u>) -> f32 {
   return U.mat[p_indices[0u]].x;
@@ -942,7 +942,7 @@
   i : i32,
 }
 
-@group(0) @binding(0) var<storage, read> S : str;
+@group(0u) @binding(0u) var<storage, read> S : str;
 
 fn a(pre : i32, post : i32) -> i32 {
   return S.i;
@@ -980,7 +980,7 @@
   arr : array<i32, 4u>,
 }
 
-@group(0) @binding(0) var<storage, read_write> S : str;
+@group(0u) @binding(0u) var<storage, read_write> S : str;
 
 fn a(pre : i32, post : i32) {
   S.arr = array<i32, 4u>();
@@ -1011,7 +1011,7 @@
 )";
 
     auto* expect = R"(
-@group(0) @binding(0) var<storage, read_write> S : array<vec4<i32>, 8u>;
+@group(0u) @binding(0u) var<storage, read_write> S : array<vec4<i32>, 8u>;
 
 fn a(pre : i32, p_indices : array<u32, 1u>, post : i32) {
   S[p_indices[0u]] = vec4<i32>();
@@ -1099,7 +1099,7 @@
   mat : mat3x4<f32>,
 }
 
-@group(0) @binding(0) var<storage, read> S : Outer;
+@group(0u) @binding(0u) var<storage, read> S : Outer;
 
 fn f0(p_indices : array<u32, 1u>) -> f32 {
   return S.mat[p_indices[0u]].x;
@@ -2117,7 +2117,7 @@
 )";
 
     auto* expect = R"(
-@group(0) @binding(0) var<storage, read> S : array<f32>;
+@group(0u) @binding(0u) var<storage, read> S : array<f32>;
 
 fn len() -> u32 {
   return arrayLength(&(S));
@@ -2242,25 +2242,25 @@
 )";
 
     auto* expect = R"(
-@group(0) @binding(0) var<uniform> U : vec4<i32>;
+@group(0u) @binding(0u) var<uniform> U : vec4<i32>;
 
 struct str {
   i : vec4<i32>,
 }
 
-@group(0) @binding(1) var<uniform> U_str : str;
+@group(0u) @binding(1u) var<uniform> U_str : str;
 
-@group(0) @binding(2) var<uniform> U_arr : array<vec4<i32>, 8u>;
+@group(0u) @binding(2u) var<uniform> U_arr : array<vec4<i32>, 8u>;
 
-@group(0) @binding(3) var<uniform> U_arr_arr : array<array<vec4<i32>, 8u>, 4u>;
+@group(0u) @binding(3u) var<uniform> U_arr_arr : array<array<vec4<i32>, 8u>, 4u>;
 
-@group(1) @binding(0) var<storage, read> S : vec4<i32>;
+@group(1u) @binding(0u) var<storage, read> S : vec4<i32>;
 
-@group(1) @binding(1) var<storage, read> S_str : str;
+@group(1u) @binding(1u) var<storage, read> S_str : str;
 
-@group(1) @binding(2) var<storage, read> S_arr : array<vec4<i32>, 8u>;
+@group(1u) @binding(2u) var<storage, read> S_arr : array<vec4<i32>, 8u>;
 
-@group(1) @binding(3) var<storage, read> S_arr_arr : array<array<vec4<i32>, 8u>, 4u>;
+@group(1u) @binding(3u) var<storage, read> S_arr_arr : array<array<vec4<i32>, 8u>, 4u>;
 
 var<workgroup> W : vec4<i32>;
 
@@ -2374,7 +2374,7 @@
 )";
 
     auto* expect = R"(
-@group(0) @binding(0) var<storage, read> S : array<array<array<array<i32, 9u>, 9u>, 9u>, 50u>;
+@group(0u) @binding(0u) var<storage, read> S : array<array<array<array<i32, 9u>, 9u>, 9u>, 50u>;
 
 fn a(i : i32) -> i32 {
   return i;
@@ -2415,7 +2415,7 @@
 )";
 
     auto* expect = R"(
-@group(0) @binding(0) var<storage, read> S : array<array<array<array<i32, 9u>, 9u>, 9u>, 50u>;
+@group(0u) @binding(0u) var<storage, read> S : array<array<array<array<i32, 9u>, 9u>, 9u>, 50u>;
 
 fn a(pre : i32, i_indices : array<u32, 4u>, post : i32) -> i32 {
   return S[i_indices[0u]][i_indices[1u]][i_indices[2u]][i_indices[3u]];
@@ -2457,9 +2457,9 @@
 )";
 
     auto* expect = R"(
-@group(0) @binding(0) var<storage, read> S : array<array<array<i32, 9u>, 9u>, 50u>;
+@group(0u) @binding(0u) var<storage, read> S : array<array<array<i32, 9u>, 9u>, 50u>;
 
-@group(0) @binding(0) var<uniform> U : array<array<array<vec4<i32>, 9u>, 9u>, 50u>;
+@group(0u) @binding(0u) var<uniform> U : array<array<array<vec4<i32>, 9u>, 9u>, 50u>;
 
 fn a(i : i32) -> i32 {
   return i;
diff --git a/src/tint/lang/core/ir/unary.h b/src/tint/lang/core/ir/unary.h
index 8ade151..739a574 100644
--- a/src/tint/lang/core/ir/unary.h
+++ b/src/tint/lang/core/ir/unary.h
@@ -57,10 +57,10 @@
     ~Unary() override;
 
     /// @returns the value for the instruction
-    Value* Val() { return operands_[kValueOperandOffset]; }
+    Value* Val() { return Operand(kValueOperandOffset); }
 
     /// @returns the value for the instruction
-    const Value* Val() const { return operands_[kValueOperandOffset]; }
+    const Value* Val() const { return Operand(kValueOperandOffset); }
 
     /// @returns the unary operator
     UnaryOp Op() const { return op_; }
diff --git a/src/tint/lang/core/ir/user_call.h b/src/tint/lang/core/ir/user_call.h
index a7a9b97..70fb815 100644
--- a/src/tint/lang/core/ir/user_call.h
+++ b/src/tint/lang/core/ir/user_call.h
@@ -66,10 +66,12 @@
     void SetArgs(VectorRef<Value*> arguments);
 
     /// @returns the called function
-    Function* Target() { return operands_[kFunctionOperandOffset]->As<ir::Function>(); }
+    Function* Target() { return tint::As<ir::Function>(Operand(kFunctionOperandOffset)); }
 
     /// @returns the called function
-    const Function* Target() const { return operands_[kFunctionOperandOffset]->As<ir::Function>(); }
+    const Function* Target() const {
+        return tint::As<ir::Function>(Operand(kFunctionOperandOffset));
+    }
 
     /// Sets called function
     /// @param target the new target of the call
diff --git a/src/tint/lang/core/ir/var.h b/src/tint/lang/core/ir/var.h
index d6a73d3..b3e06c4 100644
--- a/src/tint/lang/core/ir/var.h
+++ b/src/tint/lang/core/ir/var.h
@@ -73,9 +73,9 @@
     /// @param initializer the initializer
     void SetInitializer(Value* initializer);
     /// @returns the initializer
-    Value* Initializer() { return operands_[kInitializerOperandOffset]; }
+    Value* Initializer() { return Operand(kInitializerOperandOffset); }
     /// @returns the initializer
-    const Value* Initializer() const { return operands_[kInitializerOperandOffset]; }
+    const Value* Initializer() const { return Operand(kInitializerOperandOffset); }
 
     /// Sets the binding point
     /// @param group the group
diff --git a/src/tint/lang/hlsl/writer/ast_raise/truncate_interstage_variables.cc b/src/tint/lang/hlsl/writer/ast_raise/truncate_interstage_variables.cc
index fc9e35c..9918e7e 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/truncate_interstage_variables.cc
+++ b/src/tint/lang/hlsl/writer/ast_raise/truncate_interstage_variables.cc
@@ -112,6 +112,14 @@
 
         for (auto* member : str->Members()) {
             if (auto location = member->Attributes().location) {
+                const size_t kMaxLocation = data->interstage_locations.size() - 1u;
+                if (location.value() > kMaxLocation) {
+                    b.Diagnostics().AddError(Source{})
+                        << "The location (" << location.value() << ") of " << member->Name().Name()
+                        << " in " << str->Name().Name() << " exceeds the maximum value ("
+                        << kMaxLocation << ").";
+                    return resolver::Resolve(b);
+                }
                 if (!data->interstage_locations.test(location.value())) {
                     omit_members.Add(member);
                 }
diff --git a/src/tint/lang/hlsl/writer/ast_raise/truncate_interstage_variables_test.cc b/src/tint/lang/hlsl/writer/ast_raise/truncate_interstage_variables_test.cc
index b0e7d1a..4d42bc8 100644
--- a/src/tint/lang/hlsl/writer/ast_raise/truncate_interstage_variables_test.cc
+++ b/src/tint/lang/hlsl/writer/ast_raise/truncate_interstage_variables_test.cc
@@ -612,5 +612,32 @@
     EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(TruncateInterstageVariablesTest, LocationOutOfRange) {
+    auto* src = R"(
+struct ShaderIO {
+  @builtin(position) pos: vec4<f32>,
+  @location(0) f_0: f32,
+  @location(30) f_2: f32,
+}
+@vertex
+fn f() -> ShaderIO {
+  var io: ShaderIO;
+  io.f_0 = 1.0;
+  io.f_2 = io.f_2 + 3.0;
+  return io;
+}
+)";
+
+    // Return error when location >= 30 (maximum supported number of inter-stage shader variables)
+    auto* expect = "error: The location (30) of f_2 in ShaderIO exceeds the maximum value (29).";
+
+    TruncateInterstageVariables::Config cfg;
+    ast::transform::DataMap data;
+    data.Add<TruncateInterstageVariables::Config>(cfg);
+
+    auto got = Run<TruncateInterstageVariables>(src, data);
+    EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace tint::hlsl::writer
diff --git a/src/tint/lang/msl/writer/BUILD.bazel b/src/tint/lang/msl/writer/BUILD.bazel
index b2c8ad1..f7ccfaf 100644
--- a/src/tint/lang/msl/writer/BUILD.bazel
+++ b/src/tint/lang/msl/writer/BUILD.bazel
@@ -91,6 +91,54 @@
   visibility = ["//visibility:public"],
 )
 cc_library(
+  name = "test",
+  alwayslink = True,
+  srcs = [
+    "binary_test.cc",
+    "constant_test.cc",
+    "discard_test.cc",
+    "function_test.cc",
+    "helper_test.h",
+    "if_test.cc",
+    "let_test.cc",
+    "return_test.cc",
+    "type_test.cc",
+    "var_test.cc",
+  ],
+  deps = [
+    "//src/tint/api/common",
+    "//src/tint/api/options",
+    "//src/tint/lang/core",
+    "//src/tint/lang/core/constant",
+    "//src/tint/lang/core/intrinsic",
+    "//src/tint/lang/core/ir",
+    "//src/tint/lang/core/type",
+    "//src/tint/utils/containers",
+    "//src/tint/utils/diagnostic",
+    "//src/tint/utils/ice",
+    "//src/tint/utils/id",
+    "//src/tint/utils/macros",
+    "//src/tint/utils/math",
+    "//src/tint/utils/memory",
+    "//src/tint/utils/reflection",
+    "//src/tint/utils/result",
+    "//src/tint/utils/rtti",
+    "//src/tint/utils/symbol",
+    "//src/tint/utils/text",
+    "//src/tint/utils/traits",
+    "@gtest",
+  ] + select({
+    ":tint_build_msl_writer": [
+      "//src/tint/lang/msl/validate",
+      "//src/tint/lang/msl/writer",
+      "//src/tint/lang/msl/writer/common",
+    ],
+    "//conditions:default": [],
+  }),
+  copts = COPTS,
+  visibility = ["//visibility:public"],
+)
+cc_library(
   name = "bench",
   alwayslink = True,
   srcs = [
diff --git a/src/tint/lang/msl/writer/BUILD.cmake b/src/tint/lang/msl/writer/BUILD.cmake
index 18bec4c..0236187 100644
--- a/src/tint/lang/msl/writer/BUILD.cmake
+++ b/src/tint/lang/msl/writer/BUILD.cmake
@@ -101,6 +101,61 @@
 endif(TINT_BUILD_MSL_WRITER)
 if(TINT_BUILD_MSL_WRITER)
 ################################################################################
+# Target:    tint_lang_msl_writer_test
+# Kind:      test
+# Condition: TINT_BUILD_MSL_WRITER
+################################################################################
+tint_add_target(tint_lang_msl_writer_test test
+  lang/msl/writer/binary_test.cc
+  lang/msl/writer/constant_test.cc
+  lang/msl/writer/discard_test.cc
+  lang/msl/writer/function_test.cc
+  lang/msl/writer/helper_test.h
+  lang/msl/writer/if_test.cc
+  lang/msl/writer/let_test.cc
+  lang/msl/writer/return_test.cc
+  lang/msl/writer/type_test.cc
+  lang/msl/writer/var_test.cc
+)
+
+tint_target_add_dependencies(tint_lang_msl_writer_test test
+  tint_api_common
+  tint_api_options
+  tint_lang_core
+  tint_lang_core_constant
+  tint_lang_core_intrinsic
+  tint_lang_core_ir
+  tint_lang_core_type
+  tint_utils_containers
+  tint_utils_diagnostic
+  tint_utils_ice
+  tint_utils_id
+  tint_utils_macros
+  tint_utils_math
+  tint_utils_memory
+  tint_utils_reflection
+  tint_utils_result
+  tint_utils_rtti
+  tint_utils_symbol
+  tint_utils_text
+  tint_utils_traits
+)
+
+tint_target_add_external_dependencies(tint_lang_msl_writer_test test
+  "gtest"
+)
+
+if(TINT_BUILD_MSL_WRITER)
+  tint_target_add_dependencies(tint_lang_msl_writer_test test
+    tint_lang_msl_validate
+    tint_lang_msl_writer
+    tint_lang_msl_writer_common
+  )
+endif(TINT_BUILD_MSL_WRITER)
+
+endif(TINT_BUILD_MSL_WRITER)
+if(TINT_BUILD_MSL_WRITER)
+################################################################################
 # Target:    tint_lang_msl_writer_bench
 # Kind:      bench
 # Condition: TINT_BUILD_MSL_WRITER
diff --git a/src/tint/lang/msl/writer/BUILD.gn b/src/tint/lang/msl/writer/BUILD.gn
index e26ed1f..4fdf454 100644
--- a/src/tint/lang/msl/writer/BUILD.gn
+++ b/src/tint/lang/msl/writer/BUILD.gn
@@ -92,6 +92,55 @@
     }
   }
 }
+if (tint_build_unittests) {
+  if (tint_build_msl_writer) {
+    tint_unittests_source_set("unittests") {
+      sources = [
+        "binary_test.cc",
+        "constant_test.cc",
+        "discard_test.cc",
+        "function_test.cc",
+        "helper_test.h",
+        "if_test.cc",
+        "let_test.cc",
+        "return_test.cc",
+        "type_test.cc",
+        "var_test.cc",
+      ]
+      deps = [
+        "${tint_src_dir}:gmock_and_gtest",
+        "${tint_src_dir}/api/common",
+        "${tint_src_dir}/api/options",
+        "${tint_src_dir}/lang/core",
+        "${tint_src_dir}/lang/core/constant",
+        "${tint_src_dir}/lang/core/intrinsic",
+        "${tint_src_dir}/lang/core/ir",
+        "${tint_src_dir}/lang/core/type",
+        "${tint_src_dir}/utils/containers",
+        "${tint_src_dir}/utils/diagnostic",
+        "${tint_src_dir}/utils/ice",
+        "${tint_src_dir}/utils/id",
+        "${tint_src_dir}/utils/macros",
+        "${tint_src_dir}/utils/math",
+        "${tint_src_dir}/utils/memory",
+        "${tint_src_dir}/utils/reflection",
+        "${tint_src_dir}/utils/result",
+        "${tint_src_dir}/utils/rtti",
+        "${tint_src_dir}/utils/symbol",
+        "${tint_src_dir}/utils/text",
+        "${tint_src_dir}/utils/traits",
+      ]
+
+      if (tint_build_msl_writer) {
+        deps += [
+          "${tint_src_dir}/lang/msl/validate",
+          "${tint_src_dir}/lang/msl/writer",
+          "${tint_src_dir}/lang/msl/writer/common",
+        ]
+      }
+    }
+  }
+}
 if (tint_build_benchmarks) {
   if (tint_build_msl_writer) {
     tint_unittests_source_set("bench") {
diff --git a/src/tint/lang/msl/writer/printer/binary_test.cc b/src/tint/lang/msl/writer/binary_test.cc
similarity index 77%
rename from src/tint/lang/msl/writer/printer/binary_test.cc
rename to src/tint/lang/msl/writer/binary_test.cc
index fe0a464..32df3b4 100644
--- a/src/tint/lang/msl/writer/printer/binary_test.cc
+++ b/src/tint/lang/msl/writer/binary_test.cc
@@ -27,7 +27,7 @@
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
 #include "src/tint/lang/core/fluent_types.h"
-#include "src/tint/lang/msl/writer/printer/helper_test.h"
+#include "src/tint/lang/msl/writer/helper_test.h"
 #include "src/tint/utils/text/string_stream.h"
 
 using namespace tint::core::number_suffixes;  // NOLINT
@@ -47,8 +47,8 @@
     return out;
 }
 
-using MslPrinterBinaryTest = MslPrinterTestWithParam<BinaryData>;
-TEST_P(MslPrinterBinaryTest, Emit) {
+using MslWriterBinaryTest = MslWriterTestWithParam<BinaryData>;
+TEST_P(MslWriterBinaryTest, Emit) {
     auto params = GetParam();
 
     auto* func = b.Function("foo", ty.void_());
@@ -60,18 +60,18 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   uint const left = 1u;
   uint const right = 2u;
   uint const val = )" + params.result +
-                           R"(;
+                               R"(;
 }
 )");
 }
-INSTANTIATE_TEST_SUITE_P(MslPrinterTest,
-                         MslPrinterBinaryTest,
+INSTANTIATE_TEST_SUITE_P(MslWriterTest,
+                         MslWriterBinaryTest,
                          testing::Values(BinaryData{"(left + right)", core::BinaryOp::kAdd},
                                          BinaryData{"(left - right)", core::BinaryOp::kSubtract},
                                          BinaryData{"(left * right)", core::BinaryOp::kMultiply},
@@ -79,7 +79,7 @@
                                          BinaryData{"(left | right)", core::BinaryOp::kOr},
                                          BinaryData{"(left ^ right)", core::BinaryOp::kXor}));
 
-TEST_F(MslPrinterTest, BinaryDivU32) {
+TEST_F(MslWriterTest, BinaryDivU32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* l = b.Let("left", b.Constant(1_u));
@@ -89,8 +89,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 uint tint_div_u32(uint lhs, uint rhs) {
   return (lhs / select(rhs, 1u, (rhs == 0u)));
 }
@@ -102,7 +102,7 @@
 )");
 }
 
-TEST_F(MslPrinterTest, BinaryModU32) {
+TEST_F(MslWriterTest, BinaryModU32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* l = b.Let("left", b.Constant(1_u));
@@ -112,8 +112,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 uint tint_mod_u32(uint lhs, uint rhs) {
   uint const v = select(rhs, 1u, (rhs == 0u));
   return (lhs - ((lhs / v) * v));
@@ -126,7 +126,7 @@
 )");
 }
 
-TEST_F(MslPrinterTest, BinaryShiftLeft) {
+TEST_F(MslWriterTest, BinaryShiftLeft) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* l = b.Let("left", b.Constant(1_u));
@@ -136,8 +136,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   uint const left = 1u;
   uint const right = 2u;
@@ -146,7 +146,7 @@
 )");
 }
 
-TEST_F(MslPrinterTest, BinaryShiftRight) {
+TEST_F(MslWriterTest, BinaryShiftRight) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* l = b.Let("left", b.Constant(1_u));
@@ -156,8 +156,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   uint const left = 1u;
   uint const right = 2u;
@@ -166,8 +166,8 @@
 )");
 }
 
-using MslPrinterBinaryBoolTest = MslPrinterTestWithParam<BinaryData>;
-TEST_P(MslPrinterBinaryBoolTest, Emit) {
+using MslWriterBinaryBoolTest = MslWriterTestWithParam<BinaryData>;
+TEST_P(MslWriterBinaryBoolTest, Emit) {
     auto params = GetParam();
 
     auto* func = b.Function("foo", ty.void_());
@@ -179,19 +179,19 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   uint const left = 1u;
   uint const right = 2u;
   bool const val = )" + params.result +
-                           R"(;
+                               R"(;
 }
 )");
 }
 INSTANTIATE_TEST_SUITE_P(
-    MslPrinterTest,
-    MslPrinterBinaryBoolTest,
+    MslWriterTest,
+    MslWriterBinaryBoolTest,
     testing::Values(BinaryData{"(left == right)", core::BinaryOp::kEqual},
                     BinaryData{"(left != right)", core::BinaryOp::kNotEqual},
                     BinaryData{"(left < right)", core::BinaryOp::kLessThan},
@@ -201,8 +201,8 @@
 
 // TODO(dsinclair): Needs transform
 // TODO(dsinclair): Requires `bitcast` support
-using MslPrinterBinaryTest_SignedOverflowDefinedBehaviour = MslPrinterTestWithParam<BinaryData>;
-TEST_P(MslPrinterBinaryTest_SignedOverflowDefinedBehaviour, DISABLED_Emit) {
+using MslWriterBinaryTest_SignedOverflowDefinedBehaviour = MslWriterTestWithParam<BinaryData>;
+TEST_P(MslWriterBinaryTest_SignedOverflowDefinedBehaviour, DISABLED_Emit) {
     auto params = GetParam();
 
     auto* func = b.Function("foo", ty.void_());
@@ -215,13 +215,13 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   int const left = 1i;
   int const right = 3i;
   int const val = )" + params.result +
-                           R"(;
+                               R"(;
       }
 )");
 }
@@ -230,15 +230,14 @@
     {"as_type<int>((as_type<uint>(left) + as_type<uint>(right)))", core::BinaryOp::kAdd},
     {"as_type<int>((as_type<uint>(left) - as_type<uint>(right)))", core::BinaryOp::kSubtract},
     {"as_type<int>((as_type<uint>(left) * as_type<uint>(right)))", core::BinaryOp::kMultiply}};
-INSTANTIATE_TEST_SUITE_P(MslPrinterTest,
-                         MslPrinterBinaryTest_SignedOverflowDefinedBehaviour,
+INSTANTIATE_TEST_SUITE_P(MslWriterTest,
+                         MslWriterBinaryTest_SignedOverflowDefinedBehaviour,
                          testing::ValuesIn(signed_overflow_defined_behaviour_cases));
 
 // TODO(dsinclair): Needs transform
 // TODO(dsinclair): Requires `bitcast` support
-using MslPrinterBinaryTest_ShiftSignedOverflowDefinedBehaviour =
-    MslPrinterTestWithParam<BinaryData>;
-TEST_P(MslPrinterBinaryTest_ShiftSignedOverflowDefinedBehaviour, DISABLED_Emit) {
+using MslWriterBinaryTest_ShiftSignedOverflowDefinedBehaviour = MslWriterTestWithParam<BinaryData>;
+TEST_P(MslWriterBinaryTest_ShiftSignedOverflowDefinedBehaviour, DISABLED_Emit) {
     auto params = GetParam();
 
     auto* func = b.Function("foo", ty.void_());
@@ -250,13 +249,13 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   int const left = 1i;
   uint const right = 2u;
   int const val = )" + params.result +
-                           R"(;
+                               R"(;
       }
 )");
 }
@@ -264,15 +263,15 @@
 constexpr BinaryData shift_signed_overflow_defined_behaviour_cases[] = {
     {"as_type<int>((as_type<uint>(left) << right))", core::BinaryOp::kShiftLeft},
     {"(left >> right)", core::BinaryOp::kShiftRight}};
-INSTANTIATE_TEST_SUITE_P(MslPrinterTest,
-                         MslPrinterBinaryTest_ShiftSignedOverflowDefinedBehaviour,
+INSTANTIATE_TEST_SUITE_P(MslWriterTest,
+                         MslWriterBinaryTest_ShiftSignedOverflowDefinedBehaviour,
                          testing::ValuesIn(shift_signed_overflow_defined_behaviour_cases));
 
 // TODO(dsinclair): Needs transform
 // TODO(dsinclair): Requires `bitcast`
-using MslPrinterBinaryTest_SignedOverflowDefinedBehaviour_Chained =
-    MslPrinterTestWithParam<BinaryData>;
-TEST_P(MslPrinterBinaryTest_SignedOverflowDefinedBehaviour_Chained, DISABLED_Emit) {
+using MslWriterBinaryTest_SignedOverflowDefinedBehaviour_Chained =
+    MslWriterTestWithParam<BinaryData>;
+TEST_P(MslWriterBinaryTest_SignedOverflowDefinedBehaviour_Chained, DISABLED_Emit) {
     auto params = GetParam();
 
     auto* func = b.Function("foo", ty.void_());
@@ -286,13 +285,13 @@
         b.Let("val", expr2);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   int left;
   int right;
   int const val = )" + params.result +
-                           R"(;
+                               R"(;
 )");
 }
 constexpr BinaryData signed_overflow_defined_behaviour_chained_cases[] = {
@@ -305,15 +304,15 @@
     {R"(as_type<int>((as_type<uint>(as_type<int>((as_type<uint>(left) * as_type<uint>(right)))) *
     as_type<uint>(right))))",
      core::BinaryOp::kMultiply}};
-INSTANTIATE_TEST_SUITE_P(MslPrinterTest,
-                         MslPrinterBinaryTest_SignedOverflowDefinedBehaviour_Chained,
+INSTANTIATE_TEST_SUITE_P(MslWriterTest,
+                         MslWriterBinaryTest_SignedOverflowDefinedBehaviour_Chained,
                          testing::ValuesIn(signed_overflow_defined_behaviour_chained_cases));
 
 // TODO(dsinclair): Needs transform
 // TODO(dsinclair): Requires `bitcast`
-using MslPrinterBinaryTest_ShiftSignedOverflowDefinedBehaviour_Chained =
-    MslPrinterTestWithParam<BinaryData>;
-TEST_P(MslPrinterBinaryTest_ShiftSignedOverflowDefinedBehaviour_Chained, DISABLED_Emit) {
+using MslWriterBinaryTest_ShiftSignedOverflowDefinedBehaviour_Chained =
+    MslWriterTestWithParam<BinaryData>;
+TEST_P(MslWriterBinaryTest_ShiftSignedOverflowDefinedBehaviour_Chained, DISABLED_Emit) {
     auto params = GetParam();
 
     auto* func = b.Function("foo", ty.void_());
@@ -327,13 +326,13 @@
         b.Let("val", expr2);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   int left;
   uint right;
   int const val = )" + params.result +
-                           R"(;
+                               R"(;
 )");
 }
 constexpr BinaryData shift_signed_overflow_defined_behaviour_chained_cases[] = {
@@ -341,12 +340,12 @@
      core::BinaryOp::kShiftLeft},
     {R"(((left >> right) >> right))", core::BinaryOp::kShiftRight},
 };
-INSTANTIATE_TEST_SUITE_P(MslPrinterTest,
-                         MslPrinterBinaryTest_ShiftSignedOverflowDefinedBehaviour_Chained,
+INSTANTIATE_TEST_SUITE_P(MslWriterTest,
+                         MslWriterBinaryTest_ShiftSignedOverflowDefinedBehaviour_Chained,
                          testing::ValuesIn(shift_signed_overflow_defined_behaviour_chained_cases));
 
 // TODO(dsinclair): Needs transform
-TEST_F(MslPrinterTest, DISABLED_BinaryModF32) {
+TEST_F(MslWriterTest, DISABLED_BinaryModF32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* left = b.Var("left", ty.ptr<core::AddressSpace::kFunction, f32>());
@@ -357,8 +356,8 @@
         b.Let("val", expr1);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float left;
   float right;
@@ -367,7 +366,7 @@
 }
 
 // TODO(dsinclair): Needs transform
-TEST_F(MslPrinterTest, DISABLED_BinaryModF16) {
+TEST_F(MslWriterTest, DISABLED_BinaryModF16) {
     // Enable f16?
 
     auto* func = b.Function("foo", ty.void_());
@@ -380,8 +379,8 @@
         b.Let("val", expr1);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   half left;
   half right;
@@ -390,7 +389,7 @@
 }
 
 // TODO(dsinclair): Needs transform
-TEST_F(MslPrinterTest, DISABLED_BinaryModVec3F32) {
+TEST_F(MslWriterTest, DISABLED_BinaryModVec3F32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* left = b.Var("left", ty.ptr(core::AddressSpace::kFunction, ty.vec3<f32>()));
@@ -401,8 +400,8 @@
         b.Let("val", expr1);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float3 left;
   float3 right;
@@ -411,7 +410,7 @@
 }
 
 // TODO(dsinclair): Needs transform
-TEST_F(MslPrinterTest, DISABLED_BinaryModVec3F16) {
+TEST_F(MslWriterTest, DISABLED_BinaryModVec3F16) {
     // Enable f16?
 
     auto* func = b.Function("foo", ty.void_());
@@ -424,8 +423,8 @@
         b.Let("val", expr1);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   half3 left;
   half3 right;
@@ -434,7 +433,7 @@
 }
 
 // TODO(dsinclair): Needs transform
-TEST_F(MslPrinterTest, DISABLED_BinaryBoolAnd) {
+TEST_F(MslWriterTest, DISABLED_BinaryBoolAnd) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* left = b.Var("left", ty.ptr(core::AddressSpace::kFunction, ty.bool_()));
@@ -445,8 +444,8 @@
         b.Let("val", expr1);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float left;
   float right;
@@ -455,7 +454,7 @@
 }
 
 // TODO(dsinclair): Needs transform
-TEST_F(MslPrinterTest, DISABLED_BinaryBoolOr) {
+TEST_F(MslWriterTest, DISABLED_BinaryBoolOr) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* left = b.Var("left", ty.ptr(core::AddressSpace::kFunction, ty.bool_()));
@@ -466,8 +465,8 @@
         b.Let("val", expr1);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float left;
   float right;
diff --git a/src/tint/lang/msl/writer/common/printer_support_test.cc b/src/tint/lang/msl/writer/common/printer_support_test.cc
index 49e2063..dd781ad 100644
--- a/src/tint/lang/msl/writer/common/printer_support_test.cc
+++ b/src/tint/lang/msl/writer/common/printer_support_test.cc
@@ -49,7 +49,7 @@
 }
 
 INSTANTIATE_TEST_SUITE_P(
-    MslPrinterTest,
+    MslWriterTest,
     MslBuiltinConversionTest,
     testing::Values(
         MslBuiltinData{core::BuiltinValue::kPosition, "position"},
diff --git a/src/tint/lang/msl/writer/printer/constant_test.cc b/src/tint/lang/msl/writer/constant_test.cc
similarity index 72%
rename from src/tint/lang/msl/writer/printer/constant_test.cc
rename to src/tint/lang/msl/writer/constant_test.cc
index f4dd3b2..320a418 100644
--- a/src/tint/lang/msl/writer/printer/constant_test.cc
+++ b/src/tint/lang/msl/writer/constant_test.cc
@@ -28,7 +28,7 @@
 #include "src/tint/lang/core/fluent_types.h"
 #include "src/tint/lang/core/type/array.h"
 #include "src/tint/lang/core/type/matrix.h"
-#include "src/tint/lang/msl/writer/printer/helper_test.h"
+#include "src/tint/lang/msl/writer/helper_test.h"
 #include "src/tint/utils/text/string.h"
 
 using namespace tint::core::number_suffixes;  // NOLINT
@@ -37,7 +37,7 @@
 namespace tint::msl::writer {
 namespace {
 
-TEST_F(MslPrinterTest, Constant_Bool_True) {
+TEST_F(MslWriterTest, Constant_Bool_True) {
     auto* c = b.Constant(true);
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -45,15 +45,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   bool const a = true;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_Bool_False) {
+TEST_F(MslWriterTest, Constant_Bool_False) {
     auto* c = b.Constant(false);
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -61,15 +61,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   bool const a = false;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_i32) {
+TEST_F(MslWriterTest, Constant_i32) {
     auto* c = b.Constant(-12345_i);
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -77,15 +77,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   int const a = -12345;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_u32) {
+TEST_F(MslWriterTest, Constant_u32) {
     auto* c = b.Constant(12345_u);
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -93,15 +93,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   uint const a = 12345u;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_F32) {
+TEST_F(MslWriterTest, Constant_F32) {
     auto* c = b.Constant(f32((1 << 30) - 4));
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -109,15 +109,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float const a = 1073741824.0f;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_F16) {
+TEST_F(MslWriterTest, Constant_F16) {
     auto* c = b.Constant(f16((1 << 15) - 8));
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -125,15 +125,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   half const a = 32752.0h;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_Vector_Splat) {
+TEST_F(MslWriterTest, Constant_Vector_Splat) {
     auto* c = b.Splat<vec3<f32>>(1.5_f);
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -141,15 +141,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float3 const a = float3(1.5f);
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_Vector_Composite) {
+TEST_F(MslWriterTest, Constant_Vector_Composite) {
     auto* c = b.Composite<vec3<f32>>(1.5_f, 1.0_f, 1.5_f);
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -157,15 +157,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float3 const a = float3(1.5f, 1.0f, 1.5f);
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_Vector_Composite_AnyZero) {
+TEST_F(MslWriterTest, Constant_Vector_Composite_AnyZero) {
     auto* c = b.Composite<vec3<f32>>(1.0_f, 0.0_f, 1.5_f);
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -173,15 +173,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float3 const a = float3(1.0f, 0.0f, 1.5f);
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_Vector_Composite_AllZero) {
+TEST_F(MslWriterTest, Constant_Vector_Composite_AllZero) {
     auto* c = b.Composite<vec3<f32>>(0.0_f, 0.0_f, 0.0_f);
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -189,15 +189,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float3 const a = float3(0.0f);
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_Matrix_Splat) {
+TEST_F(MslWriterTest, Constant_Matrix_Splat) {
     auto* c = b.Splat<mat3x2<f32>>(b.Splat<vec2<f32>>(1.5_f));
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -205,15 +205,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float3x2 const a = float3x2(float2(1.5f), float2(1.5f), float2(1.5f));
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_Matrix_Composite) {
+TEST_F(MslWriterTest, Constant_Matrix_Composite) {
     auto* c = b.Composite<mat3x2<f32>>(        //
         b.Composite<vec2<f32>>(1.5_f, 1.0_f),  //
         b.Composite<vec2<f32>>(1.5_f, 2.0_f),  //
@@ -224,15 +224,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float3x2 const a = float3x2(float2(1.5f, 1.0f), float2(1.5f, 2.0f), float2(2.5f, 3.5f));
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_Matrix_Composite_AnyZero) {
+TEST_F(MslWriterTest, Constant_Matrix_Composite_AnyZero) {
     auto* c = b.Composite<mat2x2<f32>>(        //
         b.Composite<vec2<f32>>(1.0_f, 0.0_f),  //
         b.Composite<vec2<f32>>(1.5_f, 2.5_f));
@@ -242,15 +242,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float2x2 const a = float2x2(float2(1.0f, 0.0f), float2(1.5f, 2.5f));
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_Matrix_Composite_AllZero) {
+TEST_F(MslWriterTest, Constant_Matrix_Composite_AllZero) {
     auto* c = b.Composite<mat3x2<f32>>(        //
         b.Composite<vec2<f32>>(0.0_f, 0.0_f),  //
         b.Composite<vec2<f32>>(0.0_f, 0.0_f),  //
@@ -261,15 +261,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float3x2 const a = float3x2(float2(0.0f), float2(0.0f), float2(0.0f));
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_Array_Splat) {
+TEST_F(MslWriterTest, Constant_Array_Splat) {
     auto* c = b.Splat<array<f32, 3>>(1.5_f);
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -277,15 +277,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + MetalArray() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + MetalArray() + R"(
 void foo() {
   tint_array<float, 3> const a = tint_array<float, 3>{1.5f, 1.5f, 1.5f};
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_Array_Composite) {
+TEST_F(MslWriterTest, Constant_Array_Composite) {
     auto* c = b.Composite<array<f32, 3>>(1.5_f, 1.0_f, 2.0_f);
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -293,15 +293,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + MetalArray() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + MetalArray() + R"(
 void foo() {
   tint_array<float, 3> const a = tint_array<float, 3>{1.5f, 1.0f, 2.0f};
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_Array_Composite_AnyZero) {
+TEST_F(MslWriterTest, Constant_Array_Composite_AnyZero) {
     auto* c = b.Composite<array<f32, 2>>(1.0_f, 0.0_f);
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -309,15 +309,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + MetalArray() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + MetalArray() + R"(
 void foo() {
   tint_array<float, 2> const a = tint_array<float, 2>{1.0f, 0.0f};
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_Array_Composite_AllZero) {
+TEST_F(MslWriterTest, Constant_Array_Composite_AllZero) {
     auto* c = b.Composite<array<f32, 3>>(0.0_f, 0.0_f, 0.0_f);
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -325,15 +325,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + MetalArray() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + MetalArray() + R"(
 void foo() {
   tint_array<float, 3> const a = tint_array<float, 3>{};
 }
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_Struct_Splat) {
+TEST_F(MslWriterTest, Constant_Struct_Splat) {
     auto* s = ty.Struct(mod.symbols.New("S"), {
                                                   {mod.symbols.Register("a"), ty.f32()},
                                                   {mod.symbols.Register("b"), ty.f32()},
@@ -345,8 +345,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(struct S {
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(struct S {
   float a;
   float b;
 };
@@ -357,7 +357,7 @@
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_Struct_Composite) {
+TEST_F(MslWriterTest, Constant_Struct_Composite) {
     auto* s = ty.Struct(mod.symbols.New("S"), {
                                                   {mod.symbols.Register("a"), ty.f32()},
                                                   {mod.symbols.Register("b"), ty.f32()},
@@ -369,8 +369,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(struct S {
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(struct S {
   float a;
   float b;
 };
@@ -381,7 +381,7 @@
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_Struct_Composite_AnyZero) {
+TEST_F(MslWriterTest, Constant_Struct_Composite_AnyZero) {
     auto* s = ty.Struct(mod.symbols.New("S"), {
                                                   {mod.symbols.Register("a"), ty.f32()},
                                                   {mod.symbols.Register("b"), ty.f32()},
@@ -393,8 +393,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(struct S {
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(struct S {
   float a;
   float b;
 };
@@ -405,7 +405,7 @@
 )");
 }
 
-TEST_F(MslPrinterTest, Constant_Struct_Composite_AllZero) {
+TEST_F(MslWriterTest, Constant_Struct_Composite_AllZero) {
     auto* s = ty.Struct(mod.symbols.New("S"), {
                                                   {mod.symbols.Register("a"), ty.f32()},
                                                   {mod.symbols.Register("b"), ty.f32()},
@@ -417,8 +417,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(struct S {
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(struct S {
   float a;
   float b;
 };
diff --git a/src/tint/lang/msl/writer/printer/discard_test.cc b/src/tint/lang/msl/writer/discard_test.cc
similarity index 92%
rename from src/tint/lang/msl/writer/printer/discard_test.cc
rename to src/tint/lang/msl/writer/discard_test.cc
index 79a7131..c454cf8 100644
--- a/src/tint/lang/msl/writer/printer/discard_test.cc
+++ b/src/tint/lang/msl/writer/discard_test.cc
@@ -25,14 +25,14 @@
 // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
-#include "src/tint/lang/msl/writer/printer/helper_test.h"
+#include "src/tint/lang/msl/writer/helper_test.h"
 
 using namespace tint::core::number_suffixes;  // NOLINT
 
 namespace tint::msl::writer {
 namespace {
 
-TEST_F(MslPrinterTest, Discard) {
+TEST_F(MslWriterTest, Discard) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* if_ = b.If(true);
@@ -49,8 +49,8 @@
         b.Return(ep);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(struct tint_module_vars_struct {
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(struct tint_module_vars_struct {
   thread bool* continue_execution;
 };
 
diff --git a/src/tint/lang/msl/writer/printer/function_test.cc b/src/tint/lang/msl/writer/function_test.cc
similarity index 85%
rename from src/tint/lang/msl/writer/printer/function_test.cc
rename to src/tint/lang/msl/writer/function_test.cc
index 8ba60c2..21afd3a 100644
--- a/src/tint/lang/msl/writer/printer/function_test.cc
+++ b/src/tint/lang/msl/writer/function_test.cc
@@ -26,23 +26,23 @@
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
 #include "src/tint/lang/core/type/sampled_texture.h"
-#include "src/tint/lang/msl/writer/printer/helper_test.h"
+#include "src/tint/lang/msl/writer/helper_test.h"
 
 namespace tint::msl::writer {
 namespace {
 
-TEST_F(MslPrinterTest, Function_Empty) {
+TEST_F(MslWriterTest, Function_Empty) {
     auto* func = b.Function("foo", ty.void_());
     func->Block()->Append(b.Return(func));
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EntryPointParameterBufferBindingPoint) {
+TEST_F(MslWriterTest, EntryPointParameterBufferBindingPoint) {
     auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
     auto* storage = b.FunctionParam("storage", ty.ptr(core::AddressSpace::kStorage, ty.i32()));
     auto* uniform = b.FunctionParam("uniform", ty.ptr(core::AddressSpace::kUniform, ty.i32()));
@@ -51,14 +51,14 @@
     func->SetParams({storage, uniform});
     func->Block()->Append(b.Return(func));
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 fragment void foo(device int* storage [[buffer(1)]], const constant int* uniform [[buffer(2)]]) {
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EntryPointParameterHandleBindingPoint) {
+TEST_F(MslWriterTest, EntryPointParameterHandleBindingPoint) {
     auto* t = ty.Get<core::type::SampledTexture>(core::type::TextureDimension::k2d, ty.f32());
     auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kFragment);
     auto* texture = b.FunctionParam("texture", t);
@@ -68,8 +68,8 @@
     func->SetParams({texture, sampler});
     func->Block()->Append(b.Return(func));
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 fragment void foo(texture2d<float, access::sample> texture [[texture(1)]], sampler sampler [[sampler(2)]]) {
 }
 )");
diff --git a/src/tint/lang/msl/writer/printer/helper_test.h b/src/tint/lang/msl/writer/helper_test.h
similarity index 80%
rename from src/tint/lang/msl/writer/printer/helper_test.h
rename to src/tint/lang/msl/writer/helper_test.h
index bef8b61..c563d69 100644
--- a/src/tint/lang/msl/writer/printer/helper_test.h
+++ b/src/tint/lang/msl/writer/helper_test.h
@@ -25,8 +25,8 @@
 // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
-#ifndef SRC_TINT_LANG_MSL_WRITER_PRINTER_HELPER_TEST_H_
-#define SRC_TINT_LANG_MSL_WRITER_PRINTER_HELPER_TEST_H_
+#ifndef SRC_TINT_LANG_MSL_WRITER_HELPER_TEST_H_
+#define SRC_TINT_LANG_MSL_WRITER_HELPER_TEST_H_
 
 #include <iostream>
 #include <string>
@@ -35,8 +35,7 @@
 #include "src/tint/lang/core/ir/builder.h"
 #include "src/tint/lang/core/ir/validator.h"
 #include "src/tint/lang/msl/validate/validate.h"
-#include "src/tint/lang/msl/writer/printer/printer.h"
-#include "src/tint/lang/msl/writer/raise/raise.h"
+#include "src/tint/lang/msl/writer/writer.h"
 
 namespace tint::msl::writer {
 
@@ -60,9 +59,9 @@
 
 )";
 
-/// Base helper class for testing the MSL generator implementation.
+/// Base helper class for testing the MSL writer implementation.
 template <typename BASE>
-class MslPrinterTestHelperBase : public BASE {
+class MslWriterTestHelperBase : public BASE {
   public:
     /// The test module.
     core::ir::Module mod;
@@ -76,17 +75,13 @@
     std::string err_;
 
     /// Generated MSL
-    std::string output_;
+    Output output_;
 
     /// Run the writer on the IR module and validate the result.
     /// @returns true if generation and validation succeeded
     bool Generate() {
-        if (auto raised = Raise(mod, {}); raised != Success) {
-            err_ = raised.Failure().reason.Str();
-            return false;
-        }
-
-        auto result = Print(mod);
+        Options options;
+        auto result = writer::Generate(mod, options);
         if (result != Success) {
             err_ = result.Failure().reason.Str();
             return false;
@@ -94,7 +89,8 @@
         output_ = result.Get();
 
 #if TINT_BUILD_IS_MAC
-        auto msl_validation = validate::ValidateUsingMetal(output_, validate::MslVersion::kMsl_2_3);
+        auto msl_validation =
+            validate::ValidateUsingMetal(output_.msl, validate::MslVersion::kMsl_2_3);
         if (msl_validation.failed) {
             err_ = msl_validation.output;
             return false;
@@ -111,12 +107,12 @@
 };
 
 /// Printer tests
-using MslPrinterTest = MslPrinterTestHelperBase<testing::Test>;
+using MslWriterTest = MslWriterTestHelperBase<testing::Test>;
 
 /// Printer param tests
 template <typename T>
-using MslPrinterTestWithParam = MslPrinterTestHelperBase<testing::TestWithParam<T>>;
+using MslWriterTestWithParam = MslWriterTestHelperBase<testing::TestWithParam<T>>;
 
 }  // namespace tint::msl::writer
 
-#endif  // SRC_TINT_LANG_MSL_WRITER_PRINTER_HELPER_TEST_H_
+#endif  // SRC_TINT_LANG_MSL_WRITER_HELPER_TEST_H_
diff --git a/src/tint/lang/msl/writer/printer/if_test.cc b/src/tint/lang/msl/writer/if_test.cc
similarity index 82%
rename from src/tint/lang/msl/writer/printer/if_test.cc
rename to src/tint/lang/msl/writer/if_test.cc
index 2f1a583..ef798c7 100644
--- a/src/tint/lang/msl/writer/printer/if_test.cc
+++ b/src/tint/lang/msl/writer/if_test.cc
@@ -25,14 +25,14 @@
 // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
-#include "src/tint/lang/msl/writer/printer/helper_test.h"
+#include "src/tint/lang/msl/writer/helper_test.h"
 
 using namespace tint::core::number_suffixes;  // NOLINT
 
 namespace tint::msl::writer {
 namespace {
 
-TEST_F(MslPrinterTest, If) {
+TEST_F(MslWriterTest, If) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* if_ = b.If(true);
@@ -40,8 +40,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   if (true) {
   }
@@ -49,7 +49,7 @@
 )");
 }
 
-TEST_F(MslPrinterTest, IfWithElseIf) {
+TEST_F(MslWriterTest, IfWithElseIf) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* if_ = b.If(true);
@@ -62,8 +62,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   if (true) {
   } else {
@@ -74,7 +74,7 @@
 )");
 }
 
-TEST_F(MslPrinterTest, IfWithElse) {
+TEST_F(MslWriterTest, IfWithElse) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* if_ = b.If(true);
@@ -83,8 +83,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   if (true) {
   } else {
@@ -94,7 +94,7 @@
 )");
 }
 
-TEST_F(MslPrinterTest, IfBothBranchesReturn) {
+TEST_F(MslWriterTest, IfBothBranchesReturn) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* if_ = b.If(true);
@@ -103,8 +103,8 @@
         b.Unreachable();
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   if (true) {
     return;
@@ -117,7 +117,7 @@
 }
 
 // Requires a transform to turn PHIs into lets
-TEST_F(MslPrinterTest, DISABLED_IfWithSinglePhi) {
+TEST_F(MslWriterTest, DISABLED_IfWithSinglePhi) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* i = b.If(true);
@@ -131,8 +131,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   int tint_symbol;
   if (true) {
@@ -145,7 +145,7 @@
 }
 
 // Requires a transform to turn PHIs into lets
-TEST_F(MslPrinterTest, DISABLED_IfWithMultiPhi) {
+TEST_F(MslWriterTest, DISABLED_IfWithMultiPhi) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* i = b.If(true);
@@ -159,8 +159,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   int tint_symbol;
   bool tint_symbol_1;
@@ -176,7 +176,7 @@
 }
 
 // Requires a transform to turn PHIs into lets
-TEST_F(MslPrinterTest, DISABLED_IfWithMultiPhiReturn1) {
+TEST_F(MslWriterTest, DISABLED_IfWithMultiPhiReturn1) {
     auto* func = b.Function("foo", ty.i32());
     b.Append(func->Block(), [&] {
         auto* i = b.If(true);
@@ -190,8 +190,8 @@
         b.Return(func, i->Result(0));
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 int foo() {
   int tint_symbol;
   bool tint_symbol_1;
@@ -208,7 +208,7 @@
 }
 
 // Requires a transform to turn PHIs into lets
-TEST_F(MslPrinterTest, DISABLED_IfWithMultiPhiReturn2) {
+TEST_F(MslWriterTest, DISABLED_IfWithMultiPhiReturn2) {
     auto* func = b.Function("foo", ty.bool_());
     b.Append(func->Block(), [&] {
         auto* i = b.If(true);
@@ -222,8 +222,8 @@
         b.Return(func, i->Result(1));
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 bool foo() {
   int tint_symbol;
   bool tint_symbol_1;
diff --git a/src/tint/lang/msl/writer/printer/let_test.cc b/src/tint/lang/msl/writer/let_test.cc
similarity index 75%
rename from src/tint/lang/msl/writer/printer/let_test.cc
rename to src/tint/lang/msl/writer/let_test.cc
index 85248bb..7a3653c 100644
--- a/src/tint/lang/msl/writer/printer/let_test.cc
+++ b/src/tint/lang/msl/writer/let_test.cc
@@ -25,7 +25,7 @@
 // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
-#include "src/tint/lang/msl/writer/printer/helper_test.h"
+#include "src/tint/lang/msl/writer/helper_test.h"
 
 namespace tint::msl::writer {
 namespace {
@@ -33,22 +33,22 @@
 using namespace tint::core::fluent_types;     // NOLINT
 using namespace tint::core::number_suffixes;  // NOLINT
 
-TEST_F(MslPrinterTest, LetU32) {
+TEST_F(MslWriterTest, LetU32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Let("l", 42_u);
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   uint const l = 42u;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, LetDuplicate) {
+TEST_F(MslWriterTest, LetDuplicate) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Let("l1", 42_u);
@@ -56,8 +56,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   uint const l1 = 42u;
   uint const l2 = 42u;
@@ -65,37 +65,37 @@
 )");
 }
 
-TEST_F(MslPrinterTest, LetF32) {
+TEST_F(MslWriterTest, LetF32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Let("l", 42.0_f);
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float const l = 42.0f;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, LetI32) {
+TEST_F(MslWriterTest, LetI32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Let("l", 42_i);
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   int const l = 42;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, LetF16) {
+TEST_F(MslWriterTest, LetF16) {
     // Enable F16?
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -103,30 +103,30 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   half const l = 42.0h;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, LetVec3F32) {
+TEST_F(MslWriterTest, LetVec3F32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Let("l", b.Composite<vec3<f32>>(1_f, 2_f, 3_f));
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float3 const l = float3(1.0f, 2.0f, 3.0f);
 }
 )");
 }
 
-TEST_F(MslPrinterTest, LetVec3F16) {
+TEST_F(MslWriterTest, LetVec3F16) {
     // Enable f16?
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -134,15 +134,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   half3 const l = half3(1.0h, 2.0h, 3.0h);
 }
 )");
 }
 
-TEST_F(MslPrinterTest, LetMat2x3F32) {
+TEST_F(MslWriterTest, LetMat2x3F32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Let("l", b.Composite<mat2x3<f32>>(b.Composite<vec3<f32>>(1_f, 2_f, 3_f),
@@ -150,15 +150,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float2x3 const l = float2x3(float3(1.0f, 2.0f, 3.0f), float3(4.0f, 5.0f, 6.0f));
 }
 )");
 }
 
-TEST_F(MslPrinterTest, LetMat2x3F16) {
+TEST_F(MslWriterTest, LetMat2x3F16) {
     // Enable f16?
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -167,30 +167,30 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   half2x3 const l = half2x3(half3(1.0h, 2.0h, 3.0h), half3(4.0h, 5.0h, 6.0h));
 }
 )");
 }
 
-TEST_F(MslPrinterTest, LetArrF32) {
+TEST_F(MslWriterTest, LetArrF32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Let("l", b.Composite<array<f32, 3>>(1_f, 2_f, 3_f));
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + MetalArray() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + MetalArray() + R"(
 void foo() {
   tint_array<float, 3> const l = tint_array<float, 3>{1.0f, 2.0f, 3.0f};
 }
 )");
 }
 
-TEST_F(MslPrinterTest, LetArrVec2Bool) {
+TEST_F(MslWriterTest, LetArrVec2Bool) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Let("l", b.Composite<array<vec2<bool>, 3>>(b.Composite<vec2<bool>>(true, false),
@@ -199,8 +199,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + MetalArray() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + MetalArray() + R"(
 void foo() {
   tint_array<bool2, 3> const l = tint_array<bool2, 3>{bool2(true, false), bool2(false, true), bool2(true, false)};
 }
diff --git a/src/tint/lang/msl/writer/printer/BUILD.bazel b/src/tint/lang/msl/writer/printer/BUILD.bazel
index ad3ce48..5cb3a75 100644
--- a/src/tint/lang/msl/writer/printer/BUILD.bazel
+++ b/src/tint/lang/msl/writer/printer/BUILD.bazel
@@ -77,55 +77,6 @@
   copts = COPTS,
   visibility = ["//visibility:public"],
 )
-cc_library(
-  name = "test",
-  alwayslink = True,
-  srcs = [
-    "binary_test.cc",
-    "constant_test.cc",
-    "discard_test.cc",
-    "function_test.cc",
-    "helper_test.h",
-    "if_test.cc",
-    "let_test.cc",
-    "return_test.cc",
-    "type_test.cc",
-    "var_test.cc",
-  ],
-  deps = [
-    "//src/tint/api/common",
-    "//src/tint/api/options",
-    "//src/tint/lang/core",
-    "//src/tint/lang/core/constant",
-    "//src/tint/lang/core/intrinsic",
-    "//src/tint/lang/core/ir",
-    "//src/tint/lang/core/type",
-    "//src/tint/utils/containers",
-    "//src/tint/utils/diagnostic",
-    "//src/tint/utils/ice",
-    "//src/tint/utils/id",
-    "//src/tint/utils/macros",
-    "//src/tint/utils/math",
-    "//src/tint/utils/memory",
-    "//src/tint/utils/reflection",
-    "//src/tint/utils/result",
-    "//src/tint/utils/rtti",
-    "//src/tint/utils/symbol",
-    "//src/tint/utils/text",
-    "//src/tint/utils/traits",
-    "@gtest",
-  ] + select({
-    ":tint_build_msl_writer": [
-      "//src/tint/lang/msl/validate",
-      "//src/tint/lang/msl/writer/common",
-      "//src/tint/lang/msl/writer/printer",
-      "//src/tint/lang/msl/writer/raise",
-    ],
-    "//conditions:default": [],
-  }),
-  copts = COPTS,
-  visibility = ["//visibility:public"],
-)
 
 alias(
   name = "tint_build_msl_writer",
diff --git a/src/tint/lang/msl/writer/printer/BUILD.cmake b/src/tint/lang/msl/writer/printer/BUILD.cmake
index 8ae3d60..762826d 100644
--- a/src/tint/lang/msl/writer/printer/BUILD.cmake
+++ b/src/tint/lang/msl/writer/printer/BUILD.cmake
@@ -77,60 +77,4 @@
   )
 endif(TINT_BUILD_MSL_WRITER)
 
-endif(TINT_BUILD_MSL_WRITER)
-if(TINT_BUILD_MSL_WRITER)
-################################################################################
-# Target:    tint_lang_msl_writer_printer_test
-# Kind:      test
-# Condition: TINT_BUILD_MSL_WRITER
-################################################################################
-tint_add_target(tint_lang_msl_writer_printer_test test
-  lang/msl/writer/printer/binary_test.cc
-  lang/msl/writer/printer/constant_test.cc
-  lang/msl/writer/printer/discard_test.cc
-  lang/msl/writer/printer/function_test.cc
-  lang/msl/writer/printer/helper_test.h
-  lang/msl/writer/printer/if_test.cc
-  lang/msl/writer/printer/let_test.cc
-  lang/msl/writer/printer/return_test.cc
-  lang/msl/writer/printer/type_test.cc
-  lang/msl/writer/printer/var_test.cc
-)
-
-tint_target_add_dependencies(tint_lang_msl_writer_printer_test test
-  tint_api_common
-  tint_api_options
-  tint_lang_core
-  tint_lang_core_constant
-  tint_lang_core_intrinsic
-  tint_lang_core_ir
-  tint_lang_core_type
-  tint_utils_containers
-  tint_utils_diagnostic
-  tint_utils_ice
-  tint_utils_id
-  tint_utils_macros
-  tint_utils_math
-  tint_utils_memory
-  tint_utils_reflection
-  tint_utils_result
-  tint_utils_rtti
-  tint_utils_symbol
-  tint_utils_text
-  tint_utils_traits
-)
-
-tint_target_add_external_dependencies(tint_lang_msl_writer_printer_test test
-  "gtest"
-)
-
-if(TINT_BUILD_MSL_WRITER)
-  tint_target_add_dependencies(tint_lang_msl_writer_printer_test test
-    tint_lang_msl_validate
-    tint_lang_msl_writer_common
-    tint_lang_msl_writer_printer
-    tint_lang_msl_writer_raise
-  )
-endif(TINT_BUILD_MSL_WRITER)
-
 endif(TINT_BUILD_MSL_WRITER)
\ No newline at end of file
diff --git a/src/tint/lang/msl/writer/printer/BUILD.gn b/src/tint/lang/msl/writer/printer/BUILD.gn
index ec57337..c31c3dc 100644
--- a/src/tint/lang/msl/writer/printer/BUILD.gn
+++ b/src/tint/lang/msl/writer/printer/BUILD.gn
@@ -37,10 +37,6 @@
 import("../../../../../../scripts/tint_overrides_with_defaults.gni")
 
 import("${tint_src_dir}/tint.gni")
-
-if (tint_build_unittests || tint_build_benchmarks) {
-  import("//testing/test.gni")
-}
 if (tint_build_msl_writer) {
   libtint_source_set("printer") {
     sources = [
@@ -78,53 +74,3 @@
     }
   }
 }
-if (tint_build_unittests) {
-  if (tint_build_msl_writer) {
-    tint_unittests_source_set("unittests") {
-      sources = [
-        "binary_test.cc",
-        "constant_test.cc",
-        "discard_test.cc",
-        "function_test.cc",
-        "helper_test.h",
-        "if_test.cc",
-        "let_test.cc",
-        "return_test.cc",
-        "type_test.cc",
-        "var_test.cc",
-      ]
-      deps = [
-        "${tint_src_dir}:gmock_and_gtest",
-        "${tint_src_dir}/api/common",
-        "${tint_src_dir}/api/options",
-        "${tint_src_dir}/lang/core",
-        "${tint_src_dir}/lang/core/constant",
-        "${tint_src_dir}/lang/core/intrinsic",
-        "${tint_src_dir}/lang/core/ir",
-        "${tint_src_dir}/lang/core/type",
-        "${tint_src_dir}/utils/containers",
-        "${tint_src_dir}/utils/diagnostic",
-        "${tint_src_dir}/utils/ice",
-        "${tint_src_dir}/utils/id",
-        "${tint_src_dir}/utils/macros",
-        "${tint_src_dir}/utils/math",
-        "${tint_src_dir}/utils/memory",
-        "${tint_src_dir}/utils/reflection",
-        "${tint_src_dir}/utils/result",
-        "${tint_src_dir}/utils/rtti",
-        "${tint_src_dir}/utils/symbol",
-        "${tint_src_dir}/utils/text",
-        "${tint_src_dir}/utils/traits",
-      ]
-
-      if (tint_build_msl_writer) {
-        deps += [
-          "${tint_src_dir}/lang/msl/validate",
-          "${tint_src_dir}/lang/msl/writer/common",
-          "${tint_src_dir}/lang/msl/writer/printer",
-          "${tint_src_dir}/lang/msl/writer/raise",
-        ]
-      }
-    }
-  }
-}
diff --git a/src/tint/lang/msl/writer/raise/module_scope_vars.cc b/src/tint/lang/msl/writer/raise/module_scope_vars.cc
index 00a569f..6a4ecb2 100644
--- a/src/tint/lang/msl/writer/raise/module_scope_vars.cc
+++ b/src/tint/lang/msl/writer/raise/module_scope_vars.cc
@@ -174,6 +174,9 @@
         core::ir::Function* func,
         const core::ir::ReferencedModuleVars::VarSet& referenced_vars) {
         core::ir::Value* module_var_struct = nullptr;
+        core::ir::FunctionParam* workgroup_allocation_param = nullptr;
+        Vector<core::type::Manager::StructMemberDesc, 4> workgroup_struct_members;
+
         // Add parameters and insert instruction at the top of the entry point to set up the
         // module-scope variables structure.
         b.InsertBefore(func->Block()->Front(), [&] {  //
@@ -206,6 +209,24 @@
                         decl = param;
                         break;
                     }
+                    case core::AddressSpace::kWorkgroup: {
+                        // Workgroup variables are received as a function parameter (to workaround
+                        // an MSL compiler bug with threadgroup matrices), and we aggregate all
+                        // workgroup variables into a structure to avoid hitting MSL's limit for
+                        // threadgroup memory arguments.
+                        if (!workgroup_allocation_param) {
+                            workgroup_allocation_param = b.FunctionParam(nullptr);
+                            func->AppendParam(workgroup_allocation_param);
+                        }
+                        decl = b.Access(ptr, workgroup_allocation_param,
+                                        u32(workgroup_struct_members.Length()))
+                                   ->Result(0);
+                        workgroup_struct_members.Push(core::type::Manager::StructMemberDesc{
+                            ir.symbols.New(),
+                            ptr->StoreType(),
+                        });
+                        break;
+                    }
                     case core::AddressSpace::kHandle: {
                         // Handle types become function parameters and drop the pointer.
                         auto* param = b.FunctionParam(ptr->UnwrapPtr());
@@ -230,6 +251,14 @@
             auto* construct = b.Construct(struct_type, std::move(construct_args));
             module_var_struct = b.Let(kModuleVarsName, construct)->Result(0);
         });
+
+        // Create the workgroup variable structure if needed.
+        if (!workgroup_struct_members.IsEmpty()) {
+            auto* workgroup_struct =
+                ty.Struct(ir.symbols.New(), std::move(workgroup_struct_members));
+            workgroup_allocation_param->SetType(ty.ptr<workgroup>(workgroup_struct));
+        }
+
         return module_var_struct;
     }
 
diff --git a/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc b/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
index 3bef228..cf3fe76 100644
--- a/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
+++ b/src/tint/lang/msl/writer/raise/module_scope_vars_test.cc
@@ -364,6 +364,73 @@
     EXPECT_EQ(expect, str());
 }
 
+TEST_F(MslWriter_ModuleScopeVarsTest, Workgroup) {
+    auto* var_a = b.Var("a", ty.ptr<workgroup, i32>());
+    auto* var_b = b.Var("b", ty.ptr<workgroup, i32>());
+    mod.root_block->Append(var_a);
+    mod.root_block->Append(var_b);
+
+    auto* func = b.Function("foo", ty.void_(), core::ir::Function::PipelineStage::kCompute,
+                            std::array<uint32_t, 3>{1u, 1u, 1u});
+    b.Append(func->Block(), [&] {
+        auto* load_a = b.Load(var_a);
+        auto* load_b = b.Load(var_b);
+        b.Store(var_a, b.Add<i32>(load_a, load_b));
+        b.Return(func);
+    });
+
+    auto* src = R"(
+$B1: {  # root
+  %a:ptr<workgroup, i32, read_write> = var
+  %b:ptr<workgroup, i32, read_write> = var
+}
+
+%foo = @compute @workgroup_size(1, 1, 1) func():void {
+  $B2: {
+    %4:i32 = load %a
+    %5:i32 = load %b
+    %6:i32 = add %4, %5
+    store %a, %6
+    ret
+  }
+}
+)";
+    EXPECT_EQ(src, str());
+
+    auto* expect = R"(
+tint_module_vars_struct = struct @align(1) {
+  a:ptr<workgroup, i32, read_write> @offset(0)
+  b:ptr<workgroup, i32, read_write> @offset(0)
+}
+
+tint_symbol_2 = struct @align(4) {
+  tint_symbol:i32 @offset(0)
+  tint_symbol_1:i32 @offset(4)
+}
+
+%foo = @compute @workgroup_size(1, 1, 1) func(%2:ptr<workgroup, tint_symbol_2, read_write>):void {
+  $B1: {
+    %a:ptr<workgroup, i32, read_write> = access %2, 0u
+    %b:ptr<workgroup, i32, read_write> = access %2, 1u
+    %5:tint_module_vars_struct = construct %a, %b
+    %tint_module_vars:tint_module_vars_struct = let %5
+    %7:ptr<workgroup, i32, read_write> = access %tint_module_vars, 0u
+    %8:i32 = load %7
+    %9:ptr<workgroup, i32, read_write> = access %tint_module_vars, 1u
+    %10:i32 = load %9
+    %11:i32 = add %8, %10
+    %12:ptr<workgroup, i32, read_write> = access %tint_module_vars, 0u
+    store %12, %11
+    ret
+  }
+}
+)";
+
+    Run(ModuleScopeVars);
+
+    EXPECT_EQ(expect, str());
+}
+
 TEST_F(MslWriter_ModuleScopeVarsTest, MultipleAddressSpaces) {
     auto* var_a = b.Var("a", ty.ptr<uniform, i32, core::Access::kRead>());
     auto* var_b = b.Var("b", ty.ptr<storage, i32, core::Access::kReadWrite>());
diff --git a/src/tint/lang/msl/writer/printer/return_test.cc b/src/tint/lang/msl/writer/return_test.cc
similarity index 82%
rename from src/tint/lang/msl/writer/printer/return_test.cc
rename to src/tint/lang/msl/writer/return_test.cc
index f004be1..07a12c6 100644
--- a/src/tint/lang/msl/writer/printer/return_test.cc
+++ b/src/tint/lang/msl/writer/return_test.cc
@@ -25,14 +25,14 @@
 // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
-#include "src/tint/lang/msl/writer/printer/helper_test.h"
+#include "src/tint/lang/msl/writer/helper_test.h"
 
 using namespace tint::core::number_suffixes;  // NOLINT
 
 namespace tint::msl::writer {
 namespace {
 
-TEST_F(MslPrinterTest, Return) {
+TEST_F(MslWriterTest, Return) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* if_ = b.If(true);
@@ -40,8 +40,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   if (true) {
     return;
@@ -50,23 +50,23 @@
 )");
 }
 
-TEST_F(MslPrinterTest, ReturnAtEndOfVoidDropped) {
+TEST_F(MslWriterTest, ReturnAtEndOfVoidDropped) {
     auto* func = b.Function("foo", ty.void_());
     func->Block()->Append(b.Return(func));
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
 }
 )");
 }
 
-TEST_F(MslPrinterTest, ReturnWithValue) {
+TEST_F(MslWriterTest, ReturnWithValue) {
     auto* func = b.Function("foo", ty.i32());
     func->Block()->Append(b.Return(func, 123_i));
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 int foo() {
   return 123;
 }
diff --git a/src/tint/lang/msl/writer/printer/type_test.cc b/src/tint/lang/msl/writer/type_test.cc
similarity index 86%
rename from src/tint/lang/msl/writer/printer/type_test.cc
rename to src/tint/lang/msl/writer/type_test.cc
index 0af7e7d..1fd6361 100644
--- a/src/tint/lang/msl/writer/printer/type_test.cc
+++ b/src/tint/lang/msl/writer/type_test.cc
@@ -35,7 +35,7 @@
 #include "src/tint/lang/core/type/sampled_texture.h"
 #include "src/tint/lang/core/type/storage_texture.h"
 #include "src/tint/lang/core/type/struct.h"
-#include "src/tint/lang/msl/writer/printer/helper_test.h"
+#include "src/tint/lang/msl/writer/helper_test.h"
 #include "src/tint/utils/text/string.h"
 
 namespace tint::msl::writer {
@@ -44,37 +44,37 @@
 using namespace tint::core::fluent_types;     // NOLINT
 using namespace tint::core::number_suffixes;  // NOLINT
 
-TEST_F(MslPrinterTest, EmitType_Array) {
+TEST_F(MslWriterTest, EmitType_Array) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.array<bool, 4>()));
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + MetalArray() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + MetalArray() + R"(
 void foo() {
   thread tint_array<bool, 4> a = {};
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EmitType_ArrayOfArray) {
+TEST_F(MslWriterTest, EmitType_ArrayOfArray) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.array(ty.array<bool, 4>(), 5)));
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + MetalArray() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + MetalArray() + R"(
 void foo() {
   thread tint_array<tint_array<bool, 4>, 5> a = {};
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EmitType_ArrayOfArrayOfArray) {
+TEST_F(MslWriterTest, EmitType_ArrayOfArrayOfArray) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a",
@@ -82,134 +82,134 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + MetalArray() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + MetalArray() + R"(
 void foo() {
   thread tint_array<tint_array<tint_array<bool, 4>, 5>, 6> a = {};
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EmitType_RuntimeArray) {
+TEST_F(MslWriterTest, EmitType_RuntimeArray) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.array<bool, 0>()));
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + MetalArray() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + MetalArray() + R"(
 void foo() {
   thread tint_array<bool, 1> a = {};
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EmitType_Bool) {
+TEST_F(MslWriterTest, EmitType_Bool) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.bool_()));
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   thread bool a = false;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EmitType_F32) {
+TEST_F(MslWriterTest, EmitType_F32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.f32()));
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   thread float a = 0.0f;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EmitType_F16) {
+TEST_F(MslWriterTest, EmitType_F16) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.f16()));
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   thread half a = 0.0h;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EmitType_I32) {
+TEST_F(MslWriterTest, EmitType_I32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.i32()));
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   thread int a = 0;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EmitType_Matrix_F32) {
+TEST_F(MslWriterTest, EmitType_Matrix_F32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.mat2x3<f32>()));
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   thread float2x3 a = float2x3(0.0f);
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EmitType_Matrix_F16) {
+TEST_F(MslWriterTest, EmitType_Matrix_F16) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.mat2x3<f16>()));
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   thread half2x3 a = half2x3(0.0h);
 }
 )");
 }
-TEST_F(MslPrinterTest, EmitType_U32) {
+TEST_F(MslWriterTest, EmitType_U32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.u32()));
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   thread uint a = 0u;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EmitType_Atomic_U32) {
+TEST_F(MslWriterTest, EmitType_Atomic_U32) {
     auto* func = b.Function("foo", ty.void_());
     auto* param = b.FunctionParam("a", ty.ptr(core::AddressSpace::kWorkgroup, ty.atomic<u32>()));
     func->SetParams({param});
@@ -217,14 +217,14 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo(threadgroup atomic_uint* const a) {
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EmitType_Atomic_I32) {
+TEST_F(MslWriterTest, EmitType_Atomic_I32) {
     auto* func = b.Function("foo", ty.void_());
     auto* param = b.FunctionParam("a", ty.ptr(core::AddressSpace::kWorkgroup, ty.atomic<i32>()));
     func->SetParams({param});
@@ -232,67 +232,67 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo(threadgroup atomic_int* const a) {
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EmitType_Vector) {
+TEST_F(MslWriterTest, EmitType_Vector) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr(core::AddressSpace::kPrivate, ty.vec3<f32>()));
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   thread float3 a = 0.0f;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EmitType_VectorPacked) {
+TEST_F(MslWriterTest, EmitType_VectorPacked) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr(core::AddressSpace::kFunction, ty.packed_vec(ty.f32(), 3)));
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   packed_float3 a = 0.0f;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EmitType_Void) {
+TEST_F(MslWriterTest, EmitType_Void) {
     // Tested via the function return type.
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {  //
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
 }
 )");
 }
 
 // TODO(dsinclair): How do we create a pointer type ... ?
-TEST_F(MslPrinterTest, DISABLED_EmitType_Pointer_Workgroup) {
+TEST_F(MslWriterTest, DISABLED_EmitType_Pointer_Workgroup) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr<workgroup, f32, read_write>());
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   threadgroup float* a;
 }
@@ -300,22 +300,22 @@
 }
 
 // TODO(dsinclair): How do we create a pointer type ... ?
-TEST_F(MslPrinterTest, DISABLED_EmitType_Pointer_Const) {
+TEST_F(MslWriterTest, DISABLED_EmitType_Pointer_Const) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr<function, f32>());
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   const thread float* a = 0.0f;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EmitType_Struct) {
+TEST_F(MslWriterTest, EmitType_Struct) {
     auto* s = ty.Struct(mod.symbols.New("S"), {
                                                   {mod.symbols.Register("a"), ty.i32()},
                                                   {mod.symbols.Register("b"), ty.f32()},
@@ -326,8 +326,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(struct S {
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(struct S {
   int a;
   float b;
 };
@@ -338,7 +338,7 @@
 )");
 }
 
-TEST_F(MslPrinterTest, EmitType_Struct_Dedup) {
+TEST_F(MslWriterTest, EmitType_Struct_Dedup) {
     auto* s = ty.Struct(mod.symbols.New("S"), {
                                                   {mod.symbols.Register("a"), ty.i32()},
                                                   {mod.symbols.Register("b"), ty.f32()},
@@ -350,8 +350,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(struct S {
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(struct S {
   int a;
   float b;
 };
@@ -448,7 +448,7 @@
                                       tint::RoundUp(align, size), size);
 }
 
-TEST_F(MslPrinterTest, EmitType_Struct_Layout_NonComposites) {
+TEST_F(MslWriterTest, EmitType_Struct_Layout_NonComposites) {
     Vector<MemberData, 26> data = {{mod.symbols.Register("a"), ty.i32(), 32},        //
                                    {mod.symbols.Register("b"), ty.f32(), 128, 128},  //
                                    {mod.symbols.Register("c"), ty.vec2<f32>()},      //
@@ -542,8 +542,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, expect.str());
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, expect.str());
 
     // 1.4 Metal and C++14
     // The Metal programming language is a C++14-based Specification with
@@ -568,7 +568,7 @@
 #undef ALL_FIELDS
 }
 
-TEST_F(MslPrinterTest, EmitType_Struct_Layout_Structures) {
+TEST_F(MslWriterTest, EmitType_Struct_Layout_Structures) {
     // inner_x: size(1024), align(512)
     Vector<MemberData, 2> inner_x_data = {{{mod.symbols.Register("a"), ty.i32()},  //
                                            {mod.symbols.Register("b"), ty.f32(), 0, 512}}};
@@ -628,8 +628,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, expect.str());
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, expect.str());
 
     // 1.4 Metal and C++14
     // The Metal programming language is a C++14-based Specification with
@@ -667,7 +667,7 @@
 #undef ALL_FIELDS
 }
 
-TEST_F(MslPrinterTest, EmitType_Struct_Layout_ArrayDefaultStride) {
+TEST_F(MslWriterTest, EmitType_Struct_Layout_ArrayDefaultStride) {
     // inner: size(1024), align(512)
     Vector<MemberData, 2> inner_data = {{mod.symbols.Register("a"), ty.i32()},
                                         {mod.symbols.Register("b"), ty.f32(), 0, 512}};
@@ -730,8 +730,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, expect.str());
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, expect.str());
 
     // 1.4 Metal and C++14
     // The Metal programming language is a C++14-based Specification with
@@ -775,7 +775,7 @@
 #undef ALL_FIELDS
 }
 
-TEST_F(MslPrinterTest, EmitType_Struct_Layout_ArrayVec3DefaultStride) {
+TEST_F(MslWriterTest, EmitType_Struct_Layout_ArrayVec3DefaultStride) {
     // array: size(64), align(16)
     auto array = ty.array<vec3<f32>, 4>();
 
@@ -816,11 +816,11 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, expect.str());
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, expect.str());
 }
 
-TEST_F(MslPrinterTest, AttemptTintPadSymbolCollision) {
+TEST_F(MslWriterTest, AttemptTintPadSymbolCollision) {
     Vector<MemberData, 26> data = {// uses symbols tint_pad_[0..9] and tint_pad_[20..35]
                                    {mod.symbols.Register("tint_pad_2"), ty.i32(), 32},         //
                                    {mod.symbols.Register("tint_pad_20"), ty.f32(), 128, 128},  //
@@ -905,11 +905,11 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, expect);
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, expect);
 }
 
-TEST_F(MslPrinterTest, EmitType_Sampler) {
+TEST_F(MslWriterTest, EmitType_Sampler) {
     auto* func = b.Function("foo", ty.void_());
     auto* param = b.FunctionParam("a", ty.sampler());
     func->SetParams({param});
@@ -917,14 +917,14 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo(sampler a) {
 }
 )");
 }
 
-TEST_F(MslPrinterTest, EmitType_SamplerComparison) {
+TEST_F(MslWriterTest, EmitType_SamplerComparison) {
     auto* func = b.Function("foo", ty.void_());
     auto* param = b.FunctionParam("a", ty.comparison_sampler());
     func->SetParams({param});
@@ -932,8 +932,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo(sampler a) {
 }
 )");
@@ -949,8 +949,8 @@
     out << str.str();
     return out;
 }
-using MslPrinterDepthTexturesTest = MslPrinterTestWithParam<MslDepthTextureData>;
-TEST_P(MslPrinterDepthTexturesTest, Emit) {
+using MslWriterDepthTexturesTest = MslWriterTestWithParam<MslDepthTextureData>;
+TEST_P(MslWriterDepthTexturesTest, Emit) {
     auto params = GetParam();
 
     auto* t = ty.Get<core::type::DepthTexture>(params.dim);
@@ -961,16 +961,16 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo()" + params.result +
-                           R"( a) {
+                               R"( a) {
 }
 )");
 }
 INSTANTIATE_TEST_SUITE_P(
-    MslPrinterTest,
-    MslPrinterDepthTexturesTest,
+    MslWriterTest,
+    MslWriterDepthTexturesTest,
     testing::Values(MslDepthTextureData{core::type::TextureDimension::k2d,
                                         "depth2d<float, access::sample>"},
                     MslDepthTextureData{core::type::TextureDimension::k2dArray,
@@ -980,7 +980,7 @@
                     MslDepthTextureData{core::type::TextureDimension::kCubeArray,
                                         "depthcube_array<float, access::sample>"}));
 
-TEST_F(MslPrinterTest, EmitType_DepthMultisampledTexture) {
+TEST_F(MslWriterTest, EmitType_DepthMultisampledTexture) {
     auto* t = ty.Get<core::type::DepthMultisampledTexture>(core::type::TextureDimension::k2d);
     auto* func = b.Function("foo", ty.void_());
     auto* param = b.FunctionParam("a", t);
@@ -989,8 +989,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo(depth2d_ms<float, access::read> a) {
 }
 )");
@@ -1006,8 +1006,8 @@
     out << str.str();
     return out;
 }
-using MslPrinterSampledtexturesTest = MslPrinterTestWithParam<MslTextureData>;
-TEST_P(MslPrinterSampledtexturesTest, Emit) {
+using MslWriterSampledtexturesTest = MslWriterTestWithParam<MslTextureData>;
+TEST_P(MslWriterSampledtexturesTest, Emit) {
     auto params = GetParam();
 
     auto* t = ty.Get<core::type::SampledTexture>(params.dim, ty.f32());
@@ -1018,16 +1018,16 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo()" + params.result +
-                           R"( a) {
+                               R"( a) {
 }
 )");
 }
 INSTANTIATE_TEST_SUITE_P(
-    MslPrinterTest,
-    MslPrinterSampledtexturesTest,
+    MslWriterTest,
+    MslWriterSampledtexturesTest,
     testing::Values(
         MslTextureData{core::type::TextureDimension::k1d, "texture1d<float, access::sample>"},
         MslTextureData{core::type::TextureDimension::k2d, "texture2d<float, access::sample>"},
@@ -1038,7 +1038,7 @@
         MslTextureData{core::type::TextureDimension::kCubeArray,
                        "texturecube_array<float, access::sample>"}));
 
-TEST_F(MslPrinterTest, EmitType_MultisampledTexture) {
+TEST_F(MslWriterTest, EmitType_MultisampledTexture) {
     auto* ms = ty.Get<core::type::MultisampledTexture>(core::type::TextureDimension::k2d, ty.u32());
     auto* func = b.Function("foo", ty.void_());
     auto* param = b.FunctionParam("a", ms);
@@ -1047,8 +1047,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo(texture2d_ms<uint, access::read> a) {
 }
 )");
@@ -1063,8 +1063,8 @@
     str << data.dim;
     return out << str.str();
 }
-using MslPrinterStorageTexturesTest = MslPrinterTestWithParam<MslStorageTextureData>;
-TEST_P(MslPrinterStorageTexturesTest, Emit) {
+using MslWriterStorageTexturesTest = MslWriterTestWithParam<MslStorageTextureData>;
+TEST_P(MslWriterStorageTexturesTest, Emit) {
     auto params = GetParam();
 
     auto* f32 = const_cast<core::type::F32*>(ty.f32());
@@ -1077,15 +1077,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo()" + params.result +
-                           R"( a) {
+                               R"( a) {
 }
 )");
 }
-INSTANTIATE_TEST_SUITE_P(MslPrinterTest,
-                         MslPrinterStorageTexturesTest,
+INSTANTIATE_TEST_SUITE_P(MslWriterTest,
+                         MslWriterStorageTexturesTest,
                          testing::Values(MslStorageTextureData{core::type::TextureDimension::k1d,
                                                                "texture1d<float, access::write>"},
                                          MslStorageTextureData{core::type::TextureDimension::k2d,
diff --git a/src/tint/lang/msl/writer/printer/var_test.cc b/src/tint/lang/msl/writer/var_test.cc
similarity index 74%
rename from src/tint/lang/msl/writer/printer/var_test.cc
rename to src/tint/lang/msl/writer/var_test.cc
index b91c2bd..4286c89 100644
--- a/src/tint/lang/msl/writer/printer/var_test.cc
+++ b/src/tint/lang/msl/writer/var_test.cc
@@ -25,7 +25,7 @@
 // OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
 // OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
 
-#include "src/tint/lang/msl/writer/printer/helper_test.h"
+#include "src/tint/lang/msl/writer/helper_test.h"
 
 namespace tint::msl::writer {
 namespace {
@@ -33,22 +33,22 @@
 using namespace tint::core::fluent_types;     // NOLINT
 using namespace tint::core::number_suffixes;  // NOLINT
 
-TEST_F(MslPrinterTest, VarF32) {
+TEST_F(MslWriterTest, VarF32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr<core::AddressSpace::kFunction, f32>());
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float a = 0.0f;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, VarI32) {
+TEST_F(MslWriterTest, VarI32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* v = b.Var("a", ty.ptr<core::AddressSpace::kFunction, i32>());
@@ -56,15 +56,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   int a = 1;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, VarU32) {
+TEST_F(MslWriterTest, VarU32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* v = b.Var("a", ty.ptr<core::AddressSpace::kFunction, u32>());
@@ -72,30 +72,30 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   uint a = 1u;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, VarArrayF32) {
+TEST_F(MslWriterTest, VarArrayF32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr<core::AddressSpace::kFunction, array<f32, 5>>());
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + MetalArray() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + MetalArray() + R"(
 void foo() {
   tint_array<float, 5> a = {};
 }
 )");
 }
 
-TEST_F(MslPrinterTest, VarStruct) {
+TEST_F(MslWriterTest, VarStruct) {
     auto* s = ty.Struct(mod.symbols.New("MyStruct"), {{mod.symbols.Register("a"), ty.f32()},  //
                                                       {mod.symbols.Register("b"), ty.vec4<i32>()}});
 
@@ -105,8 +105,8 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(struct MyStruct {
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(struct MyStruct {
   float a;
   int4 b;
 };
@@ -117,22 +117,22 @@
 )");
 }
 
-TEST_F(MslPrinterTest, VarVecF32) {
+TEST_F(MslWriterTest, VarVecF32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr<core::AddressSpace::kFunction, vec2<f32>>());
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float2 a = 0.0f;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, VarVecF16) {
+TEST_F(MslWriterTest, VarVecF16) {
     // Enable f16?
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -140,30 +140,30 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   half2 a = 0.0h;
 }
 )");
 }
 
-TEST_F(MslPrinterTest, VarMatF32) {
+TEST_F(MslWriterTest, VarMatF32) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         b.Var("a", ty.ptr<core::AddressSpace::kFunction, mat3x2<f32>>());
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float3x2 a = float3x2(0.0f);
 }
 )");
 }
 
-TEST_F(MslPrinterTest, VarMatF16) {
+TEST_F(MslWriterTest, VarMatF16) {
     // Enable f16?
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -171,15 +171,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   half3x2 a = half3x2(0.0h);
 }
 )");
 }
 
-TEST_F(MslPrinterTest, VarVecF32SplatZero) {
+TEST_F(MslWriterTest, VarVecF32SplatZero) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* v = b.Var("a", ty.ptr<core::AddressSpace::kFunction, vec3<f32>>());
@@ -187,15 +187,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float3 a = float3(0.0f);
 }
 )");
 }
 
-TEST_F(MslPrinterTest, VarVecF16SplatZero) {
+TEST_F(MslWriterTest, VarVecF16SplatZero) {
     // Enable f16
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -204,15 +204,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   half3 a = half3(0.0h);
 }
 )");
 }
 
-TEST_F(MslPrinterTest, VarMatF32SplatZero) {
+TEST_F(MslWriterTest, VarMatF32SplatZero) {
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
         auto* v = b.Var("a", ty.ptr<core::AddressSpace::kFunction, mat2x3<f32>>());
@@ -221,15 +221,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   float2x3 a = float2x3(float3(0.0f), float3(0.0f));
 }
 )");
 }
 
-TEST_F(MslPrinterTest, VarMatF16SplatZero) {
+TEST_F(MslWriterTest, VarMatF16SplatZero) {
     // Enable f16?
     auto* func = b.Function("foo", ty.void_());
     b.Append(func->Block(), [&] {
@@ -239,15 +239,15 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(
 void foo() {
   half2x3 a = half2x3(half3(0.0h), half3(0.0h));
 }
 )");
 }
 
-TEST_F(MslPrinterTest, VarGlobalPrivate) {
+TEST_F(MslWriterTest, VarGlobalPrivate) {
     core::ir::Var* v = nullptr;
     b.Append(mod.root_block, [&] {  //
         v = b.Var("v", ty.ptr<core::AddressSpace::kPrivate, f32>());
@@ -266,8 +266,8 @@
         b.Return(frag);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(struct tint_module_vars_struct {
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(struct tint_module_vars_struct {
   thread float* v;
 };
 
@@ -282,8 +282,7 @@
 )");
 }
 
-// TODO(jrprice): Requires ModuleScopeVarToEntryPointParam transform
-TEST_F(MslPrinterTest, DISABLED_VarGlobalWorkgroup) {
+TEST_F(MslWriterTest, VarGlobalWorkgroup) {
     core::ir::Var* v = nullptr;
     b.Append(mod.root_block,
              [&] { v = b.Var("v", ty.ptr<core::AddressSpace::kWorkgroup, f32>()); });
@@ -296,11 +295,13 @@
         b.Return(func);
     });
 
-    ASSERT_TRUE(Generate()) << err_ << output_;
-    EXPECT_EQ(output_, MetalHeader() + R"(
-threadgroup float v;
-void foo() {
-  float a = v;
+    ASSERT_TRUE(Generate()) << err_ << output_.msl;
+    EXPECT_EQ(output_.msl, MetalHeader() + R"(struct tint_module_vars_struct {
+  threadgroup float* v;
+};
+
+void foo(tint_module_vars_struct tint_module_vars) {
+  float a = (*tint_module_vars.v);
 }
 )");
 }
diff --git a/src/tint/lang/wgsl/ir_roundtrip_test.cc b/src/tint/lang/wgsl/ir_roundtrip_test.cc
index 39c7ca9..60345a5 100644
--- a/src/tint/lang/wgsl/ir_roundtrip_test.cc
+++ b/src/tint/lang/wgsl/ir_roundtrip_test.cc
@@ -426,7 +426,7 @@
 
 TEST_F(IRToProgramRoundtripTest, CoreBuiltinCall_PtrArg) {
     RUN_TEST(R"(
-@group(0) @binding(0) var<storage, read> v : array<u32>;
+@group(0u) @binding(0u) var<storage, read> v : array<u32>;
 
 fn foo() -> u32 {
   return arrayLength(&(v));
@@ -1888,7 +1888,7 @@
 
 TEST_F(IRToProgramRoundtripTest, PhonyAssign_HandleVar) {
     RUN_TEST(R"(
-@group(0) @binding(0) var t : texture_2d<f32>;
+@group(0u) @binding(0u) var t : texture_2d<f32>;
 
 fn f() {
   _ = t;
@@ -2143,36 +2143,36 @@
 
 TEST_F(IRToProgramRoundtripTest, ModuleScopeVar_Uniform_vec4i) {
     RUN_TEST(R"(
-@group(10) @binding(20) var<uniform> v : vec4<i32>;
+@group(10u) @binding(20u) var<uniform> v : vec4<i32>;
 )");
 }
 
 TEST_F(IRToProgramRoundtripTest, ModuleScopeVar_StorageRead_u32) {
     RUN_TEST(R"(
-@group(10) @binding(20) var<storage, read> v : u32;
+@group(10u) @binding(20u) var<storage, read> v : u32;
 )");
 }
 
 TEST_F(IRToProgramRoundtripTest, ModuleScopeVar_StorageReadWrite_i32) {
     RUN_TEST(R"(
-@group(10) @binding(20) var<storage, read_write> v : i32;
+@group(10u) @binding(20u) var<storage, read_write> v : i32;
 )");
 }
 TEST_F(IRToProgramRoundtripTest, ModuleScopeVar_Handle_Texture2D) {
     RUN_TEST(R"(
-@group(0) @binding(0) var t : texture_2d<f32>;
+@group(0u) @binding(0u) var t : texture_2d<f32>;
 )");
 }
 
 TEST_F(IRToProgramRoundtripTest, ModuleScopeVar_Handle_Sampler) {
     RUN_TEST(R"(
-@group(0) @binding(0) var s : sampler;
+@group(0u) @binding(0u) var s : sampler;
 )");
 }
 
 TEST_F(IRToProgramRoundtripTest, ModuleScopeVar_Handle_SamplerCmp) {
     RUN_TEST(R"(
-@group(0) @binding(0) var s : sampler_comparison;
+@group(0u) @binding(0u) var s : sampler_comparison;
 )");
 }
 
@@ -3314,5 +3314,42 @@
 )");
 }
 
+TEST_F(IRToProgramRoundtripTest, WorkgroupSizeLargerThanI32) {
+    RUN_TEST(R"(
+@compute @workgroup_size(4294967295u, 1u, 1u)
+fn main() {
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BindingLargerThanI32) {
+    RUN_TEST(R"(
+@group(0u) @binding(4000000000u) var s : sampler;
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, GroupLargerThanI32) {
+    RUN_TEST(R"(
+@group(4000000000u) @binding(0u) var s : sampler;
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, LocationInputLargerThanI32) {
+    RUN_TEST(R"(
+@fragment
+fn main(@location(4000000000u) color : vec4<f32>) {
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, LocationOutputLargerThanI32) {
+    RUN_TEST(R"(
+@fragment
+fn main() -> @location(4000000000u) vec4<f32> {
+  return vec4<f32>();
+}
+)");
+}
+
 }  // namespace
 }  // namespace tint::wgsl
diff --git a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
index 7243f25..4482524 100644
--- a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
+++ b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program.cc
@@ -252,7 +252,7 @@
                 }
             }
             if (auto loc = param->Location()) {
-                attrs.Push(b.Location(AInt(loc->value)));
+                attrs.Push(b.Location(u32(loc->value)));
                 if (auto interp = loc->interpolation) {
                     attrs.Push(b.Interpolate(interp->type, interp->sampling));
                 }
@@ -277,7 +277,7 @@
             case core::ir::Function::PipelineStage::kCompute: {
                 auto wgsize = fn->WorkgroupSize().value();
                 attrs.Push(b.Stage(ast::PipelineStage::kCompute));
-                attrs.Push(b.WorkgroupSize(AInt(wgsize[0]), AInt(wgsize[1]), AInt(wgsize[2])));
+                attrs.Push(b.WorkgroupSize(u32(wgsize[0]), u32(wgsize[1]), u32(wgsize[2])));
                 break;
             }
             case core::ir::Function::PipelineStage::kFragment:
@@ -305,7 +305,7 @@
             }
         }
         if (auto loc = fn->ReturnLocation()) {
-            ret_attrs.Push(b.Location(AInt(loc->value)));
+            ret_attrs.Push(b.Location(u32(loc->value)));
             if (auto interp = loc->interpolation) {
                 ret_attrs.Push(b.Interpolate(interp->type, interp->sampling));
             }
@@ -557,8 +557,8 @@
 
         Vector<const ast::Attribute*, 4> attrs;
         if (auto bp = var->BindingPoint()) {
-            attrs.Push(b.Group(AInt(bp->group)));
-            attrs.Push(b.Binding(AInt(bp->binding)));
+            attrs.Push(b.Group(u32(bp->group)));
+            attrs.Push(b.Binding(u32(bp->binding)));
         }
 
         const ast::Expression* init = nullptr;
diff --git a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
index 766469f..3619a0d 100644
--- a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
+++ b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
@@ -129,7 +129,7 @@
     fn->Block()->Append(b.Return(fn));
 
     EXPECT_WGSL(R"(
-@compute @workgroup_size(3, 4, 5)
+@compute @workgroup_size(3u, 4u, 5u)
 fn f() {
 }
 )");
@@ -212,7 +212,7 @@
 
     EXPECT_WGSL(R"(
 @fragment
-fn f() -> @location(1) vec4<f32> {
+fn f() -> @location(1u) vec4<f32> {
   return vec4<f32>();
 }
 )");
@@ -246,7 +246,7 @@
     EXPECT_WGSL(R"(
 enable chromium_experimental_subgroups;
 
-@compute @workgroup_size(3, 4, 5)
+@compute @workgroup_size(3u, 4u, 5u)
 fn f(@builtin(local_invocation_id) v : vec3<u32>, @builtin(local_invocation_index) v_1 : u32, @builtin(global_invocation_id) v_2 : vec3<u32>, @builtin(workgroup_id) v_3 : vec3<u32>, @builtin(num_workgroups) v_4 : vec3<u32>, @builtin(subgroup_invocation_id) v_5 : u32, @builtin(subgroup_size) v_6 : u32) {
 }
 )");
@@ -2226,7 +2226,7 @@
     });
 
     EXPECT_WGSL(R"(
-@group(0) @binding(0) var<storage, read_write> v : u32;
+@group(0u) @binding(0u) var<storage, read_write> v : u32;
 
 fn f() {
   for(v = (v + 1u); (v < 10u); ) {
@@ -2647,7 +2647,7 @@
     EXPECT_WGSL(R"(
 enable chromium_internal_graphite;
 
-@group(0) @binding(0) var T : texture_storage_2d<r8unorm, read>;
+@group(0u) @binding(0u) var T : texture_storage_2d<r8unorm, read>;
 )");
 }
 
diff --git a/src/tint/tint.gni b/src/tint/tint.gni
index 177fa29..8fdccd7 100644
--- a/src/tint/tint.gni
+++ b/src/tint/tint.gni
@@ -79,6 +79,7 @@
   }
   proto_library(target_name) {
     forward_variables_from(invoker, "*", [ "configs" ])
+    proto_in_dir = "${tint_root_dir}"
     generate_cc = true
     generate_python = false
     use_protobuf_full = true