diff --git a/src/tint/lang/core/type/manager.h b/src/tint/lang/core/type/manager.h
index 0e8b324..6c12a6c 100644
--- a/src/tint/lang/core/type/manager.h
+++ b/src/tint/lang/core/type/manager.h
@@ -65,13 +65,13 @@
 /// @returns the default access control for a memory view with the given address space.
 static constexpr inline core::Access DefaultAccessFor(core::AddressSpace space) {
     switch (space) {
+        case core::AddressSpace::kIn:
         case core::AddressSpace::kPushConstant:
         case core::AddressSpace::kUniform:
         case core::AddressSpace::kHandle:
             return core::Access::kRead;
 
         case core::AddressSpace::kUndefined:
-        case core::AddressSpace::kIn:
         case core::AddressSpace::kOut:
         case core::AddressSpace::kFunction:
         case core::AddressSpace::kPixelLocal:
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
index 63dcf81..718770e 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.cc
@@ -2122,8 +2122,7 @@
     };
 
     TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
-    Line() << "while (true) {";
-    EmitLoopPreserver();
+    Line() << IsolateUB() << " while(true) {";
     {
         ScopedIndent si(this);
         if (!EmitStatements(stmt->body->statements)) {
@@ -2193,8 +2192,7 @@
         };
 
         TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
-        Line() << "while (true) {";
-        EmitLoopPreserver();
+        Line() << IsolateUB() << " while(true) {";
         IncrementIndent();
         TINT_DEFER({
             DecrementIndent();
@@ -2217,7 +2215,7 @@
         // For-loop can be generated.
         {
             auto out = Line();
-            out << "for";
+            out << IsolateUB() << " for";
             {
                 ScopedParen sp(out);
 
@@ -2235,7 +2233,6 @@
             }
             out << " {";
         }
-        EmitLoopPreserver();
         {
             auto emit_continuing = [] { return true; };
             TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing);
@@ -2268,8 +2265,7 @@
     // as a regular while in MSL. Instead we need to generate a `while(true)` loop.
     bool emit_as_loop = cond_pre.lines.size() > 0;
     if (emit_as_loop) {
-        Line() << "while (true) {";
-        EmitLoopPreserver();
+        Line() << IsolateUB() << " while(true) {";
         IncrementIndent();
         TINT_DEFER({
             DecrementIndent();
@@ -2283,16 +2279,7 @@
         }
     } else {
         // While can be generated.
-        {
-            auto out = Line();
-            out << "while";
-            {
-                ScopedParen sp(out);
-                out << cond_buf.str();
-            }
-            out << " {";
-        }
-        EmitLoopPreserver();
+        Line() << IsolateUB() << " while(" << cond_buf.str() << ") {";
         if (!EmitStatementsWithIndent(stmt->body->statements)) {
             return false;
         }
@@ -3031,32 +3018,16 @@
     return true;
 }
 
-void ASTPrinter::EmitLoopPreserver() {
-    IncrementIndent();
-    // This statement prevents the MSL compiler from erasing a loop during
-    // optimizations.  In the AIR dialiect of LLVM IR, WGSL loops should compile
-    // to a loop that contains an 'asm' call with a 'sideeffect' annotation.
-    //
-    // For example, compile a WGSL file with a trivial while(1) loop to 'a.metal',
-    // then compile that to AIR (LLVM IR dialect):
-    //
-    //    xcrun metal a.metal -S -o -
-    //
-    // The loop in the AIR should look something like this:
-    //
-    //    1: ...
-    //      br label %2
-    //
-    //    2:                                      ; preds = %1, %2
-    //      tail call void asm sideeffect "", ""() #1, !srcloc !27
-    //      br label %2, !llvm.loop !28
-    //
-    // It is important that the 'sideeffect' annotation exist. That tells the
-    // optimizer that the instruction has side effects invisible to the
-    // optimizer, and therefore the loop should not be eliminated.
-    Line() << R"(__asm__("");)";
-
-    DecrementIndent();
+std::string_view ASTPrinter::IsolateUB() {
+    if (isolate_ub_macro_name_.empty()) {
+        isolate_ub_macro_name_ = UniqueIdentifier("TINT_ISOLATE_UB");
+        auto volatile_true = UniqueIdentifier("tint_volatile_true");
+        Line(&helpers_) << "#define " << isolate_ub_macro_name_ << " \\";
+        Line(&helpers_) << "  if (volatile bool " << volatile_true << " = true; " << volatile_true
+                        << ")";
+        Line(&helpers_);
+    }
+    return isolate_ub_macro_name_;
 }
 
 template <typename F>
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer.h b/src/tint/lang/msl/writer/ast_printer/ast_printer.h
index af251db..70269c9 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer.h
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer.h
@@ -378,9 +378,10 @@
                               const ast::CallExpression* expr,
                               const sem::BuiltinFn* builtin);
 
-    /// Emits a code sequence that preserves a loop during
-    /// optimizations even if the loop is infinite.
-    void EmitLoopPreserver();
+    /// Lazilly generates the TINT_ISOLATE_UB macro, used to prevent UB statements from affecting
+    /// later logic.
+    /// @return the unique name of the TINT_ISOLATE_UB macro.
+    std::string_view IsolateUB();
 
     /// Handles generating a builtin name
     /// @param builtin the semantic info for the builtin
@@ -437,6 +438,10 @@
 
     std::function<bool()> emit_continuing_;
 
+    /// The name of the macro used to prevent UB affecting later control flow.
+    /// Do not use this directly, instead call IsolateUB().
+    std::string isolate_ub_macro_name_;
+
     /// Name of atomicCompareExchangeWeak() helper for the given pointer storage
     /// class and struct return type
     using ACEWKeyType =
diff --git a/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc b/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
index 0477e48..74d8fdd 100644
--- a/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/ast_printer_test.cc
@@ -226,13 +226,15 @@
     T elements[N];
 };
 
+#define TINT_ISOLATE_UB \
+  if (volatile bool tint_volatile_true = true; tint_volatile_true)
+
 struct tint_symbol_3 {
   tint_array<float2x2, 4> m;
 };
 
 void comp_main_inner(uint local_invocation_index, threadgroup tint_array<float2x2, 4>* const tint_symbol) {
-  for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
-    __asm__("");
+  TINT_ISOLATE_UB for(uint idx = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
     uint const i = idx;
     (*(tint_symbol))[i] = float2x2(float2(0.0f), float2(0.0f));
   }
diff --git a/src/tint/lang/msl/writer/ast_printer/continue_test.cc b/src/tint/lang/msl/writer/ast_printer/continue_test.cc
index a79ddc3..727f05c 100644
--- a/src/tint/lang/msl/writer/ast_printer/continue_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/continue_test.cc
@@ -39,16 +39,24 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
+    ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
+    EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
 
-    ASSERT_TRUE(gen.EmitStatement(loop)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while (true) {
-    __asm__("");
+using namespace metal;
+
+#define TINT_ISOLATE_UB \
+  if (volatile bool tint_volatile_true = true; tint_volatile_true)
+
+kernel void test_function() {
+  TINT_ISOLATE_UB while(true) {
     if (false) {
       break;
     }
     continue;
   }
+  return;
+}
+
 )");
 }
 
diff --git a/src/tint/lang/msl/writer/ast_printer/loop_test.cc b/src/tint/lang/msl/writer/ast_printer/loop_test.cc
index 450592b..7aed227 100644
--- a/src/tint/lang/msl/writer/ast_printer/loop_test.cc
+++ b/src/tint/lang/msl/writer/ast_printer/loop_test.cc
@@ -46,13 +46,21 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
+    ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
+    EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
 
-    ASSERT_TRUE(gen.EmitStatement(l)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while (true) {
-    __asm__("");
+using namespace metal;
+
+#define TINT_ISOLATE_UB \
+  if (volatile bool tint_volatile_true = true; tint_volatile_true)
+
+fragment void F() {
+  TINT_ISOLATE_UB while(true) {
     break;
   }
+  return;
+}
+
 )");
 }
 
@@ -67,16 +75,27 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
+    ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
+    EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
 
-    ASSERT_TRUE(gen.EmitStatement(l)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while (true) {
-    __asm__("");
+using namespace metal;
+
+#define TINT_ISOLATE_UB \
+  if (volatile bool tint_volatile_true = true; tint_volatile_true)
+
+void a_statement() {
+}
+
+fragment void F() {
+  TINT_ISOLATE_UB while(true) {
     break;
     {
       a_statement();
     }
   }
+  return;
+}
+
 )");
 }
 
@@ -91,17 +110,28 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
+    ASSERT_TRUE(gen.Generate()) << gen.Diagnostics();
+    EXPECT_EQ(gen.Result(), R"(#include <metal_stdlib>
 
-    ASSERT_TRUE(gen.EmitStatement(l)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while (true) {
-    __asm__("");
+using namespace metal;
+
+#define TINT_ISOLATE_UB \
+  if (volatile bool tint_volatile_true = true; tint_volatile_true)
+
+void a_statement() {
+}
+
+fragment void F() {
+  TINT_ISOLATE_UB while(true) {
     break;
     {
       a_statement();
       if (true) { break; }
     }
   }
+  return;
+}
+
 )");
 }
 
@@ -125,23 +155,19 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(outer)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while (true) {
-    __asm__("");
-    while (true) {
-      __asm__("");
-      break;
-      {
-        a_statement();
-      }
-    }
+    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while(true) {
+  TINT_ISOLATE_UB while(true) {
+    break;
     {
-      lhs = rhs;
-      if (true) { break; }
+      a_statement();
     }
   }
+  {
+    lhs = rhs;
+    if (true) { break; }
+  }
+}
 )");
 }
 
@@ -167,18 +193,15 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(outer)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while (true) {
-    __asm__("");
-    float lhs = 2.5f;
-    float other = 0.0f;
-    break;
-    {
-      lhs = rhs;
-    }
+    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while(true) {
+  float lhs = 2.5f;
+  float other = 0.0f;
+  break;
+  {
+    lhs = rhs;
   }
+}
 )");
 }
 
@@ -193,13 +216,10 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  for(; ; ) {
-    __asm__("");
-    return;
-  }
+    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB for(; ; ) {
+  return;
+}
 )");
 }
 
@@ -214,13 +234,10 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  for(int i = 0; ; ) {
-    __asm__("");
-    return;
-  }
+    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB for(int i = 0; ; ) {
+  return;
+}
 )");
 }
 
@@ -243,19 +260,16 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(loop)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  {
-    {
-      f(1);
-      f(2);
-    }
-    for(; ; ) {
-      __asm__("");
-      return;
-    }
+    EXPECT_EQ(gen.Result(), R"({
+  {
+    f(1);
+    f(2);
   }
+  TINT_ISOLATE_UB for(; ; ) {
+    return;
+  }
+}
 )");
 }
 
@@ -270,13 +284,10 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  for(; true; ) {
-    __asm__("");
-    return;
-  }
+    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB for(; true; ) {
+  return;
+}
 )");
 }
 
@@ -292,14 +303,11 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
     EXPECT_EQ(gen.Result(),
-              R"(  for(; ; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
-    return;
-  }
+              R"(TINT_ISOLATE_UB for(; ; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  return;
+}
 )");
 }
 
@@ -322,17 +330,14 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(loop)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while (true) {
-    __asm__("");
-    return;
-    {
-      f(1);
-      f(2);
-    }
+    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while(true) {
+  return;
+  {
+    f(1);
+    f(2);
   }
+}
 )");
 }
 
@@ -349,14 +354,12 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(),
-              R"(  for(int i = 0; true; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
-    __asm__("");
-    a_statement();
-  }
+    EXPECT_EQ(
+        gen.Result(),
+        R"(TINT_ISOLATE_UB for(int i = 0; true; i = as_type<int>((as_type<uint>(i) + as_type<uint>(1)))) {
+  a_statement();
+}
 )");
 }
 
@@ -380,24 +383,21 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(loop)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  {
+    EXPECT_EQ(gen.Result(), R"({
+  {
+    f(1);
+    f(2);
+  }
+  TINT_ISOLATE_UB while(true) {
+    if (!(true)) { break; }
+    return;
     {
-      f(1);
-      f(2);
-    }
-    while (true) {
-      __asm__("");
-      if (!(true)) { break; }
-      return;
-      {
-        f(3);
-        f(4);
-      }
+      f(3);
+      f(4);
     }
   }
+}
 )");
 }
 
@@ -411,13 +411,10 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while(true) {
-    __asm__("");
-    return;
-  }
+    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while(true) {
+  return;
+}
 )");
 }
 
@@ -431,13 +428,10 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while(true) {
-    __asm__("");
-    continue;
-  }
+    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while(true) {
+  continue;
+}
 )");
 }
 
@@ -454,13 +448,10 @@
 
     ASTPrinter& gen = Build();
 
-    gen.IncrementIndent();
-
     ASSERT_TRUE(gen.EmitStatement(f)) << gen.Diagnostics();
-    EXPECT_EQ(gen.Result(), R"(  while((t && false)) {
-    __asm__("");
-    return;
-  }
+    EXPECT_EQ(gen.Result(), R"(TINT_ISOLATE_UB while((t && false)) {
+  return;
+}
 )");
 }
 
diff --git a/src/tint/lang/msl/writer/ast_raise/packed_vec3.cc b/src/tint/lang/msl/writer/ast_raise/packed_vec3.cc
index 9d651c6..a61e3f8 100644
--- a/src/tint/lang/msl/writer/ast_raise/packed_vec3.cc
+++ b/src/tint/lang/msl/writer/ast_raise/packed_vec3.cc
@@ -58,6 +58,10 @@
 
 namespace tint::msl::writer {
 
+/// Arrays larger than this will be packed/unpacked with a for loop.
+/// Arrays up to this size will be packed/unpacked with a sequence of statements.
+static constexpr uint32_t kMaxSeriallyUnpackedArraySize = 8;
+
 /// PIMPL state for the transform
 struct PackedVec3::State {
     /// Constructor
@@ -242,7 +246,6 @@
         const std::function<ast::Type()>& out_type) {
         // Allocate a variable to hold the return value of the function.
         tint::Vector<const ast::Statement*, 4> statements;
-        statements.Push(b.Decl(b.Var("result", out_type())));
 
         // Helper that generates a loop to copy and pack/unpack elements of an array to the result:
         //   for (var i = 0u; i < num_elements; i = i + 1) {
@@ -250,15 +253,28 @@
         //   }
         auto copy_array_elements = [&](uint32_t num_elements,
                                        const core::type::Type* element_type) {
-            // Generate an expression for packing or unpacking an element of the array.
-            auto* element = pack_or_unpack_element(b.IndexAccessor("in", "i"), element_type);
-            statements.Push(b.For(                   //
-                b.Decl(b.Var("i", b.ty.u32())),      //
-                b.LessThan("i", u32(num_elements)),  //
-                b.Assign("i", b.Add("i", 1_a)),      //
-                b.Block(tint::Vector{
-                    b.Assign(b.IndexAccessor("result", "i"), element),
-                })));
+            // Generate code for unpacking the array.
+            if (num_elements <= kMaxSeriallyUnpackedArraySize) {
+                // Generate a variable with an explicit initializer.
+                tint::Vector<const ast::Expression*, 8> elements;
+                for (uint32_t i = 0; i < num_elements; i++) {
+                    elements.Push(pack_or_unpack_element(
+                        b.IndexAccessor("in", b.Expr(core::AInt(i))), element_type));
+                }
+                statements.Push(b.Decl(b.Var("result", b.Call(out_type(), b.ExprList(elements)))));
+            } else {
+                statements.Push(b.Decl(b.Var("result", out_type())));
+                // Generate a for loop.
+                // Generate an expression for packing or unpacking an element of the array.
+                auto* element = pack_or_unpack_element(b.IndexAccessor("in", "i"), element_type);
+                statements.Push(b.For(                   //
+                    b.Decl(b.Var("i", b.ty.u32())),      //
+                    b.LessThan("i", u32(num_elements)),  //
+                    b.Assign("i", b.Add("i", 1_a)),      //
+                    b.Block(tint::Vector{
+                        b.Assign(b.IndexAccessor("result", "i"), element),
+                    })));
+            }
         };
 
         // Copy the elements of the value over to the result.
@@ -272,6 +288,7 @@
                 copy_array_elements(mat->columns(), mat->ColumnType());
             },
             [&](const core::type::Struct* str) {
+                statements.Push(b.Decl(b.Var("result", out_type())));
                 // Copy the struct members over one at a time, packing/unpacking as necessary.
                 for (auto* member : str->Members()) {
                     const ast::Expression* element =
diff --git a/src/tint/lang/msl/writer/ast_raise/packed_vec3_test.cc b/src/tint/lang/msl/writer/ast_raise/packed_vec3_test.cc
index 220dabb..cc002ba 100644
--- a/src/tint/lang/msl/writer/ast_raise/packed_vec3_test.cc
+++ b/src/tint/lang/msl/writer/ast_raise/packed_vec3_test.cc
@@ -502,7 +502,7 @@
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(PackedVec3Test, ArrayOfVec3_ReadArray) {
+TEST_F(PackedVec3Test, ArrayOfVec3_ReadArray_Small) {
     auto* src = R"(
 @group(0) @binding(0) var<storage> arr : array<vec3<f32>, 4>;
 
@@ -520,14 +520,49 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
-  var result : array<vec3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
+  var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
+  return result;
+}
+
+@group(0) @binding(0) var<storage> arr : array<tint_packed_vec3_f32_array_element, 4u>;
+
+fn f() {
+  let x = tint_unpack_vec3_in_composite(arr);
+}
+)";
+
+    ast::transform::DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ArrayOfVec3_ReadArray_Large) {
+    auto* src = R"(
+@group(0) @binding(0) var<storage> arr : array<vec3<f32>, 9>;
+
+fn f() {
+  let x = arr;
+}
+)";
+
+    auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+  @align(16)
+  elements : __packed_vec3<f32>,
+}
+
+fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 9u>) -> array<vec3<f32>, 9u> {
+  var result : array<vec3<f32>, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
     result[i] = vec3<f32>(in[i].elements);
   }
   return result;
 }
 
-@group(0) @binding(0) var<storage> arr : array<tint_packed_vec3_f32_array_element, 4u>;
+@group(0) @binding(0) var<storage> arr : array<tint_packed_vec3_f32_array_element, 9u>;
 
 fn f() {
   let x = tint_unpack_vec3_in_composite(arr);
@@ -630,7 +665,7 @@
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(PackedVec3Test, ArrayOfVec3_WriteArray_ValueRHS) {
+TEST_F(PackedVec3Test, ArrayOfVec3_WriteArray_ValueRHS_Small) {
     auto* src = R"(
 @group(0) @binding(0) var<storage, read_write> arr : array<vec3<f32>, 2>;
 
@@ -648,10 +683,7 @@
 }
 
 fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 2u>) -> array<tint_packed_vec3_f32_array_element, 2u> {
-  var result : array<tint_packed_vec3_f32_array_element, 2u>;
-  for(var i : u32; (i < 2u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 2u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])));
   return result;
 }
 
@@ -668,6 +700,52 @@
     EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(PackedVec3Test, ArrayOfVec3_WriteArray_ValueRHS_Large) {
+    auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> arr : array<vec3<f32>, 9>;
+
+fn f() {
+  arr = array(vec3(1.5, 2.5, 3.5),
+              vec3(4.5, 5.5, 6.5),
+              vec3(7.5, 8.5, 9.5),
+              vec3(7.5, 8.5, 9.5),
+              vec3(7.5, 8.5, 9.5),
+              vec3(7.5, 8.5, 9.5),
+              vec3(7.5, 8.5, 9.5),
+              vec3(7.5, 8.5, 9.5),
+              vec3(7.5, 8.5, 9.5));
+}
+)";
+
+    auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+  @align(16)
+  elements : __packed_vec3<f32>,
+}
+
+fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 9u>) -> array<tint_packed_vec3_f32_array_element, 9u> {
+  var result : array<tint_packed_vec3_f32_array_element, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
+    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
+  }
+  return result;
+}
+
+@group(0) @binding(0) var<storage, read_write> arr : array<tint_packed_vec3_f32_array_element, 9u>;
+
+fn f() {
+  arr = tint_pack_vec3_in_composite(array(vec3(1.5, 2.5, 3.5), vec3(4.5, 5.5, 6.5), vec3(7.5, 8.5, 9.5), vec3(7.5, 8.5, 9.5), vec3(7.5, 8.5, 9.5), vec3(7.5, 8.5, 9.5), vec3(7.5, 8.5, 9.5), vec3(7.5, 8.5, 9.5), vec3(7.5, 8.5, 9.5)));
+}
+)";
+
+    ast::transform::DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(PackedVec3Test, ArrayOfVec3_WriteArray_RefRHS) {
     auto* src = R"(
 @group(0) @binding(0) var<storage, read_write> arr : array<vec3<f32>, 2>;
@@ -731,7 +809,7 @@
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(PackedVec3Test, ArrayOfVec3_WriteVector_RefRHS) {
+TEST_F(PackedVec3Test, ArrayOfVec3_WriteVector_RefRHS_Small) {
     auto* src = R"(
 @group(0) @binding(0) var<storage, read_write> arr : array<vec3<f32>, 4>;
 @group(0) @binding(1) var<uniform> in_arr : array<vec3<f32>, 4>;
@@ -769,6 +847,44 @@
     EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(PackedVec3Test, ArrayOfVec3_WriteVector_RefRHS_Large) {
+    auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> arr : array<vec3<f32>, 9>;
+@group(0) @binding(1) var<uniform> in_arr : array<vec3<f32>, 9>;
+@group(0) @binding(2) var<uniform> in_vec : vec3<f32>;
+
+fn f() {
+  arr[0] = in_arr[0];
+  arr[1] = in_vec;
+}
+)";
+
+    auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+  @align(16)
+  elements : __packed_vec3<f32>,
+}
+
+@group(0) @binding(0) var<storage, read_write> arr : array<tint_packed_vec3_f32_array_element, 9u>;
+
+@group(0) @binding(1) var<uniform> in_arr : array<tint_packed_vec3_f32_array_element, 9u>;
+
+@group(0) @binding(2) var<uniform> in_vec : __packed_vec3<f32>;
+
+fn f() {
+  arr[0].elements = in_arr[0].elements;
+  arr[1].elements = in_vec;
+}
+)";
+
+    ast::transform::DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(PackedVec3Test, ArrayOfVec3_WriteComponent_MemberAccessor) {
     auto* src = R"(
 @group(0) @binding(0) var<storage, read_write> arr : array<vec3<f32>, 4>;
@@ -847,10 +963,7 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
@@ -975,10 +1088,7 @@
 }
 
 fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
-  var result : array<tint_packed_vec3_f32_array_element, 3u>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
   return result;
 }
 
@@ -1156,7 +1266,7 @@
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(PackedVec3Test, ArrayOfMatrix_ReadArray) {
+TEST_F(PackedVec3Test, ArrayOfMatrix_ReadArray_Small) {
     auto* src = R"(
 @group(0) @binding(0) var<storage> arr : array<mat3x3<f32>, 4>;
 
@@ -1174,22 +1284,59 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
-  var result : array<mat3x3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
+  var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
+  return result;
+}
+
+@group(0) @binding(0) var<storage> arr : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>;
+
+fn f() {
+  let x = tint_unpack_vec3_in_composite_1(arr);
+}
+)";
+
+    ast::transform::DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, ArrayOfMatrix_ReadArray_Large) {
+    auto* src = R"(
+@group(0) @binding(0) var<storage> arr : array<mat3x3<f32>, 9>;
+
+fn f() {
+  let x = arr;
+}
+)";
+
+    auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+  @align(16)
+  elements : __packed_vec3<f32>,
+}
+
+fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
+  return result;
+}
+
+fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>) -> array<mat3x3<f32>, 9u> {
+  var result : array<mat3x3<f32>, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
     result[i] = tint_unpack_vec3_in_composite(in[i]);
   }
   return result;
 }
 
-@group(0) @binding(0) var<storage> arr : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>;
+@group(0) @binding(0) var<storage> arr : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>;
 
 fn f() {
   let x = tint_unpack_vec3_in_composite_1(arr);
@@ -1220,10 +1367,7 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
@@ -1330,7 +1474,7 @@
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(PackedVec3Test, ArrayOfMatrix_WriteArray_ValueRHS) {
+TEST_F(PackedVec3Test, ArrayOfMatrix_WriteArray_ValueRHS_Small) {
     auto* src = R"(
 @group(0) @binding(0) var<storage, read_write> arr : array<mat3x3<f32>, 2>;
 
@@ -1348,18 +1492,12 @@
 }
 
 fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
-  var result : array<tint_packed_vec3_f32_array_element, 3u>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
   return result;
 }
 
 fn tint_pack_vec3_in_composite_1(in : array<mat3x3<f32>, 2u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 2u> {
-  var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 2u>;
-  for(var i : u32; (i < 2u); i = (i + 1)) {
-    result[i] = tint_pack_vec3_in_composite(in[i]);
-  }
+  var result = array<array<tint_packed_vec3_f32_array_element, 3u>, 2u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]));
   return result;
 }
 
@@ -1376,6 +1514,57 @@
     EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(PackedVec3Test, ArrayOfMatrix_WriteArray_ValueRHS_Large) {
+    auto* src = R"(
+@group(0) @binding(0) var<storage, read_write> arr : array<mat3x3<f32>, 9>;
+
+fn f() {
+  arr = array(mat3x3<f32>(),
+              mat3x3(1.5, 2.5, 3.5, 4.5, 5.5, 6.5, 7.5, 8.5, 9.5),
+              mat3x3f(),
+              mat3x3f(),
+              mat3x3f(),
+              mat3x3f(),
+              mat3x3f(),
+              mat3x3f(),
+              mat3x3f());
+}
+)";
+
+    auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+  @align(16)
+  elements : __packed_vec3<f32>,
+}
+
+fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
+  var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
+  return result;
+}
+
+fn tint_pack_vec3_in_composite_1(in : array<mat3x3<f32>, 9u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 9u> {
+  var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
+    result[i] = tint_pack_vec3_in_composite(in[i]);
+  }
+  return result;
+}
+
+@group(0) @binding(0) var<storage, read_write> arr : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>;
+
+fn f() {
+  arr = tint_pack_vec3_in_composite_1(array(mat3x3<f32>(), mat3x3(1.5, 2.5, 3.5, 4.5, 5.5, 6.5, 7.5, 8.5, 9.5), mat3x3f(), mat3x3f(), mat3x3f(), mat3x3f(), mat3x3f(), mat3x3f(), mat3x3f()));
+}
+)";
+
+    ast::transform::DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(PackedVec3Test, ArrayOfMatrix_WriteArray_RefRHS) {
     auto* src = R"(
 @group(0) @binding(0) var<storage, read_write> arr : array<mat3x3<f32>, 2>;
@@ -1427,10 +1616,7 @@
 }
 
 fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
-  var result : array<tint_packed_vec3_f32_array_element, 3u>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
   return result;
 }
 
@@ -2025,7 +2211,7 @@
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_ReadStruct) {
+TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_ReadStruct_Small) {
     auto* src = R"(
 struct S {
   arr : array<vec3<f32>, 4>,
@@ -2052,10 +2238,7 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
-  var result : array<vec3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
   return result;
 }
 
@@ -2082,7 +2265,64 @@
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_ReadArray) {
+TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_ReadStruct_Large) {
+    auto* src = R"(
+struct S {
+  arr : array<vec3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = P;
+}
+)";
+
+    auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+  @align(16)
+  elements : __packed_vec3<f32>,
+}
+
+struct S_tint_packed_vec3 {
+  @align(16)
+  arr : array<tint_packed_vec3_f32_array_element, 9u>,
+}
+
+fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 9u>) -> array<vec3<f32>, 9u> {
+  var result : array<vec3<f32>, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
+    result[i] = vec3<f32>(in[i].elements);
+  }
+  return result;
+}
+
+fn tint_unpack_vec3_in_composite_1(in : S_tint_packed_vec3) -> S {
+  var result : S;
+  result.arr = tint_unpack_vec3_in_composite(in.arr);
+  return result;
+}
+
+struct S {
+  arr : array<vec3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage> P : S_tint_packed_vec3;
+
+fn f() {
+  let x = tint_unpack_vec3_in_composite_1(P);
+}
+)";
+
+    ast::transform::DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_ReadArray_Small) {
     auto* src = R"(
 struct S {
   arr : array<vec3<f32>, 4>,
@@ -2109,15 +2349,63 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
-  var result : array<vec3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
+  var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
+  return result;
+}
+
+struct S {
+  arr : array<vec3<f32>, 4>,
+}
+
+@group(0) @binding(0) var<storage> P : S_tint_packed_vec3;
+
+fn f() {
+  let x = tint_unpack_vec3_in_composite(P.arr);
+}
+)";
+
+    ast::transform::DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_ReadArray_Large) {
+    auto* src = R"(
+struct S {
+  arr : array<vec3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = P.arr;
+}
+)";
+
+    auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+  @align(16)
+  elements : __packed_vec3<f32>,
+}
+
+struct S_tint_packed_vec3 {
+  @align(16)
+  arr : array<tint_packed_vec3_f32_array_element, 9u>,
+}
+
+fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 9u>) -> array<vec3<f32>, 9u> {
+  var result : array<vec3<f32>, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
     result[i] = vec3<f32>(in[i].elements);
   }
   return result;
 }
 
 struct S {
-  arr : array<vec3<f32>, 4>,
+  arr : array<vec3<f32>, 9>,
 }
 
 @group(0) @binding(0) var<storage> P : S_tint_packed_vec3;
@@ -2262,7 +2550,7 @@
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_WriteStruct_ValueRHS) {
+TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_WriteStruct_ValueRHS_Small) {
     auto* src = R"(
 struct S {
   arr : array<vec3<f32>, 2>,
@@ -2290,10 +2578,7 @@
 }
 
 fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 2u>) -> array<tint_packed_vec3_f32_array_element, 2u> {
-  var result : array<tint_packed_vec3_f32_array_element, 2u>;
-  for(var i : u32; (i < 2u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 2u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])));
   return result;
 }
 
@@ -2320,6 +2605,72 @@
     EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_WriteStruct_ValueRHS_Large) {
+    auto* src = R"(
+struct S {
+  arr : array<vec3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S;
+
+fn f() {
+  P = S(array(vec3(1.5, 4.5, 7.5),
+              vec3(9.5, 6.5, 3.5),
+              vec3(9.5, 6.5, 3.5),
+              vec3(9.5, 6.5, 3.5),
+              vec3(9.5, 6.5, 3.5),
+              vec3(9.5, 6.5, 3.5),
+              vec3(9.5, 6.5, 3.5),
+              vec3(9.5, 6.5, 3.5),
+              vec3(9.5, 6.5, 3.5)));
+}
+)";
+
+    auto* expect =
+        R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+  @align(16)
+  elements : __packed_vec3<f32>,
+}
+
+struct S_tint_packed_vec3 {
+  @align(16)
+  arr : array<tint_packed_vec3_f32_array_element, 9u>,
+}
+
+fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 9u>) -> array<tint_packed_vec3_f32_array_element, 9u> {
+  var result : array<tint_packed_vec3_f32_array_element, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
+    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
+  }
+  return result;
+}
+
+fn tint_pack_vec3_in_composite_1(in : S) -> S_tint_packed_vec3 {
+  var result : S_tint_packed_vec3;
+  result.arr = tint_pack_vec3_in_composite(in.arr);
+  return result;
+}
+
+struct S {
+  arr : array<vec3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S_tint_packed_vec3;
+
+fn f() {
+  P = tint_pack_vec3_in_composite_1(S(array(vec3(1.5, 4.5, 7.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5))));
+}
+)";
+
+    ast::transform::DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_WriteStruct_RefRHS) {
     auto* src = R"(
 struct S {
@@ -2366,7 +2717,7 @@
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_WriteArray_ValueRHS) {
+TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_WriteArray_ValueRHS_Small) {
     auto* src = R"(
 struct S {
   arr : array<vec3<f32>, 2>,
@@ -2393,10 +2744,7 @@
 }
 
 fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 2u>) -> array<tint_packed_vec3_f32_array_element, 2u> {
-  var result : array<tint_packed_vec3_f32_array_element, 2u>;
-  for(var i : u32; (i < 2u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 2u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])));
   return result;
 }
 
@@ -2417,6 +2765,65 @@
     EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_WriteArray_ValueRHS_Large) {
+    auto* src = R"(
+struct S {
+  arr : array<vec3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S;
+
+fn f() {
+  P.arr = array(vec3(1.5, 4.5, 7.5),
+                vec3(9.5, 6.5, 3.5),
+                vec3(9.5, 6.5, 3.5),
+                vec3(9.5, 6.5, 3.5),
+                vec3(9.5, 6.5, 3.5),
+                vec3(9.5, 6.5, 3.5),
+                vec3(9.5, 6.5, 3.5),
+                vec3(9.5, 6.5, 3.5),
+                vec3(9.5, 6.5, 3.5));
+}
+)";
+
+    auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+  @align(16)
+  elements : __packed_vec3<f32>,
+}
+
+struct S_tint_packed_vec3 {
+  @align(16)
+  arr : array<tint_packed_vec3_f32_array_element, 9u>,
+}
+
+fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 9u>) -> array<tint_packed_vec3_f32_array_element, 9u> {
+  var result : array<tint_packed_vec3_f32_array_element, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
+    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
+  }
+  return result;
+}
+
+struct S {
+  arr : array<vec3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage, read_write> P : S_tint_packed_vec3;
+
+fn f() {
+  P.arr = tint_pack_vec3_in_composite(array(vec3(1.5, 4.5, 7.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5), vec3(9.5, 6.5, 3.5)));
+}
+)";
+
+    ast::transform::DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(PackedVec3Test, StructMember_ArrayOfVec3_WriteArray_RefRHS) {
     auto* src = R"(
 struct S {
@@ -2684,10 +3091,7 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
@@ -2741,10 +3145,7 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
@@ -2922,10 +3323,7 @@
 }
 
 fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
-  var result : array<tint_packed_vec3_f32_array_element, 3u>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
   return result;
 }
 
@@ -3025,10 +3423,7 @@
 }
 
 fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
-  var result : array<tint_packed_vec3_f32_array_element, 3u>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
   return result;
 }
 
@@ -3289,7 +3684,7 @@
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(PackedVec3Test, StructMember_ArrayOfMatrix_ReadStruct) {
+TEST_F(PackedVec3Test, StructMember_ArrayOfMatrix_ReadStruct_Small) {
     auto* src = R"(
 struct S {
   arr : array<mat3x3<f32>, 4>,
@@ -3316,18 +3711,12 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
-  var result : array<mat3x3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite(in[i]);
-  }
+  var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
   return result;
 }
 
@@ -3354,10 +3743,72 @@
     EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(PackedVec3Test, StructMember_ArrayOfMatrix_ReadStruct_Large) {
+    auto* src = R"(
+struct S {
+  arr : array<mat3x3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage> P : S;
+
+fn f() {
+  let x = P;
+}
+)";
+
+    auto* expect = R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+  @align(16)
+  elements : __packed_vec3<f32>,
+}
+
+struct S_tint_packed_vec3 {
+  @align(16)
+  arr : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>,
+}
+
+fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
+  return result;
+}
+
+fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>) -> array<mat3x3<f32>, 9u> {
+  var result : array<mat3x3<f32>, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
+    result[i] = tint_unpack_vec3_in_composite(in[i]);
+  }
+  return result;
+}
+
+fn tint_unpack_vec3_in_composite_2(in : S_tint_packed_vec3) -> S {
+  var result : S;
+  result.arr = tint_unpack_vec3_in_composite_1(in.arr);
+  return result;
+}
+
+struct S {
+  arr : array<mat3x3<f32>, 9>,
+}
+
+@group(0) @binding(0) var<storage> P : S_tint_packed_vec3;
+
+fn f() {
+  let x = tint_unpack_vec3_in_composite_2(P);
+}
+)";
+
+    ast::transform::DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(PackedVec3Test, StructMember_ArrayOfMatrix_ReadArray) {
     auto* src = R"(
 struct S {
-  arr : array<mat3x3<f32>, 4>,
+  arr : array<mat3x3<f32>, 9>,
 }
 
 @group(0) @binding(0) var<storage> P : S;
@@ -3377,27 +3828,24 @@
 
 struct S_tint_packed_vec3 {
   @align(16)
-  arr : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>,
+  arr : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>,
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
-fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
-  var result : array<mat3x3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
+fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>) -> array<mat3x3<f32>, 9u> {
+  var result : array<mat3x3<f32>, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
     result[i] = tint_unpack_vec3_in_composite(in[i]);
   }
   return result;
 }
 
 struct S {
-  arr : array<mat3x3<f32>, 4>,
+  arr : array<mat3x3<f32>, 9>,
 }
 
 @group(0) @binding(0) var<storage> P : S_tint_packed_vec3;
@@ -3440,10 +3888,7 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
@@ -3621,18 +4066,12 @@
 }
 
 fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
-  var result : array<tint_packed_vec3_f32_array_element, 3u>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
   return result;
 }
 
 fn tint_pack_vec3_in_composite_1(in : array<mat3x3<f32>, 2u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 2u> {
-  var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 2u>;
-  for(var i : u32; (i < 2u); i = (i + 1)) {
-    result[i] = tint_pack_vec3_in_composite(in[i]);
-  }
+  var result = array<array<tint_packed_vec3_f32_array_element, 3u>, 2u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]));
   return result;
 }
 
@@ -3732,18 +4171,12 @@
 }
 
 fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
-  var result : array<tint_packed_vec3_f32_array_element, 3u>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
   return result;
 }
 
 fn tint_pack_vec3_in_composite_1(in : array<mat3x3<f32>, 2u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 2u> {
-  var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 2u>;
-  for(var i : u32; (i < 2u); i = (i + 1)) {
-    result[i] = tint_pack_vec3_in_composite(in[i]);
-  }
+  var result = array<array<tint_packed_vec3_f32_array_element, 3u>, 2u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]));
   return result;
 }
 
@@ -3846,10 +4279,7 @@
 }
 
 fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
-  var result : array<tint_packed_vec3_f32_array_element, 3u>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
   return result;
 }
 
@@ -4510,10 +4940,7 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
-  var result : array<vec3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
   return result;
 }
 
@@ -4528,10 +4955,7 @@
 }
 
 fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
-  var result : array<tint_packed_vec3_f32_array_element, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
   return result;
 }
 
@@ -4759,10 +5183,7 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
-  var result : array<vec3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
   return result;
 }
 
@@ -4965,18 +5386,12 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
 fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
-  var result : array<tint_packed_vec3_f32_array_element, 3u>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
   return result;
 }
 
@@ -5047,18 +5462,12 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
-  var result : array<vec3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
   return result;
 }
 
 fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
-  var result : array<tint_packed_vec3_f32_array_element, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
   return result;
 }
 
@@ -5120,34 +5529,22 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
-  var result : array<mat3x3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite(in[i]);
-  }
+  var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
   return result;
 }
 
 fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
-  var result : array<tint_packed_vec3_f32_array_element, 3u>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
   return result;
 }
 
 fn tint_pack_vec3_in_composite_1(in : array<mat3x3<f32>, 4u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 4u> {
-  var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_pack_vec3_in_composite(in[i]);
-  }
+  var result = array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]), tint_pack_vec3_in_composite(in[2]), tint_pack_vec3_in_composite(in[3]));
   return result;
 }
 
@@ -5213,26 +5610,17 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_1(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
-  var result : array<vec3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_2(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
-  var result : array<mat3x3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite(in[i]);
-  }
+  var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
   return result;
 }
 
@@ -5246,26 +5634,17 @@
 }
 
 fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
-  var result : array<tint_packed_vec3_f32_array_element, 3u>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
   return result;
 }
 
 fn tint_pack_vec3_in_composite_1(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
-  var result : array<tint_packed_vec3_f32_array_element, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
   return result;
 }
 
 fn tint_pack_vec3_in_composite_2(in : array<mat3x3<f32>, 4u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 4u> {
-  var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_pack_vec3_in_composite(in[i]);
-  }
+  var result = array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]), tint_pack_vec3_in_composite(in[2]), tint_pack_vec3_in_composite(in[3]));
   return result;
 }
 
@@ -5460,18 +5839,12 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
 fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
-  var result : array<tint_packed_vec3_f32_array_element, 3u>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
   return result;
 }
 
@@ -5551,18 +5924,12 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
-  var result : array<vec3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
   return result;
 }
 
 fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
-  var result : array<tint_packed_vec3_f32_array_element, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
   return result;
 }
 
@@ -5635,34 +6002,22 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
-  var result : array<mat3x3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite(in[i]);
-  }
+  var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
   return result;
 }
 
 fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
-  var result : array<tint_packed_vec3_f32_array_element, 3u>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
   return result;
 }
 
 fn tint_pack_vec3_in_composite_1(in : array<mat3x3<f32>, 4u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 4u> {
-  var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_pack_vec3_in_composite(in[i]);
-  }
+  var result = array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]), tint_pack_vec3_in_composite(in[2]), tint_pack_vec3_in_composite(in[3]));
   return result;
 }
 
@@ -5741,26 +6096,17 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_1(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
-  var result : array<vec3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_2(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
-  var result : array<mat3x3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite(in[i]);
-  }
+  var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
   return result;
 }
 
@@ -5774,26 +6120,17 @@
 }
 
 fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
-  var result : array<tint_packed_vec3_f32_array_element, 3u>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
   return result;
 }
 
 fn tint_pack_vec3_in_composite_1(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
-  var result : array<tint_packed_vec3_f32_array_element, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
   return result;
 }
 
 fn tint_pack_vec3_in_composite_2(in : array<mat3x3<f32>, 4u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 4u> {
-  var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_pack_vec3_in_composite(in[i]);
-  }
+  var result = array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]), tint_pack_vec3_in_composite(in[2]), tint_pack_vec3_in_composite(in[3]));
   return result;
 }
 
@@ -5868,10 +6205,7 @@
 }
 
 fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
-  var result : array<tint_packed_vec3_f32_array_element, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
   return result;
 }
 
@@ -5947,10 +6281,7 @@
 }
 
 fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
-  var result : array<tint_packed_vec3_f32_array_element, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
   return result;
 }
 
@@ -6041,10 +6372,7 @@
 }
 
 fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
-  var result : array<tint_packed_vec3_f32_array_element, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
   return result;
 }
 
@@ -6133,26 +6461,17 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_1(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
-  var result : array<vec3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_2(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
-  var result : array<mat3x3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite(in[i]);
-  }
+  var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
   return result;
 }
 
@@ -6239,26 +6558,17 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_1(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
-  var result : array<vec3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_2(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
-  var result : array<mat3x3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite(in[i]);
-  }
+  var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
   return result;
 }
 
@@ -6484,10 +6794,7 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
@@ -6706,10 +7013,7 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
@@ -6933,10 +7237,7 @@
 }
 
 fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
-  var result : array<tint_packed_vec3_f32_array_element, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
   return result;
 }
 
@@ -7000,18 +7301,12 @@
 }
 
 fn tint_pack_vec3_in_composite(in : array<vec3<f32>, 4u>) -> array<tint_packed_vec3_f32_array_element, 4u> {
-  var result : array<tint_packed_vec3_f32_array_element, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 4u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[3])));
   return result;
 }
 
 fn tint_pack_vec3_in_composite_1(in : array<array<vec3<f32>, 4u>, 4u>) -> array<array<tint_packed_vec3_f32_array_element, 4u>, 4u> {
-  var result : array<array<tint_packed_vec3_f32_array_element, 4u>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_pack_vec3_in_composite(in[i]);
-  }
+  var result = array<array<tint_packed_vec3_f32_array_element, 4u>, 4u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]), tint_pack_vec3_in_composite(in[2]), tint_pack_vec3_in_composite(in[3]));
   return result;
 }
 
@@ -7075,10 +7370,7 @@
 }
 
 fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
-  var result : array<tint_packed_vec3_f32_array_element, 3u>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
   return result;
 }
 
@@ -7140,18 +7432,12 @@
 }
 
 fn tint_pack_vec3_in_composite(in : mat3x3<f32>) -> array<tint_packed_vec3_f32_array_element, 3u> {
-  var result : array<tint_packed_vec3_f32_array_element, 3u>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[i]));
-  }
+  var result = array<tint_packed_vec3_f32_array_element, 3u>(tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[0])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[1])), tint_packed_vec3_f32_array_element(__packed_vec3<f32>(in[2])));
   return result;
 }
 
 fn tint_pack_vec3_in_composite_1(in : array<mat3x3<f32>, 4u>) -> array<array<tint_packed_vec3_f32_array_element, 3u>, 4u> {
-  var result : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_pack_vec3_in_composite(in[i]);
-  }
+  var result = array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>(tint_pack_vec3_in_composite(in[0]), tint_pack_vec3_in_composite(in[1]), tint_pack_vec3_in_composite(in[2]), tint_pack_vec3_in_composite(in[3]));
   return result;
 }
 
@@ -7222,42 +7508,27 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
-  var result : array<vec3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 4u>, 4u>) -> array<array<vec3<f32>, 4u>, 4u> {
-  var result : array<array<vec3<f32>, 4u>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite(in[i]);
-  }
+  var result = array<array<vec3<f32>, 4u>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_2(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_3(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
-  var result : array<mat3x3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite_2(in[i]);
-  }
+  var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite_2(in[0]), tint_unpack_vec3_in_composite_2(in[1]), tint_unpack_vec3_in_composite_2(in[2]), tint_unpack_vec3_in_composite_2(in[3]));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_4(in : array<array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>, 4u>) -> array<array<mat3x3<f32>, 4u>, 4u> {
-  var result : array<array<mat3x3<f32>, 4u>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite_3(in[i]);
-  }
+  var result = array<array<mat3x3<f32>, 4u>, 4u>(tint_unpack_vec3_in_composite_3(in[0]), tint_unpack_vec3_in_composite_3(in[1]), tint_unpack_vec3_in_composite_3(in[2]), tint_unpack_vec3_in_composite_3(in[3]));
   return result;
 }
 
@@ -7269,10 +7540,7 @@
 }
 
 fn tint_unpack_vec3_in_composite_6(in : array<S_tint_packed_vec3, 4u>) -> array<S, 4u> {
-  var result : array<S, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite_5(in[i]);
-  }
+  var result = array<S, 4u>(tint_unpack_vec3_in_composite_5(in[0]), tint_unpack_vec3_in_composite_5(in[1]), tint_unpack_vec3_in_composite_5(in[2]), tint_unpack_vec3_in_composite_5(in[3]));
   return result;
 }
 
@@ -7304,7 +7572,7 @@
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(PackedVec3Test, NestedArrays_VarInit) {
+TEST_F(PackedVec3Test, NestedArrays_VarInit_Small) {
     auto* src = R"(
 struct S {
   arr_v : array<array<vec3<f32>, 4>, 4>,
@@ -7345,42 +7613,27 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
-  var result : array<vec3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 4u>, 4u>) -> array<array<vec3<f32>, 4u>, 4u> {
-  var result : array<array<vec3<f32>, 4u>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite(in[i]);
-  }
+  var result = array<array<vec3<f32>, 4u>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_2(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_3(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
-  var result : array<mat3x3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite_2(in[i]);
-  }
+  var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite_2(in[0]), tint_unpack_vec3_in_composite_2(in[1]), tint_unpack_vec3_in_composite_2(in[2]), tint_unpack_vec3_in_composite_2(in[3]));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_4(in : array<array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>, 4u>) -> array<array<mat3x3<f32>, 4u>, 4u> {
-  var result : array<array<mat3x3<f32>, 4u>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite_3(in[i]);
-  }
+  var result = array<array<mat3x3<f32>, 4u>, 4u>(tint_unpack_vec3_in_composite_3(in[0]), tint_unpack_vec3_in_composite_3(in[1]), tint_unpack_vec3_in_composite_3(in[2]), tint_unpack_vec3_in_composite_3(in[3]));
   return result;
 }
 
@@ -7392,10 +7645,7 @@
 }
 
 fn tint_unpack_vec3_in_composite_6(in : array<S_tint_packed_vec3, 4u>) -> array<S, 4u> {
-  var result : array<S, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite_5(in[i]);
-  }
+  var result = array<S, 4u>(tint_unpack_vec3_in_composite_5(in[0]), tint_unpack_vec3_in_composite_5(in[1]), tint_unpack_vec3_in_composite_5(in[2]), tint_unpack_vec3_in_composite_5(in[3]));
   return result;
 }
 
@@ -7427,7 +7677,127 @@
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(PackedVec3Test, NestedArrays_VarAssignment) {
+TEST_F(PackedVec3Test, NestedArrays_VarInit_Large) {
+    auto* src = R"(
+struct S {
+  arr_v : array<array<vec3<f32>, 9>, 9>,
+  arr_m : array<array<mat3x3<f32>, 9>, 9>,
+}
+
+@group(0) @binding(0) var<storage, read_write> arr_s : array<S, 9>;
+
+fn f() {
+  var full_var : array<S, 9> = arr_s;
+  var struct_var : S = arr_s[0];
+  var outer_arr_v_var : array<array<vec3<f32>, 9>, 9> = arr_s[0].arr_v;
+  var inner_arr_v_var : array<vec3<f32>, 9> = arr_s[0].arr_v[1];
+  var v_var : vec3<f32> = arr_s[0].arr_v[1][2];
+  var v_element_var : f32 = arr_s[0].arr_v[1][2].y;
+  var outer_arr_m_var : array<array<mat3x3<f32>, 9>, 9> = arr_s[0].arr_m;
+  var inner_arr_m_var : array<mat3x3<f32>, 9> = arr_s[0].arr_m[1];
+  var m_var : mat3x3<f32> = arr_s[0].arr_m[1][2];
+  var m_col_var : vec3<f32> = arr_s[0].arr_m[1][2][0];
+  var m_element_var : f32 = arr_s[0].arr_m[1][2][0].y;
+}
+)";
+
+    auto* expect =
+        R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+  @align(16)
+  elements : __packed_vec3<f32>,
+}
+
+struct S_tint_packed_vec3 {
+  @align(16)
+  arr_v : array<array<tint_packed_vec3_f32_array_element, 9u>, 9u>,
+  @align(16)
+  arr_m : array<array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>, 9u>,
+}
+
+fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 9u>) -> array<vec3<f32>, 9u> {
+  var result : array<vec3<f32>, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
+    result[i] = vec3<f32>(in[i].elements);
+  }
+  return result;
+}
+
+fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 9u>, 9u>) -> array<array<vec3<f32>, 9u>, 9u> {
+  var result : array<array<vec3<f32>, 9u>, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
+    result[i] = tint_unpack_vec3_in_composite(in[i]);
+  }
+  return result;
+}
+
+fn tint_unpack_vec3_in_composite_2(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
+  return result;
+}
+
+fn tint_unpack_vec3_in_composite_3(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>) -> array<mat3x3<f32>, 9u> {
+  var result : array<mat3x3<f32>, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
+    result[i] = tint_unpack_vec3_in_composite_2(in[i]);
+  }
+  return result;
+}
+
+fn tint_unpack_vec3_in_composite_4(in : array<array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>, 9u>) -> array<array<mat3x3<f32>, 9u>, 9u> {
+  var result : array<array<mat3x3<f32>, 9u>, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
+    result[i] = tint_unpack_vec3_in_composite_3(in[i]);
+  }
+  return result;
+}
+
+fn tint_unpack_vec3_in_composite_5(in : S_tint_packed_vec3) -> S {
+  var result : S;
+  result.arr_v = tint_unpack_vec3_in_composite_1(in.arr_v);
+  result.arr_m = tint_unpack_vec3_in_composite_4(in.arr_m);
+  return result;
+}
+
+fn tint_unpack_vec3_in_composite_6(in : array<S_tint_packed_vec3, 9u>) -> array<S, 9u> {
+  var result : array<S, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
+    result[i] = tint_unpack_vec3_in_composite_5(in[i]);
+  }
+  return result;
+}
+
+struct S {
+  arr_v : array<array<vec3<f32>, 9>, 9>,
+  arr_m : array<array<mat3x3<f32>, 9>, 9>,
+}
+
+@group(0) @binding(0) var<storage, read_write> arr_s : array<S_tint_packed_vec3, 9u>;
+
+fn f() {
+  var full_var : array<S, 9> = tint_unpack_vec3_in_composite_6(arr_s);
+  var struct_var : S = tint_unpack_vec3_in_composite_5(arr_s[0]);
+  var outer_arr_v_var : array<array<vec3<f32>, 9>, 9> = tint_unpack_vec3_in_composite_1(arr_s[0].arr_v);
+  var inner_arr_v_var : array<vec3<f32>, 9> = tint_unpack_vec3_in_composite(arr_s[0].arr_v[1]);
+  var v_var : vec3<f32> = vec3<f32>(arr_s[0].arr_v[1][2].elements);
+  var v_element_var : f32 = arr_s[0].arr_v[1][2].elements.y;
+  var outer_arr_m_var : array<array<mat3x3<f32>, 9>, 9> = tint_unpack_vec3_in_composite_4(arr_s[0].arr_m);
+  var inner_arr_m_var : array<mat3x3<f32>, 9> = tint_unpack_vec3_in_composite_3(arr_s[0].arr_m[1]);
+  var m_var : mat3x3<f32> = tint_unpack_vec3_in_composite_2(arr_s[0].arr_m[1][2]);
+  var m_col_var : vec3<f32> = vec3<f32>(arr_s[0].arr_m[1][2][0].elements);
+  var m_element_var : f32 = arr_s[0].arr_m[1][2][0].elements.y;
+}
+)";
+
+    ast::transform::DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PackedVec3Test, NestedArrays_VarAssignment_Small) {
     auto* src = R"(
 struct S {
   arr_v : array<array<vec3<f32>, 4>, 4>,
@@ -7480,42 +7850,27 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
-  var result : array<vec3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 4u>, 4u>) -> array<array<vec3<f32>, 4u>, 4u> {
-  var result : array<array<vec3<f32>, 4u>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite(in[i]);
-  }
+  var result = array<array<vec3<f32>, 4u>, 4u>(tint_unpack_vec3_in_composite(in[0]), tint_unpack_vec3_in_composite(in[1]), tint_unpack_vec3_in_composite(in[2]), tint_unpack_vec3_in_composite(in[3]));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_2(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_3(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>) -> array<mat3x3<f32>, 4u> {
-  var result : array<mat3x3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite_2(in[i]);
-  }
+  var result = array<mat3x3<f32>, 4u>(tint_unpack_vec3_in_composite_2(in[0]), tint_unpack_vec3_in_composite_2(in[1]), tint_unpack_vec3_in_composite_2(in[2]), tint_unpack_vec3_in_composite_2(in[3]));
   return result;
 }
 
 fn tint_unpack_vec3_in_composite_4(in : array<array<array<tint_packed_vec3_f32_array_element, 3u>, 4u>, 4u>) -> array<array<mat3x3<f32>, 4u>, 4u> {
-  var result : array<array<mat3x3<f32>, 4u>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite_3(in[i]);
-  }
+  var result = array<array<mat3x3<f32>, 4u>, 4u>(tint_unpack_vec3_in_composite_3(in[0]), tint_unpack_vec3_in_composite_3(in[1]), tint_unpack_vec3_in_composite_3(in[2]), tint_unpack_vec3_in_composite_3(in[3]));
   return result;
 }
 
@@ -7527,10 +7882,7 @@
 }
 
 fn tint_unpack_vec3_in_composite_6(in : array<S_tint_packed_vec3, 4u>) -> array<S, 4u> {
-  var result : array<S, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = tint_unpack_vec3_in_composite_5(in[i]);
-  }
+  var result = array<S, 4u>(tint_unpack_vec3_in_composite_5(in[0]), tint_unpack_vec3_in_composite_5(in[1]), tint_unpack_vec3_in_composite_5(in[2]), tint_unpack_vec3_in_composite_5(in[3]));
   return result;
 }
 
@@ -7573,6 +7925,149 @@
     EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(PackedVec3Test, NestedArrays_VarAssignment_Large) {
+    auto* src = R"(
+struct S {
+  arr_v : array<array<vec3<f32>, 9>, 9>,
+  arr_m : array<array<mat3x3<f32>, 9>, 9>,
+}
+
+@group(0) @binding(0) var<storage, read_write> arr_s : array<S, 9>;
+
+fn f() {
+  var full_var : array<S, 9>;
+  var struct_var : S;
+  var outer_arr_v_var : array<array<vec3<f32>, 9>, 9>;
+  var inner_arr_v_var : array<vec3<f32>, 9>;
+  var v_var : vec3<f32>;
+  var v_element_var : f32;
+  var outer_arr_m_var : array<array<mat3x3<f32>, 9>, 9>;
+  var inner_arr_m_var : array<mat3x3<f32>, 9>;
+  var m_var : mat3x3<f32>;
+  var m_col_var : vec3<f32>;
+  var m_element_var : f32;
+
+  full_var = arr_s;
+  struct_var = arr_s[0];
+  outer_arr_v_var = arr_s[0].arr_v;
+  inner_arr_v_var = arr_s[0].arr_v[1];
+  v_var = arr_s[0].arr_v[1][2];
+  v_element_var = arr_s[0].arr_v[1][2].y;
+  outer_arr_m_var = arr_s[0].arr_m;
+  inner_arr_m_var = arr_s[0].arr_m[1];
+  m_var = arr_s[0].arr_m[1][2];
+  m_col_var = arr_s[0].arr_m[1][2][0];
+  m_element_var = arr_s[0].arr_m[1][2][0].y;
+}
+)";
+
+    auto* expect =
+        R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+  @align(16)
+  elements : __packed_vec3<f32>,
+}
+
+struct S_tint_packed_vec3 {
+  @align(16)
+  arr_v : array<array<tint_packed_vec3_f32_array_element, 9u>, 9u>,
+  @align(16)
+  arr_m : array<array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>, 9u>,
+}
+
+fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 9u>) -> array<vec3<f32>, 9u> {
+  var result : array<vec3<f32>, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
+    result[i] = vec3<f32>(in[i].elements);
+  }
+  return result;
+}
+
+fn tint_unpack_vec3_in_composite_1(in : array<array<tint_packed_vec3_f32_array_element, 9u>, 9u>) -> array<array<vec3<f32>, 9u>, 9u> {
+  var result : array<array<vec3<f32>, 9u>, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
+    result[i] = tint_unpack_vec3_in_composite(in[i]);
+  }
+  return result;
+}
+
+fn tint_unpack_vec3_in_composite_2(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
+  return result;
+}
+
+fn tint_unpack_vec3_in_composite_3(in : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>) -> array<mat3x3<f32>, 9u> {
+  var result : array<mat3x3<f32>, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
+    result[i] = tint_unpack_vec3_in_composite_2(in[i]);
+  }
+  return result;
+}
+
+fn tint_unpack_vec3_in_composite_4(in : array<array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>, 9u>) -> array<array<mat3x3<f32>, 9u>, 9u> {
+  var result : array<array<mat3x3<f32>, 9u>, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
+    result[i] = tint_unpack_vec3_in_composite_3(in[i]);
+  }
+  return result;
+}
+
+fn tint_unpack_vec3_in_composite_5(in : S_tint_packed_vec3) -> S {
+  var result : S;
+  result.arr_v = tint_unpack_vec3_in_composite_1(in.arr_v);
+  result.arr_m = tint_unpack_vec3_in_composite_4(in.arr_m);
+  return result;
+}
+
+fn tint_unpack_vec3_in_composite_6(in : array<S_tint_packed_vec3, 9u>) -> array<S, 9u> {
+  var result : array<S, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
+    result[i] = tint_unpack_vec3_in_composite_5(in[i]);
+  }
+  return result;
+}
+
+struct S {
+  arr_v : array<array<vec3<f32>, 9>, 9>,
+  arr_m : array<array<mat3x3<f32>, 9>, 9>,
+}
+
+@group(0) @binding(0) var<storage, read_write> arr_s : array<S_tint_packed_vec3, 9u>;
+
+fn f() {
+  var full_var : array<S, 9>;
+  var struct_var : S;
+  var outer_arr_v_var : array<array<vec3<f32>, 9>, 9>;
+  var inner_arr_v_var : array<vec3<f32>, 9>;
+  var v_var : vec3<f32>;
+  var v_element_var : f32;
+  var outer_arr_m_var : array<array<mat3x3<f32>, 9>, 9>;
+  var inner_arr_m_var : array<mat3x3<f32>, 9>;
+  var m_var : mat3x3<f32>;
+  var m_col_var : vec3<f32>;
+  var m_element_var : f32;
+  full_var = tint_unpack_vec3_in_composite_6(arr_s);
+  struct_var = tint_unpack_vec3_in_composite_5(arr_s[0]);
+  outer_arr_v_var = tint_unpack_vec3_in_composite_1(arr_s[0].arr_v);
+  inner_arr_v_var = tint_unpack_vec3_in_composite(arr_s[0].arr_v[1]);
+  v_var = vec3<f32>(arr_s[0].arr_v[1][2].elements);
+  v_element_var = arr_s[0].arr_v[1][2].elements.y;
+  outer_arr_m_var = tint_unpack_vec3_in_composite_4(arr_s[0].arr_m);
+  inner_arr_m_var = tint_unpack_vec3_in_composite_3(arr_s[0].arr_m[1]);
+  m_var = tint_unpack_vec3_in_composite_2(arr_s[0].arr_m[1][2]);
+  m_col_var = vec3<f32>(arr_s[0].arr_m[1][2][0].elements);
+  m_element_var = arr_s[0].arr_m[1][2][0].elements.y;
+}
+)";
+
+    ast::transform::DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(PackedVec3Test, RuntimeSizedArray) {
     auto* src = R"(
 struct S {
@@ -7817,10 +8312,7 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
@@ -7939,10 +8431,7 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 3u>) -> mat3x3<f32> {
-  var result : mat3x3<f32>;
-  for(var i : u32; (i < 3u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = mat3x3<f32>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements));
   return result;
 }
 
@@ -8048,7 +8537,7 @@
     EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(PackedVec3Test, Aliases) {
+TEST_F(PackedVec3Test, Aliases_Small) {
     auto* src = R"(
 alias VecArray = array<vec3<f32>, 4>;
 alias MatArray = array<mat3x3<f32>, 4>;
@@ -8097,10 +8586,7 @@
 }
 
 fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 4u>) -> array<vec3<f32>, 4u> {
-  var result : array<vec3<f32>, 4u>;
-  for(var i : u32; (i < 4u); i = (i + 1)) {
-    result[i] = vec3<f32>(in[i].elements);
-  }
+  var result = array<vec3<f32>, 4u>(vec3<f32>(in[0].elements), vec3<f32>(in[1].elements), vec3<f32>(in[2].elements), vec3<f32>(in[3].elements));
   return result;
 }
 
@@ -8142,6 +8628,100 @@
     EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(PackedVec3Test, Aliases_Large) {
+    auto* src = R"(
+alias VecArray = array<vec3<f32>, 9>;
+alias MatArray = array<mat3x3<f32>, 9>;
+alias NestedArray = array<VecArray, 9>;
+
+struct S {
+  v : VecArray,
+  m : MatArray,
+  n : NestedArray,
+}
+
+@group(0) @binding(0) var<storage, read_write> s : S;
+@group(0) @binding(1) var<storage, read_write> arr_v : VecArray;
+@group(0) @binding(2) var<storage, read_write> arr_m : MatArray;
+@group(0) @binding(3) var<storage, read_write> arr_n : NestedArray;
+
+fn g(p : ptr<function, VecArray>) {
+}
+
+fn f() {
+  var f_arr_v : VecArray = s.v;
+  g(&f_arr_v);
+
+  arr_v = s.v;
+  arr_m[0] = s.m[0];
+  arr_n[1][2] = s.n[1][2];
+}
+)";
+
+    auto* expect =
+        R"(
+enable chromium_internal_relaxed_uniform_layout;
+
+struct tint_packed_vec3_f32_array_element {
+  @align(16)
+  elements : __packed_vec3<f32>,
+}
+
+struct S_tint_packed_vec3 {
+  @align(16)
+  v : array<tint_packed_vec3_f32_array_element, 9u>,
+  @align(16)
+  m : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>,
+  @align(16)
+  n : array<array<tint_packed_vec3_f32_array_element, 9u>, 9u>,
+}
+
+fn tint_unpack_vec3_in_composite(in : array<tint_packed_vec3_f32_array_element, 9u>) -> array<vec3<f32>, 9u> {
+  var result : array<vec3<f32>, 9u>;
+  for(var i : u32; (i < 9u); i = (i + 1)) {
+    result[i] = vec3<f32>(in[i].elements);
+  }
+  return result;
+}
+
+alias VecArray = array<vec3<f32>, 9>;
+
+alias MatArray = array<mat3x3<f32>, 9>;
+
+alias NestedArray = array<VecArray, 9>;
+
+struct S {
+  v : VecArray,
+  m : MatArray,
+  n : NestedArray,
+}
+
+@group(0) @binding(0) var<storage, read_write> s : S_tint_packed_vec3;
+
+@group(0) @binding(1) var<storage, read_write> arr_v : array<tint_packed_vec3_f32_array_element, 9u>;
+
+@group(0) @binding(2) var<storage, read_write> arr_m : array<array<tint_packed_vec3_f32_array_element, 3u>, 9u>;
+
+@group(0) @binding(3) var<storage, read_write> arr_n : array<array<tint_packed_vec3_f32_array_element, 9u>, 9u>;
+
+fn g(p : ptr<function, VecArray>) {
+}
+
+fn f() {
+  var f_arr_v : VecArray = tint_unpack_vec3_in_composite(s.v);
+  g(&(f_arr_v));
+  arr_v = s.v;
+  arr_m[0] = s.m[0];
+  arr_n[1][2].elements = s.n[1][2].elements;
+}
+)";
+
+    ast::transform::DataMap data;
+    auto got = Run<PackedVec3>(src, data);
+
+    EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(PackedVec3Test, Vec3Bool) {
     // Make sure that we don't rewrite vec3<bool> types, as the `packed_bool<n>` types are reserved
     // in MSL and might not be supported everywhere.
diff --git a/src/tint/lang/spirv/reader/parser/memory_test.cc b/src/tint/lang/spirv/reader/parser/memory_test.cc
index 9f8c75f..acb3972 100644
--- a/src/tint/lang/spirv/reader/parser/memory_test.cc
+++ b/src/tint/lang/spirv/reader/parser/memory_test.cc
@@ -760,5 +760,59 @@
 )");
 }
 
+TEST_F(SpirvParserTest, StorageBufferAccessMode) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpExtension "SPV_KHR_storage_buffer_storage_class"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpDecorate %str Block
+               OpMemberDecorate %str 0 Offset 0
+               OpDecorate %ro_var NonWritable
+               OpDecorate %ro_var DescriptorSet 1
+               OpDecorate %ro_var Binding 2
+               OpDecorate %rw_var DescriptorSet 1
+               OpDecorate %rw_var Binding 3
+       %void = OpTypeVoid
+        %u32 = OpTypeInt 32 0
+        %str = OpTypeStruct %u32
+    %u32_ptr = OpTypePointer StorageBuffer %u32
+    %str_ptr = OpTypePointer StorageBuffer %str
+    %ep_type = OpTypeFunction %void
+      %u32_0 = OpConstant %u32 0
+     %ro_var = OpVariable %str_ptr StorageBuffer
+     %rw_var = OpVariable %str_ptr StorageBuffer
+       %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+  %ro_access = OpAccessChain %u32_ptr %ro_var %u32_0
+  %rw_access = OpAccessChain %u32_ptr %rw_var %u32_0
+       %load = OpLoad %u32 %ro_access
+               OpStore %rw_access %load
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+tint_symbol_1 = struct @align(4) {
+  tint_symbol:u32 @offset(0)
+}
+
+%b1 = block {  # root
+  %1:ptr<storage, tint_symbol_1, read> = var @binding_point(1, 2)
+  %2:ptr<storage, tint_symbol_1, read_write> = var @binding_point(1, 3)
+}
+
+%main = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
+  %b2 = block {
+    %4:ptr<storage, u32, read> = access %1, 0u
+    %5:ptr<storage, u32, read_write> = access %2, 0u
+    %6:u32 = load %4
+    store %5, %6
+    ret
+  }
+}
+)");
+}
+
 }  // namespace
 }  // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index 6fe4eb9..cfbc8cc 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -76,6 +76,14 @@
             return Failure("failed to build the internal representation of the module");
         }
 
+        // Check for unsupported extensions.
+        for (const auto& ext : spirv_context_->extensions()) {
+            auto name = ext.GetOperand(0).AsString();
+            if (name != "SPV_KHR_storage_buffer_storage_class") {
+                return Failure("SPIR-V extension '" + name + "' is not supported");
+            }
+        }
+
         {
             TINT_SCOPED_ASSIGNMENT(current_block_, ir_.root_block);
             EmitModuleScopeVariables();
@@ -95,10 +103,16 @@
     /// @returns the Tint address space for a SPIR-V storage class
     core::AddressSpace AddressSpace(spv::StorageClass sc) {
         switch (sc) {
+            case spv::StorageClass::Input:
+                return core::AddressSpace::kIn;
+            case spv::StorageClass::Output:
+                return core::AddressSpace::kOut;
             case spv::StorageClass::Function:
                 return core::AddressSpace::kFunction;
             case spv::StorageClass::Private:
                 return core::AddressSpace::kPrivate;
+            case spv::StorageClass::StorageBuffer:
+                return core::AddressSpace::kStorage;
             case spv::StorageClass::Uniform:
                 return core::AddressSpace::kUniform;
             default:
@@ -108,10 +122,50 @@
         }
     }
 
+    /// @param b a SPIR-V BuiltIn
+    /// @returns the Tint builtin value for a SPIR-V BuiltIn decoration
+    core::BuiltinValue Builtin(spv::BuiltIn b) {
+        switch (b) {
+            case spv::BuiltIn::FragCoord:
+                return core::BuiltinValue::kPosition;
+            case spv::BuiltIn::FragDepth:
+                return core::BuiltinValue::kFragDepth;
+            case spv::BuiltIn::FrontFacing:
+                return core::BuiltinValue::kFrontFacing;
+            case spv::BuiltIn::GlobalInvocationId:
+                return core::BuiltinValue::kGlobalInvocationId;
+            case spv::BuiltIn::InstanceIndex:
+                return core::BuiltinValue::kInstanceIndex;
+            case spv::BuiltIn::LocalInvocationId:
+                return core::BuiltinValue::kLocalInvocationId;
+            case spv::BuiltIn::LocalInvocationIndex:
+                return core::BuiltinValue::kLocalInvocationIndex;
+            case spv::BuiltIn::NumWorkgroups:
+                return core::BuiltinValue::kNumWorkgroups;
+            case spv::BuiltIn::PointSize:
+                return core::BuiltinValue::kPointSize;
+            case spv::BuiltIn::Position:
+                return core::BuiltinValue::kPosition;
+            case spv::BuiltIn::SampleId:
+                return core::BuiltinValue::kSampleIndex;
+            case spv::BuiltIn::SampleMask:
+                return core::BuiltinValue::kSampleMask;
+            case spv::BuiltIn::VertexIndex:
+                return core::BuiltinValue::kVertexIndex;
+            case spv::BuiltIn::WorkgroupId:
+                return core::BuiltinValue::kWorkgroupId;
+            default:
+                TINT_UNIMPLEMENTED() << "unhandled SPIR-V BuiltIn: " << static_cast<uint32_t>(b);
+                return core::BuiltinValue::kUndefined;
+        }
+    }
+
     /// @param type a SPIR-V type object
+    /// @param access_mode an optional access mode (for pointers)
     /// @returns a Tint type object
-    const core::type::Type* Type(const spvtools::opt::analysis::Type* type) {
-        return types_.GetOrCreate(type, [&]() -> const core::type::Type* {
+    const core::type::Type* Type(const spvtools::opt::analysis::Type* type,
+                                 core::Access access_mode = core::Access::kUndefined) {
+        return types_.GetOrCreate(TypeKey{type, access_mode}, [&]() -> const core::type::Type* {
             switch (type->kind()) {
                 case spvtools::opt::analysis::Type::kVoid:
                     return ty_.void_();
@@ -156,7 +210,7 @@
                 case spvtools::opt::analysis::Type::kPointer: {
                     auto* ptr_ty = type->AsPointer();
                     return ty_.ptr(AddressSpace(ptr_ty->storage_class()),
-                                   Type(ptr_ty->pointee_type()));
+                                   Type(ptr_ty->pointee_type()), access_mode);
                 }
                 default:
                     TINT_UNIMPLEMENTED() << "unhandled SPIR-V type: " << type->str();
@@ -166,9 +220,10 @@
     }
 
     /// @param id a SPIR-V result ID for a type declaration instruction
+    /// @param access_mode an optional access mode (for pointers)
     /// @returns a Tint type object
-    const core::type::Type* Type(uint32_t id) {
-        return Type(spirv_context_->get_type_mgr()->GetType(id));
+    const core::type::Type* Type(uint32_t id, core::Access access_mode = core::Access::kUndefined) {
+        return Type(spirv_context_->get_type_mgr()->GetType(id), access_mode);
     }
 
     /// @param arr_ty a SPIR-V array object
@@ -209,6 +264,15 @@
             uint32_t align = std::max<uint32_t>(member_ty->Align(), 1u);
             uint32_t offset = tint::RoundUp(align, current_size);
             core::type::StructMemberAttributes attributes;
+            auto interpolation = [&]() -> core::Interpolation& {
+                // Create the interpolation field with the default values on first call.
+                if (!attributes.interpolation.has_value()) {
+                    attributes.interpolation =
+                        core::Interpolation{core::InterpolationType::kPerspective,
+                                            core::InterpolationSampling::kCenter};
+                }
+                return attributes.interpolation.value();
+            };
 
             // Handle member decorations that affect layout or attributes.
             if (struct_ty->element_decorations().count(i)) {
@@ -217,6 +281,28 @@
                         case spv::Decoration::Offset:
                             offset = deco[1];
                             break;
+                        case spv::Decoration::BuiltIn:
+                            attributes.builtin = Builtin(spv::BuiltIn(deco[1]));
+                            break;
+                        case spv::Decoration::Invariant:
+                            attributes.invariant = true;
+                            break;
+                        case spv::Decoration::Location:
+                            attributes.location = deco[1];
+                            break;
+                        case spv::Decoration::NoPerspective:
+                            interpolation().type = core::InterpolationType::kLinear;
+                            break;
+                        case spv::Decoration::Flat:
+                            interpolation().type = core::InterpolationType::kFlat;
+                            break;
+                        case spv::Decoration::Centroid:
+                            interpolation().sampling = core::InterpolationSampling::kCentroid;
+                            break;
+                        case spv::Decoration::Sample:
+                            interpolation().sampling = core::InterpolationSampling::kSample;
+                            break;
+
                         default:
                             TINT_UNIMPLEMENTED() << "unhandled member decoration: " << deco[0];
                             break;
@@ -459,7 +545,14 @@
             indices.Push(Value(inst.GetSingleWordOperand(i)));
         }
         auto* base = Value(inst.GetSingleWordOperand(2));
-        auto* access = b_.Access(Type(inst.type_id()), base, std::move(indices));
+
+        // Propagate the access mode of the base object.
+        auto access_mode = core::Access::kUndefined;
+        if (auto* ptr = base->Type()->As<core::type::Pointer>()) {
+            access_mode = ptr->Access();
+        }
+
+        auto* access = b_.Access(Type(inst.type_id(), access_mode), base, std::move(indices));
         Emit(access, inst.result_id());
     }
 
@@ -496,19 +589,25 @@
 
     /// @param inst the SPIR-V instruction for OpVariable
     void EmitVar(const spvtools::opt::Instruction& inst) {
-        auto* var = b_.Var(Type(inst.type_id())->As<core::type::Pointer>());
-        if (inst.NumOperands() > 3) {
-            var->SetInitializer(Value(inst.GetSingleWordOperand(3)));
-        }
-
         // Handle decorations.
         std::optional<uint32_t> group;
         std::optional<uint32_t> binding;
+        core::Access access_mode = core::Access::kUndefined;
+        core::ir::IOAttributes io_attributes;
+        auto interpolation = [&]() -> core::Interpolation& {
+            // Create the interpolation field with the default values on first call.
+            if (!io_attributes.interpolation.has_value()) {
+                io_attributes.interpolation = core::Interpolation{
+                    core::InterpolationType::kPerspective, core::InterpolationSampling::kCenter};
+            }
+            return io_attributes.interpolation.value();
+        };
         for (auto* deco :
              spirv_context_->get_decoration_mgr()->GetDecorationsFor(inst.result_id(), false)) {
             auto d = deco->GetSingleWordOperand(1);
             switch (spv::Decoration(d)) {
                 case spv::Decoration::NonWritable:
+                    access_mode = core::Access::kRead;
                     break;
                 case spv::Decoration::DescriptorSet:
                     group = deco->GetSingleWordOperand(2);
@@ -516,20 +615,70 @@
                 case spv::Decoration::Binding:
                     binding = deco->GetSingleWordOperand(2);
                     break;
+                case spv::Decoration::BuiltIn:
+                    io_attributes.builtin = Builtin(spv::BuiltIn(deco->GetSingleWordOperand(2)));
+                    break;
+                case spv::Decoration::Invariant:
+                    io_attributes.invariant = true;
+                    break;
+                case spv::Decoration::Location:
+                    io_attributes.location = deco->GetSingleWordOperand(2);
+                    break;
+                case spv::Decoration::NoPerspective:
+                    interpolation().type = core::InterpolationType::kLinear;
+                    break;
+                case spv::Decoration::Flat:
+                    interpolation().type = core::InterpolationType::kFlat;
+                    break;
+                case spv::Decoration::Centroid:
+                    interpolation().sampling = core::InterpolationSampling::kCentroid;
+                    break;
+                case spv::Decoration::Sample:
+                    interpolation().sampling = core::InterpolationSampling::kSample;
+                    break;
                 default:
                     TINT_UNIMPLEMENTED() << "unhandled decoration " << d;
                     break;
             }
         }
+
+        auto* var = b_.Var(Type(inst.type_id(), access_mode)->As<core::type::Pointer>());
+        if (inst.NumOperands() > 3) {
+            var->SetInitializer(Value(inst.GetSingleWordOperand(3)));
+        }
+
         if (group || binding) {
             TINT_ASSERT(group && binding);
             var->SetBindingPoint(group.value(), binding.value());
         }
+        var->SetAttributes(std::move(io_attributes));
 
         Emit(var, inst.result_id());
     }
 
   private:
+    /// TypeKey describes a SPIR-V type with an access mode.
+    struct TypeKey {
+        /// The SPIR-V type object.
+        const spvtools::opt::analysis::Type* type;
+        /// The access mode.
+        core::Access access_mode;
+
+        // Equality operator for TypeKey.
+        bool operator==(const TypeKey& other) const {
+            return type == other.type && access_mode == other.access_mode;
+        }
+
+        /// Hasher provides a hash function for the TypeKey.
+        struct Hasher {
+            /// @param tk the TypeKey to create a hash for
+            /// @return the hash value
+            inline std::size_t operator()(const TypeKey& tk) const {
+                return HashCombine(Hash(tk.type), tk.access_mode);
+            }
+        };
+    };
+
     /// The generated IR module.
     core::ir::Module ir_;
     /// The Tint IR builder.
@@ -541,8 +690,8 @@
     core::ir::Function* current_function_ = nullptr;
     /// The Tint IR block that is currently being emitted.
     core::ir::Block* current_block_ = nullptr;
-    /// A map from a SPIR-V type declaration result ID to the corresponding Tint type object.
-    Hashmap<const spvtools::opt::analysis::Type*, const core::type::Type*, 16> types_;
+    /// A map from a SPIR-V type declaration to the corresponding Tint type object.
+    Hashmap<TypeKey, const core::type::Type*, 16, TypeKey::Hasher> types_;
     /// A map from a SPIR-V function definition result ID to the corresponding Tint function object.
     Hashmap<uint32_t, core::ir::Function*, 8> functions_;
     /// A map from a SPIR-V result ID to the corresponding Tint value object.
diff --git a/src/tint/lang/spirv/reader/parser/struct_test.cc b/src/tint/lang/spirv/reader/parser/struct_test.cc
index afbaa39..bedb4e2 100644
--- a/src/tint/lang/spirv/reader/parser/struct_test.cc
+++ b/src/tint/lang/spirv/reader/parser/struct_test.cc
@@ -229,4 +229,132 @@
 )");
 }
 
+TEST_F(SpirvParserTest, Struct_Builtin) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpCapability SampleRateShading
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main"
+               OpExecutionMode %main OriginUpperLeft
+               OpMemberDecorate %str 0 BuiltIn Position
+       %void = OpTypeVoid
+        %f32 = OpTypeFloat 32
+      %vec4f = OpTypeVector %f32 4
+        %str = OpTypeStruct %vec4f
+    %fn_type = OpTypeFunction %void
+
+%_ptr_Output = OpTypePointer Output %str
+        %var = OpVariable %_ptr_Output Output
+
+       %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+tint_symbol_1 = struct @align(16) {
+  tint_symbol:vec4<f32> @offset(0), @builtin(position)
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Struct_Builtin_WithInvariant) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpCapability SampleRateShading
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main"
+               OpExecutionMode %main OriginUpperLeft
+               OpMemberDecorate %str 0 BuiltIn Position
+               OpMemberDecorate %str 0 Invariant
+       %void = OpTypeVoid
+        %f32 = OpTypeFloat 32
+      %vec4f = OpTypeVector %f32 4
+        %str = OpTypeStruct %vec4f
+    %fn_type = OpTypeFunction %void
+
+%_ptr_Output = OpTypePointer Output %str
+        %var = OpVariable %_ptr_Output Output
+
+       %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+tint_symbol_1 = struct @align(16) {
+  tint_symbol:vec4<f32> @offset(0), @invariant, @builtin(position)
+}
+)");
+}
+
+struct LocationCase {
+    std::string spirv_decorations;
+    std::string ir;
+};
+
+using LocationStructTest = SpirvParserTestWithParam<LocationCase>;
+
+TEST_P(LocationStructTest, MemberDecorations) {
+    auto params = GetParam();
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpCapability SampleRateShading
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main"
+               OpExecutionMode %main OriginUpperLeft
+          )" + params.spirv_decorations +
+                  R"(
+       %void = OpTypeVoid
+        %f32 = OpTypeFloat 32
+      %vec4f = OpTypeVector %f32 4
+        %str = OpTypeStruct %vec4f
+    %fn_type = OpTypeFunction %void
+
+%_ptr_Input = OpTypePointer Input %str
+        %var = OpVariable %_ptr_Input Input
+
+       %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+)",
+              params.ir);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    SpirvParser,
+    LocationStructTest,
+    testing::Values(
+        LocationCase{
+            "OpMemberDecorate %str 0 Location 1 ",
+            "tint_symbol:vec4<f32> @offset(0), @location(1)",
+        },
+        LocationCase{
+            "OpMemberDecorate %str 0 Location 2 "
+            "OpMemberDecorate %str 0 NoPerspective ",
+            "tint_symbol:vec4<f32> @offset(0), @location(2), @interpolate(linear, center)",
+        },
+        LocationCase{
+            "OpMemberDecorate %str 0 Location 3 "
+            "OpMemberDecorate %str 0 Flat ",
+            "tint_symbol:vec4<f32> @offset(0), @location(3), @interpolate(flat, center)",
+        },
+        LocationCase{
+            "OpMemberDecorate %str 0 Location 4 "
+            "OpMemberDecorate %str 0 Centroid ",
+            "tint_symbol:vec4<f32> @offset(0), @location(4), @interpolate(perspective, centroid)",
+        },
+        LocationCase{
+            "OpMemberDecorate %str 0 Location 5 "
+            "OpMemberDecorate %str 0 Sample ",
+            "tint_symbol:vec4<f32> @offset(0), @location(5), @interpolate(perspective, sample)",
+        },
+        LocationCase{
+            "OpMemberDecorate %str 0 Location 6 "
+            "OpMemberDecorate %str 0 NoPerspective "
+            "OpMemberDecorate %str 0 Centroid ",
+            "tint_symbol:vec4<f32> @offset(0), @location(6), @interpolate(linear, centroid)",
+        }));
+
 }  // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/parser/var_test.cc b/src/tint/lang/spirv/reader/parser/var_test.cc
index 88a6347..f524b23 100644
--- a/src/tint/lang/spirv/reader/parser/var_test.cc
+++ b/src/tint/lang/spirv/reader/parser/var_test.cc
@@ -142,6 +142,129 @@
 )");
 }
 
+TEST_F(SpirvParserTest, StorageVar_ReadOnly) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpExtension "SPV_KHR_storage_buffer_storage_class"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %1 "main"
+               OpExecutionMode %1 LocalSize 1 1 1
+               OpDecorate %str Block
+               OpMemberDecorate %str 0 Offset 0
+               OpDecorate %6 NonWritable
+               OpDecorate %6 DescriptorSet 1
+               OpDecorate %6 Binding 2
+       %void = OpTypeVoid
+       %uint = OpTypeInt 32 0
+        %str = OpTypeStruct %uint
+%_ptr_StorageBuffer_str = OpTypePointer StorageBuffer %str
+          %5 = OpTypeFunction %void
+          %6 = OpVariable %_ptr_StorageBuffer_str StorageBuffer
+          %1 = OpFunction %void None %5
+          %7 = OpLabel
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+tint_symbol_1 = struct @align(4) {
+  tint_symbol:u32 @offset(0)
+}
+
+%b1 = block {  # root
+  %1:ptr<storage, tint_symbol_1, read> = var @binding_point(1, 2)
+}
+
+%main = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
+  %b2 = block {
+    ret
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, StorageVar_ReadWrite) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpExtension "SPV_KHR_storage_buffer_storage_class"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %1 "main"
+               OpExecutionMode %1 LocalSize 1 1 1
+               OpDecorate %str Block
+               OpMemberDecorate %str 0 Offset 0
+               OpDecorate %6 DescriptorSet 1
+               OpDecorate %6 Binding 2
+       %void = OpTypeVoid
+       %uint = OpTypeInt 32 0
+        %str = OpTypeStruct %uint
+%_ptr_StorageBuffer_str = OpTypePointer StorageBuffer %str
+          %5 = OpTypeFunction %void
+          %6 = OpVariable %_ptr_StorageBuffer_str StorageBuffer
+          %1 = OpFunction %void None %5
+          %7 = OpLabel
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+tint_symbol_1 = struct @align(4) {
+  tint_symbol:u32 @offset(0)
+}
+
+%b1 = block {  # root
+  %1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(1, 2)
+}
+
+%main = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
+  %b2 = block {
+    ret
+  }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, StorageVar_ReadOnly_And_ReadWrite) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpExtension "SPV_KHR_storage_buffer_storage_class"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %1 "main"
+               OpExecutionMode %1 LocalSize 1 1 1
+               OpDecorate %str Block
+               OpMemberDecorate %str 0 Offset 0
+               OpDecorate %6 NonWritable
+               OpDecorate %6 DescriptorSet 1
+               OpDecorate %6 Binding 2
+               OpDecorate %7 DescriptorSet 1
+               OpDecorate %7 Binding 3
+       %void = OpTypeVoid
+       %uint = OpTypeInt 32 0
+        %str = OpTypeStruct %uint
+%_ptr_StorageBuffer_str = OpTypePointer StorageBuffer %str
+          %5 = OpTypeFunction %void
+          %6 = OpVariable %_ptr_StorageBuffer_str StorageBuffer
+          %7 = OpVariable %_ptr_StorageBuffer_str StorageBuffer
+          %1 = OpFunction %void None %5
+          %8 = OpLabel
+               OpReturn
+               OpFunctionEnd
+)",
+              R"(
+tint_symbol_1 = struct @align(4) {
+  tint_symbol:u32 @offset(0)
+}
+
+%b1 = block {  # root
+  %1:ptr<storage, tint_symbol_1, read> = var @binding_point(1, 2)
+  %2:ptr<storage, tint_symbol_1, read_write> = var @binding_point(1, 3)
+}
+
+%main = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
+  %b2 = block {
+    ret
+  }
+}
+)");
+}
+
 TEST_F(SpirvParserTest, UniformVar) {
     EXPECT_IR(R"(
                OpCapability Shader
@@ -181,5 +304,278 @@
 )");
 }
 
+struct BuiltinCase {
+    std::string spirv_type;
+    std::string spirv_builtin;
+    std::string ir;
+};
+std::string PrintBuiltinCase(testing::TestParamInfo<BuiltinCase> bc) {
+    return bc.param.spirv_builtin;
+}
+
+using BuiltinInputTest = SpirvParserTestWithParam<BuiltinCase>;
+
+TEST_P(BuiltinInputTest, Enum) {
+    auto params = GetParam();
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpCapability SampleRateShading
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpDecorate %var BuiltIn )" +
+                  params.spirv_builtin + R"(
+       %void = OpTypeVoid
+       %bool = OpTypeBool
+        %u32 = OpTypeInt 32 0
+      %vec3u = OpTypeVector %u32 3
+        %f32 = OpTypeFloat 32
+      %vec4f = OpTypeVector %f32 4
+      %u32_1 = OpConstant %u32 1
+  %arr_u32_1 = OpTypeArray %u32 %u32_1
+    %fn_type = OpTypeFunction %void
+
+ %_ptr_Input = OpTypePointer Input %)" +
+                  params.spirv_type + R"(
+        %var = OpVariable %_ptr_Input Input
+
+       %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+)",
+              params.ir);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    SpirvParser,
+    BuiltinInputTest,
+    testing::Values(
+        BuiltinCase{
+            "vec4f",
+            "FragCoord",
+            "%1:ptr<__in, vec4<f32>, read> = var @builtin(position)",
+        },
+        BuiltinCase{
+            "bool",
+            "FrontFacing",
+            "%1:ptr<__in, bool, read> = var @builtin(front_facing)",
+        },
+        BuiltinCase{
+            "vec3u",
+            "GlobalInvocationId",
+            "%1:ptr<__in, vec3<u32>, read> = var @builtin(global_invocation_id)",
+        },
+        BuiltinCase{
+            "u32",
+            "InstanceIndex",
+            "%1:ptr<__in, u32, read> = var @builtin(instance_index)",
+        },
+        BuiltinCase{
+            "vec3u",
+            "LocalInvocationId",
+            "%1:ptr<__in, vec3<u32>, read> = var @builtin(local_invocation_id)",
+        },
+        BuiltinCase{
+            "u32",
+            "LocalInvocationIndex",
+            "%1:ptr<__in, u32, read> = var @builtin(local_invocation_index)",
+        },
+        BuiltinCase{
+            "vec3u",
+            "NumWorkgroups",
+            "%1:ptr<__in, vec3<u32>, read> = var @builtin(num_workgroups)",
+        },
+        BuiltinCase{
+            "u32",
+            "SampleId",
+            "%1:ptr<__in, u32, read> = var @builtin(sample_index)",
+        },
+        BuiltinCase{
+            "arr_u32_1",
+            "SampleMask",
+            "%1:ptr<__in, array<u32, 1>, read> = var @builtin(sample_mask)",
+        },
+        BuiltinCase{
+            "u32",
+            "VertexIndex",
+            "%1:ptr<__in, u32, read> = var @builtin(vertex_index)",
+        },
+        BuiltinCase{
+            "vec3u",
+            "WorkgroupId",
+            "%1:ptr<__in, vec3<u32>, read> = var @builtin(workgroup_id)",
+        }),
+    PrintBuiltinCase);
+
+using BuiltinOutputTest = SpirvParserTestWithParam<BuiltinCase>;
+
+TEST_P(BuiltinOutputTest, Enum) {
+    auto params = GetParam();
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpDecorate %var BuiltIn )" +
+                  params.spirv_builtin + R"(
+       %void = OpTypeVoid
+        %u32 = OpTypeInt 32 0
+        %f32 = OpTypeFloat 32
+      %vec4f = OpTypeVector %f32 4
+      %u32_1 = OpConstant %u32 1
+  %arr_u32_1 = OpTypeArray %u32 %u32_1
+    %fn_type = OpTypeFunction %void
+
+%_ptr_Output = OpTypePointer Output %)" +
+                  params.spirv_type + R"(
+        %var = OpVariable %_ptr_Output Output
+
+       %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+)",
+              params.ir);
+}
+
+INSTANTIATE_TEST_SUITE_P(
+    SpirvParser,
+    BuiltinOutputTest,
+    testing::Values(
+        BuiltinCase{
+            "f32",
+            "PointSize",
+            "%1:ptr<__out, f32, read_write> = var @builtin(__point_size)",
+        },
+        BuiltinCase{
+            "vec4f",
+            "Position",
+            "%1:ptr<__out, vec4<f32>, read_write> = var @builtin(position)",
+        },
+        BuiltinCase{
+            "arr_u32_1",
+            "SampleMask",
+            "%1:ptr<__out, array<u32, 1>, read_write> = var @builtin(sample_mask)",
+        }),
+    PrintBuiltinCase);
+
+TEST_F(SpirvParserTest, Invariant_OnVariable) {
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpDecorate %var BuiltIn Position
+               OpDecorate %var Invariant
+       %void = OpTypeVoid
+        %f32 = OpTypeFloat 32
+      %vec4f = OpTypeVector %f32 4
+    %fn_type = OpTypeFunction %void
+
+%_ptr_Output = OpTypePointer Output %vec4f
+        %var = OpVariable %_ptr_Output Output
+
+       %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+)",
+              "%1:ptr<__out, vec4<f32>, read_write> = var @invariant @builtin(position)");
+}
+
+struct LocationCase {
+    std::string spirv_decorations;
+    std::string ir;
+};
+
+using LocationVarTest = SpirvParserTestWithParam<LocationCase>;
+
+TEST_P(LocationVarTest, Input) {
+    auto params = GetParam();
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpCapability SampleRateShading
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+          )" + params.spirv_decorations +
+                  R"(
+       %void = OpTypeVoid
+        %f32 = OpTypeFloat 32
+      %vec4f = OpTypeVector %f32 4
+    %fn_type = OpTypeFunction %void
+
+%_ptr_Input = OpTypePointer Input %vec4f
+        %var = OpVariable %_ptr_Input Input
+
+       %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+)",
+              "%1:ptr<__in, vec4<f32>, read> = " + params.ir);
+}
+
+TEST_P(LocationVarTest, Output) {
+    auto params = GetParam();
+    EXPECT_IR(R"(
+               OpCapability Shader
+               OpCapability SampleRateShading
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+          )" + params.spirv_decorations +
+                  R"(
+       %void = OpTypeVoid
+        %f32 = OpTypeFloat 32
+      %vec4f = OpTypeVector %f32 4
+    %fn_type = OpTypeFunction %void
+
+%_ptr_Output = OpTypePointer Output %vec4f
+        %var = OpVariable %_ptr_Output Output
+
+       %main = OpFunction %void None %fn_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+)",
+              "%1:ptr<__out, vec4<f32>, read_write> = " + params.ir);
+}
+
+INSTANTIATE_TEST_SUITE_P(SpirvParser,
+                         LocationVarTest,
+                         testing::Values(
+                             LocationCase{
+                                 "OpDecorate %var Location 1 ",
+                                 "var @location(1)",
+                             },
+                             LocationCase{
+                                 "OpDecorate %var Location 2 "
+                                 "OpDecorate %var NoPerspective ",
+                                 "var @location(2) @interpolate(linear, center)",
+                             },
+                             LocationCase{
+                                 "OpDecorate %var Location 3 "
+                                 "OpDecorate %var Flat ",
+                                 "var @location(3) @interpolate(flat, center)",
+                             },
+                             LocationCase{
+                                 "OpDecorate %var Location 4 "
+                                 "OpDecorate %var Centroid ",
+                                 "var @location(4) @interpolate(perspective, centroid)",
+                             },
+                             LocationCase{
+                                 "OpDecorate %var Location 5 "
+                                 "OpDecorate %var Sample ",
+                                 "var @location(5) @interpolate(perspective, sample)",
+                             },
+                             LocationCase{
+                                 "OpDecorate %var Location 6 "
+                                 "OpDecorate %var NoPerspective "
+                                 "OpDecorate %var Centroid ",
+                                 "var @location(6) @interpolate(linear, centroid)",
+                             }));
+
 }  // namespace
 }  // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/reader_test.cc b/src/tint/lang/spirv/reader/reader_test.cc
index 1a086a8..5313aaf 100644
--- a/src/tint/lang/spirv/reader/reader_test.cc
+++ b/src/tint/lang/spirv/reader/reader_test.cc
@@ -68,6 +68,25 @@
     }
 };
 
+TEST_F(SpirvReaderTest, UnsupportedExtension) {
+    auto got = Run(R"(
+               OpCapability Shader
+               OpExtension "SPV_KHR_variable_pointers"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+       %void = OpTypeVoid
+    %ep_type = OpTypeFunction %void
+       %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+               OpReturn
+               OpFunctionEnd
+)");
+    ASSERT_NE(got, Success);
+    EXPECT_EQ(got.Failure().reason.str(),
+              "error: SPIR-V extension 'SPV_KHR_variable_pointers' is not supported");
+}
+
 TEST_F(SpirvReaderTest, Load_VectorComponent) {
     auto got = Run(R"(
                OpCapability Shader
diff --git a/src/tint/lang/wgsl/inspector/inspector.cc b/src/tint/lang/wgsl/inspector/inspector.cc
index 630ef0f..d7fe564 100644
--- a/src/tint/lang/wgsl/inspector/inspector.cc
+++ b/src/tint/lang/wgsl/inspector/inspector.cc
@@ -27,7 +27,7 @@
 
 #include "src/tint/lang/wgsl/inspector/inspector.h"
 
-#include <limits>
+#include <unordered_set>
 #include <utility>
 
 #include "src/tint/lang/core/builtin_value.h"
@@ -1009,4 +1009,139 @@
     }
 }
 
+std::vector<Inspector::LevelSampleInfo> Inspector::GetTextureQueries(const std::string& ep_name) {
+    std::vector<LevelSampleInfo> res;
+
+    std::unordered_set<BindingPoint> seen = {};
+
+    auto sample_type_for_call_and_type = [](wgsl::BuiltinFn builtin, const core::type::Type* ty) {
+        if (builtin == wgsl::BuiltinFn::kTextureNumLevels) {
+            return TextureQueryType::kTextureNumLevels;
+        }
+        if (builtin == wgsl::BuiltinFn::kTextureLoad) {
+            if (!ty->UnwrapRef()
+                     ->IsAnyOf<core::type::MultisampledTexture,
+                               core::type::DepthMultisampledTexture>()) {
+                return TextureQueryType::kTextureNumLevels;
+            }
+        }
+
+        return TextureQueryType::kTextureNumSamples;
+    };
+
+    Hashmap<const sem::Function*, Hashmap<const ast::Parameter*, TextureQueryType, 4>, 8>
+        fn_to_data;
+
+    auto record_function_param = [&fn_to_data](const sem::Function* func,
+                                               const ast::Parameter* param, TextureQueryType type) {
+        auto& param_to_type = *fn_to_data.GetOrZero(func);
+
+        auto entry = param_to_type.Get(param);
+        if (entry.has_value()) {
+            return;
+        }
+
+        param_to_type.Add(param, type);
+    };
+
+    auto save_if_needed = [&res, &seen](const sem::GlobalVariable* global, TextureQueryType type) {
+        auto binding = global->Attributes().binding_point.value();
+        if (seen.insert(binding).second) {
+            res.emplace_back(LevelSampleInfo{type, binding.group, binding.binding});
+        }
+    };
+
+    auto& sem = program_.Sem();
+
+    const auto* ep = FindEntryPointByName(ep_name);
+    if (!ep) {
+        return {};
+    }
+
+    // This works in dependency order such that we'll see the texture call first and can record
+    // any function parameter information and then as we walk up the function chain we can look
+    // the call data.
+    for (auto* fn_decl : sem.Module()->DependencyOrderedDeclarations()) {
+        auto* fn = sem.Get<sem::Function>(fn_decl);
+        if (!fn) {
+            continue;
+        }
+
+        // This is an entrypoint, make sure it's the requested entry point
+        if (fn->Declaration()->IsEntryPoint()) {
+            if (fn->Declaration() != ep) {
+                continue;
+            }
+        } else {
+            // Not an entry point, make sure it was called from the requested entry point
+            if (!fn->HasAncestorEntryPoint(ep->name->symbol)) {
+                continue;
+            }
+        }
+
+        for (auto* call : fn->DirectCalls()) {
+            // Builtin function call, record the texture information. If the used texture maps
+            // back up to a function parameter just store the type of the call and we'll track the
+            // function callback up in the `sem::Function` branch.
+            tint::Switch(
+                call->Target(),
+                [&](const sem::BuiltinFn* builtin) {
+                    if (builtin->Fn() != wgsl::BuiltinFn::kTextureNumLevels &&
+                        builtin->Fn() != wgsl::BuiltinFn::kTextureNumSamples &&
+                        builtin->Fn() != wgsl::BuiltinFn::kTextureLoad) {
+                        return;
+                    }
+
+                    auto* texture_expr = call->Declaration()->args[0];
+                    auto* texture_sem = sem.GetVal(texture_expr)->RootIdentifier();
+                    TINT_ASSERT(texture_sem);
+
+                    auto type = sample_type_for_call_and_type(builtin->Fn(), texture_sem->Type());
+
+                    tint::Switch(
+                        texture_sem,  //
+                        [&](const sem::GlobalVariable* global) { save_if_needed(global, type); },
+                        [&](const sem::Parameter* param) {
+                            record_function_param(fn, param->Declaration(), type);
+                        },
+                        TINT_ICE_ON_NO_MATCH);
+                },
+                [&](const sem::Function* func) {
+                    // A function call, check to see if any params needed to be tracked back to a
+                    // global texture.
+
+                    auto param_to_type = fn_to_data.Find(func);
+                    if (!param_to_type) {
+                        return;
+                    }
+                    TINT_ASSERT(call->Arguments().Length() == func->Declaration()->params.Length());
+
+                    for (size_t i = 0; i < call->Arguments().Length(); i++) {
+                        auto param = func->Declaration()->params[i];
+
+                        // Determine if this had a texture we cared about
+                        auto type = param_to_type->Get(param);
+                        if (!type.has_value()) {
+                            continue;
+                        }
+
+                        auto* arg = call->Arguments()[i];
+                        auto* texture_sem = arg->RootIdentifier();
+
+                        tint::Switch(
+                            texture_sem,
+                            [&](const sem::GlobalVariable* global) {
+                                save_if_needed(global, type.value());
+                            },
+                            [&](const sem::Parameter* p) {
+                                record_function_param(fn, p->Declaration(), type.value());
+                            },
+                            TINT_ICE_ON_NO_MATCH);
+                    }
+                });
+        }
+    }
+    return res;
+}
+
 }  // namespace tint::inspector
diff --git a/src/tint/lang/wgsl/inspector/inspector.h b/src/tint/lang/wgsl/inspector/inspector.h
index 14c5ef6..6604832 100644
--- a/src/tint/lang/wgsl/inspector/inspector.h
+++ b/src/tint/lang/wgsl/inspector/inspector.h
@@ -156,6 +156,31 @@
     /// extension.
     std::vector<std::pair<std::string, Source>> GetEnableDirectives();
 
+    /// The information needed to be supplied.
+    enum class TextureQueryType : uint8_t {
+        /// Texture Num Levels
+        kTextureNumLevels,
+        /// Texture Num Samples
+        kTextureNumSamples,
+    };
+    /// Information on level and sample calls by a given texture binding point
+    struct LevelSampleInfo {
+        /// The type of function
+        TextureQueryType type = TextureQueryType::kTextureNumLevels;
+        /// The group number
+        uint32_t group = 0;
+        /// The binding number
+        uint32_t binding = 0;
+    };
+
+    /// @param ep the entry point ot get the information for
+    /// @returns a vector of information for textures which call textureNumLevels and
+    /// textureNumSamples for backends which require additional support for those methods. Each
+    /// binding point will only be returned once regardless of the number of calls made. The
+    /// texture types for `textureNumSamples` is disjoint from the texture types in
+    /// `textureNumLevels` so the binding point will always be one or the other.
+    std::vector<LevelSampleInfo> GetTextureQueries(const std::string& ep);
+
   private:
     const Program& program_;
     diag::List diagnostics_;
diff --git a/src/tint/lang/wgsl/inspector/inspector_test.cc b/src/tint/lang/wgsl/inspector/inspector_test.cc
index c33f950..76d514e 100644
--- a/src/tint/lang/wgsl/inspector/inspector_test.cc
+++ b/src/tint/lang/wgsl/inspector/inspector_test.cc
@@ -3701,5 +3701,286 @@
     inspector.GetSamplerTextureUses("main");
 }
 
+class InspectorTextureTest : public InspectorRunner, public testing::Test {};
+
+TEST_F(InspectorTextureTest, TextureLevelInEP) {
+    std::string shader = R"(
+@group(2) @binding(3) var myTexture: texture_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+  let num = textureNumLevels(myTexture);
+})";
+
+    Inspector& inspector = Initialize(shader);
+    auto info = inspector.GetTextureQueries("main");
+
+    ASSERT_EQ(1u, info.size());
+    EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info[0].type);
+    EXPECT_EQ(2u, info[0].group);
+    EXPECT_EQ(3u, info[0].binding);
+}
+
+TEST_F(InspectorTextureTest, TextureLevelInEPNoDups) {
+    std::string shader = R"(
+@group(0) @binding(0) var myTexture: texture_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+  let num1 = textureNumLevels(myTexture);
+  let num2 = textureNumLevels(myTexture);
+})";
+
+    Inspector& inspector = Initialize(shader);
+    auto info = inspector.GetTextureQueries("main");
+
+    ASSERT_EQ(1u, info.size());
+}
+
+TEST_F(InspectorTextureTest, TextureLevelInEPMultiple) {
+    std::string shader = R"(
+@group(2) @binding(3) var tex1: texture_2d<f32>;
+@group(1) @binding(2) var tex2: texture_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+  let num1 = textureNumLevels(tex1);
+  let num2 = textureNumLevels(tex2);
+})";
+
+    Inspector& inspector = Initialize(shader);
+    auto info = inspector.GetTextureQueries("main");
+
+    ASSERT_EQ(2u, info.size());
+
+    EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info[0].type);
+    EXPECT_EQ(2u, info[0].group);
+    EXPECT_EQ(3u, info[0].binding);
+
+    EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info[1].type);
+    EXPECT_EQ(1u, info[1].group);
+    EXPECT_EQ(2u, info[1].binding);
+}
+
+TEST_F(InspectorTextureTest, TextureSamplesInEP) {
+    std::string shader = R"(
+@group(2) @binding(3) var myTexture: texture_multisampled_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+  let num = textureNumSamples(myTexture);
+})";
+
+    Inspector& inspector = Initialize(shader);
+    auto info = inspector.GetTextureQueries("main");
+
+    ASSERT_EQ(1u, info.size());
+    EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info[0].type);
+    EXPECT_EQ(2u, info[0].group);
+    EXPECT_EQ(3u, info[0].binding);
+}
+
+TEST_F(InspectorTextureTest, TextureSamplesInEPNoDups) {
+    std::string shader = R"(
+@group(0) @binding(0) var myTexture: texture_multisampled_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+  let num1 = textureNumSamples(myTexture);
+  let num2 = textureNumSamples(myTexture);
+})";
+
+    Inspector& inspector = Initialize(shader);
+    auto info = inspector.GetTextureQueries("main");
+
+    ASSERT_EQ(1u, info.size());
+}
+
+TEST_F(InspectorTextureTest, TextureSamplesInEPMultiple) {
+    std::string shader = R"(
+@group(2) @binding(3) var tex1: texture_multisampled_2d<f32>;
+@group(1) @binding(2) var tex2: texture_multisampled_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+  let num1 = textureNumSamples(tex1);
+  let num2 = textureNumSamples(tex2);
+})";
+
+    Inspector& inspector = Initialize(shader);
+    auto info = inspector.GetTextureQueries("main");
+
+    ASSERT_EQ(2u, info.size());
+
+    EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info[0].type);
+    EXPECT_EQ(2u, info[0].group);
+    EXPECT_EQ(3u, info[0].binding);
+
+    EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info[1].type);
+    EXPECT_EQ(1u, info[1].group);
+    EXPECT_EQ(2u, info[1].binding);
+}
+
+TEST_F(InspectorTextureTest, TextureLoadInEP) {
+    std::string shader = R"(
+@group(2) @binding(3) var tex1: texture_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+  let num1 = textureLoad(tex1, vec2(0, 0), 0);
+})";
+
+    Inspector& inspector = Initialize(shader);
+    auto info = inspector.GetTextureQueries("main");
+
+    ASSERT_EQ(1u, info.size());
+
+    EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info[0].type);
+    EXPECT_EQ(2u, info[0].group);
+    EXPECT_EQ(3u, info[0].binding);
+}
+
+TEST_F(InspectorTextureTest, TextureLoadMultisampledInEP) {
+    std::string shader = R"(
+@group(2) @binding(3) var tex1: texture_multisampled_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+  let num1 = textureLoad(tex1, vec2(0, 0), 0);
+})";
+
+    Inspector& inspector = Initialize(shader);
+    auto info = inspector.GetTextureQueries("main");
+
+    ASSERT_EQ(1u, info.size());
+
+    EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info[0].type);
+    EXPECT_EQ(2u, info[0].group);
+    EXPECT_EQ(3u, info[0].binding);
+}
+
+TEST_F(InspectorTextureTest, TextureLoadMultipleInEP) {
+    std::string shader = R"(
+@group(2) @binding(3) var tex1: texture_2d<f32>;
+@group(1) @binding(4) var tex2: texture_multisampled_2d<f32>;
+
+@compute @workgroup_size(1)
+fn main() {
+  let num1 = textureLoad(tex1, vec2(0, 0), 0);
+  let num2 = textureLoad(tex2, vec2(0, 0), 0);
+})";
+
+    Inspector& inspector = Initialize(shader);
+    auto info = inspector.GetTextureQueries("main");
+
+    ASSERT_EQ(2u, info.size());
+
+    EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info[0].type);
+    EXPECT_EQ(2u, info[0].group);
+    EXPECT_EQ(3u, info[0].binding);
+    EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info[1].type);
+    EXPECT_EQ(1u, info[1].group);
+    EXPECT_EQ(4u, info[1].binding);
+}
+
+TEST_F(InspectorTextureTest, TextureInSubfunction) {
+    std::string shader = R"(
+@group(2) @binding(3) var tex1: texture_2d<f32>;
+@group(1) @binding(4) var tex2: texture_multisampled_2d<f32>;
+@group(1) @binding(3) var tex3: texture_2d<f32>;
+
+fn b(tx1: texture_2d<f32>, tx2: texture_multisampled_2d<f32>, tx3: texture_2d<f32>, tx4: texture_2d<f32>) {
+  let v1 = textureNumLevels(tx1);
+  let v2 = textureNumSamples(tx2);
+  let v3 = textureLoad(tx3, vec2(0, 0), 0);
+  let v4 = textureNumLevels(tx4);
+}
+
+fn a(tx1: texture_2d<f32>, tx2: texture_multisampled_2d<f32>, tx3: texture_2d<f32>) {
+  b(tx1, tx2, tx3, tx1);
+}
+
+@compute @workgroup_size(1)
+fn main() {
+  a(tex1, tex2, tex3);
+})";
+
+    Inspector& inspector = Initialize(shader);
+    auto info = inspector.GetTextureQueries("main");
+
+    ASSERT_EQ(3u, info.size());
+
+    EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info[0].type);
+    EXPECT_EQ(2u, info[0].group);
+    EXPECT_EQ(3u, info[0].binding);
+    EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info[1].type);
+    EXPECT_EQ(1u, info[1].group);
+    EXPECT_EQ(4u, info[1].binding);
+    EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info[2].type);
+    EXPECT_EQ(1u, info[2].group);
+    EXPECT_EQ(3u, info[2].binding);
+}
+
+TEST_F(InspectorTextureTest, TextureMultipleEPs) {
+    std::string shader = R"(
+@group(0) @binding(0) var<storage, read_write> dstBuf : array<u32>;
+@group(0) @binding(1) var tex1 : texture_2d_array<f32>;
+@group(0) @binding(4) var tex2 : texture_multisampled_2d<f32>;
+@group(1) @binding(3) var tex3 : texture_2d_array<f32>;
+
+@compute @workgroup_size(1, 1, 1) fn main1() {
+    dstBuf[0] = textureNumLayers(tex1);
+    dstBuf[1] = textureNumLevels(tex1);
+    dstBuf[2] = textureNumSamples(tex2);
+    dstBuf[3] = textureNumLevels(tex3);
+}
+
+@compute @workgroup_size(1, 1, 1) fn main2() {
+    dstBuf[0] = textureNumLayers(tex1);
+    dstBuf[1] = textureNumLevels(tex1);
+    dstBuf[2] = textureNumSamples(tex2);
+}
+    )";
+    Inspector& inspector = Initialize(shader);
+    {
+        auto info1 = inspector.GetTextureQueries("main1");
+        ASSERT_EQ(3u, info1.size());
+
+        EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info1[0].type);
+        EXPECT_EQ(0u, info1[0].group);
+        EXPECT_EQ(1u, info1[0].binding);
+        EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info1[1].type);
+        EXPECT_EQ(0u, info1[1].group);
+        EXPECT_EQ(4u, info1[1].binding);
+        EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info1[2].type);
+        EXPECT_EQ(1u, info1[2].group);
+        EXPECT_EQ(3u, info1[2].binding);
+    }
+    {
+        auto info2 = inspector.GetTextureQueries("main2");
+        ASSERT_EQ(2u, info2.size());
+
+        EXPECT_EQ(Inspector::TextureQueryType::kTextureNumLevels, info2[0].type);
+        EXPECT_EQ(0u, info2[0].group);
+        EXPECT_EQ(1u, info2[0].binding);
+        EXPECT_EQ(Inspector::TextureQueryType::kTextureNumSamples, info2[1].type);
+        EXPECT_EQ(0u, info2[1].group);
+        EXPECT_EQ(4u, info2[1].binding);
+    }
+}
+
 }  // namespace
+
+static std::ostream& operator<<(std::ostream& out, const Inspector::TextureQueryType& ty) {
+    switch (ty) {
+        case Inspector::TextureQueryType::kTextureNumLevels:
+            out << "textureNumLevels";
+            break;
+        case Inspector::TextureQueryType::kTextureNumSamples:
+            out << "textureNumSamples";
+            break;
+    }
+    return out;
+}
+
 }  // namespace tint::inspector