diff --git a/src/transform/decompose_memory_access.cc b/src/transform/decompose_memory_access.cc
index 06251de..f06b3ae 100644
--- a/src/transform/decompose_memory_access.cc
+++ b/src/transform/decompose_memory_access.cc
@@ -303,6 +303,10 @@
 
 /// State holds the current transform state
 struct DecomposeMemoryAccess::State {
+  /// The clone context
+  CloneContext& ctx;
+  /// Alias to `*ctx.dst`
+  ProgramBuilder& b;
   /// Map of AST expression to storage or uniform buffer access
   /// This map has entries added when encountered, and removed when outer
   /// expressions chain the access.
@@ -322,6 +326,10 @@
   /// Allocations for offsets
   BlockAllocator<Offset> offsets_;
 
+  /// Constructor
+  /// @param context the CloneContext
+  explicit State(CloneContext& context) : ctx(context), b(*ctx.dst) {}
+
   /// @param offset the offset value to wrap in an Offset
   /// @returns an Offset for the given literal value
   const Offset* ToOffset(uint32_t offset) {
@@ -440,13 +448,11 @@
   /// of type `el_ty` from a storage or uniform buffer of type `buf_ty`.
   /// The emitted function has the signature:
   ///   `fn load(buf : buf_ty, offset : u32) -> el_ty`
-  /// @param ctx the CloneContext
   /// @param buf_ty the storage or uniform buffer type
   /// @param el_ty the storage or uniform buffer element type
   /// @param var_user the variable user
   /// @return the name of the function that performs the load
-  Symbol LoadFunc(CloneContext& ctx,
-                  const sem::Type* buf_ty,
+  Symbol LoadFunc(const sem::Type* buf_ty,
                   const sem::Type* el_ty,
                   const sem::VariableUser* var_user) {
     auto storage_class = var_user->Variable()->StorageClass();
@@ -454,70 +460,89 @@
         load_funcs, LoadStoreKey{storage_class, buf_ty, el_ty}, [&] {
           auto* buf_ast_ty = CreateASTTypeFor(ctx, buf_ty);
           auto* disable_validation =
-              ctx.dst->ASTNodes().Create<ast::DisableValidationDecoration>(
-                  ctx.dst->ID(), ast::DisabledValidation::
-                                     kIgnoreConstructibleFunctionParameter);
+              b.ASTNodes().Create<ast::DisableValidationDecoration>(
+                  b.ID(), ast::DisabledValidation::
+                              kIgnoreConstructibleFunctionParameter);
 
           ast::VariableList params = {
               // Note: The buffer parameter requires the StorageClass in
               // order for HLSL to emit this as a ByteAddressBuffer or cbuffer
               // array.
-              ctx.dst->create<ast::Variable>(
-                  ctx.dst->Sym("buffer"), storage_class,
-                  var_user->Variable()->Access(), buf_ast_ty, true, nullptr,
-                  ast::DecorationList{disable_validation}),
-              ctx.dst->Param("offset", ctx.dst->ty.u32()),
+              b.create<ast::Variable>(b.Sym("buffer"), storage_class,
+                                      var_user->Variable()->Access(),
+                                      buf_ast_ty, true, nullptr,
+                                      ast::DecorationList{disable_validation}),
+              b.Param("offset", b.ty.u32()),
           };
 
-          ast::Function* func = nullptr;
+          auto name = b.Sym();
+
           if (auto* intrinsic =
                   IntrinsicLoadFor(ctx.dst, storage_class, el_ty)) {
             auto* el_ast_ty = CreateASTTypeFor(ctx, el_ty);
-            func = ctx.dst->create<ast::Function>(
-                ctx.dst->Sym(), params, el_ast_ty, nullptr,
+            auto* func = b.create<ast::Function>(
+                name, params, el_ast_ty, nullptr,
                 ast::DecorationList{
                     intrinsic,
-                    ctx.dst->ASTNodes()
-                        .Create<ast::DisableValidationDecoration>(
-                            ctx.dst->ID(),
-                            ast::DisabledValidation::kFunctionHasNoBody),
+                    b.ASTNodes().Create<ast::DisableValidationDecoration>(
+                        b.ID(), ast::DisabledValidation::kFunctionHasNoBody),
                 },
                 ast::DecorationList{});
+            b.AST().AddFunction(func);
+          } else if (auto* arr_ty = el_ty->As<sem::Array>()) {
+            // fn load_func(buf : buf_ty, offset : u32) -> array<T, N> {
+            //   var arr : array<T, N>;
+            //   for (var i = 0u; i < array_count; i = i + 1) {
+            //     arr[i] = el_load_func(buf, offset + i * array_stride)
+            //   }
+            //   return arr;
+            // }
+            auto load =
+                LoadFunc(buf_ty, arr_ty->ElemType()->UnwrapRef(), var_user);
+            auto* arr =
+                b.Var(b.Symbols().New("arr"), CreateASTTypeFor(ctx, arr_ty));
+            auto* i = b.Var(b.Symbols().New("i"), nullptr, b.Expr(0u));
+            auto* for_init = b.Decl(i);
+            auto* for_cond = b.create<ast::BinaryExpression>(
+                ast::BinaryOp::kLessThan, b.Expr(i), b.Expr(arr_ty->Count()));
+            auto* for_cont = b.Assign(i, b.Add(i, 1u));
+            auto* arr_el = b.IndexAccessor(arr, i);
+            auto* el_offset =
+                b.Add(b.Expr("offset"), b.Mul(i, arr_ty->Stride()));
+            auto* el_val = b.Call(load, "buffer", el_offset);
+            auto* for_loop = b.For(for_init, for_cond, for_cont,
+                                   b.Block(b.Assign(arr_el, el_val)));
+
+            b.Func(name, params, CreateASTTypeFor(ctx, arr_ty),
+                   {
+                       b.Decl(arr),
+                       for_loop,
+                       b.Return(arr),
+                   });
           } else {
             ast::ExpressionList values;
             if (auto* mat_ty = el_ty->As<sem::Matrix>()) {
               auto* vec_ty = mat_ty->ColumnType();
-              Symbol load = LoadFunc(ctx, buf_ty, vec_ty, var_user);
+              Symbol load = LoadFunc(buf_ty, vec_ty, var_user);
               for (uint32_t i = 0; i < mat_ty->columns(); i++) {
-                auto* offset =
-                    ctx.dst->Add("offset", i * MatrixColumnStride(mat_ty));
-                values.emplace_back(ctx.dst->Call(load, "buffer", offset));
+                auto* offset = b.Add("offset", i * MatrixColumnStride(mat_ty));
+                values.emplace_back(b.Call(load, "buffer", offset));
               }
             } else if (auto* str = el_ty->As<sem::Struct>()) {
               for (auto* member : str->Members()) {
-                auto* offset = ctx.dst->Add("offset", member->Offset());
-                Symbol load = LoadFunc(ctx, buf_ty, member->Type()->UnwrapRef(),
-                                       var_user);
-                values.emplace_back(ctx.dst->Call(load, "buffer", offset));
-              }
-            } else if (auto* arr = el_ty->As<sem::Array>()) {
-              for (uint32_t i = 0; i < arr->Count(); i++) {
-                auto* offset = ctx.dst->Add("offset", arr->Stride() * i);
-                Symbol load = LoadFunc(ctx, buf_ty,
-                                       arr->ElemType()->UnwrapRef(), var_user);
-                values.emplace_back(ctx.dst->Call(load, "buffer", offset));
+                auto* offset = b.Add("offset", member->Offset());
+                Symbol load =
+                    LoadFunc(buf_ty, member->Type()->UnwrapRef(), var_user);
+                values.emplace_back(b.Call(load, "buffer", offset));
               }
             }
-            auto* el_ast_ty = CreateASTTypeFor(ctx, el_ty);
-            func = ctx.dst->create<ast::Function>(
-                ctx.dst->Sym(), params, el_ast_ty,
-                ctx.dst->Block(ctx.dst->Return(
-                    ctx.dst->create<ast::TypeConstructorExpression>(
-                        CreateASTTypeFor(ctx, el_ty), values))),
-                ast::DecorationList{}, ast::DecorationList{});
+            b.Func(name, params, CreateASTTypeFor(ctx, el_ty),
+                   {
+                       b.Return(b.create<ast::TypeConstructorExpression>(
+                           CreateASTTypeFor(ctx, el_ty), values)),
+                   });
           }
-          ctx.dst->AST().AddFunction(func);
-          return func->symbol();
+          return name;
         });
   }
 
@@ -525,13 +550,11 @@
   /// element of type `el_ty` to a storage buffer of type `buf_ty`.
   /// The function has the signature:
   ///   `fn store(buf : buf_ty, offset : u32, value : el_ty)`
-  /// @param ctx the CloneContext
   /// @param buf_ty the storage buffer type
   /// @param el_ty the storage buffer element type
   /// @param var_user the variable user
   /// @return the name of the function that performs the store
-  Symbol StoreFunc(CloneContext& ctx,
-                   const sem::Type* buf_ty,
+  Symbol StoreFunc(const sem::Type* buf_ty,
                    const sem::Type* el_ty,
                    const sem::VariableUser* var_user) {
     auto storage_class = var_user->Variable()->StorageClass();
@@ -540,75 +563,87 @@
           auto* buf_ast_ty = CreateASTTypeFor(ctx, buf_ty);
           auto* el_ast_ty = CreateASTTypeFor(ctx, el_ty);
           auto* disable_validation =
-              ctx.dst->ASTNodes().Create<ast::DisableValidationDecoration>(
-                  ctx.dst->ID(), ast::DisabledValidation::
-                                     kIgnoreConstructibleFunctionParameter);
+              b.ASTNodes().Create<ast::DisableValidationDecoration>(
+                  b.ID(), ast::DisabledValidation::
+                              kIgnoreConstructibleFunctionParameter);
           ast::VariableList params{
               // Note: The buffer parameter requires the StorageClass in
               // order for HLSL to emit this as a ByteAddressBuffer.
 
-              ctx.dst->create<ast::Variable>(
-                  ctx.dst->Sym("buffer"), storage_class,
-                  var_user->Variable()->Access(), buf_ast_ty, true, nullptr,
-                  ast::DecorationList{disable_validation}),
-              ctx.dst->Param("offset", ctx.dst->ty.u32()),
-              ctx.dst->Param("value", el_ast_ty),
+              b.create<ast::Variable>(b.Sym("buffer"), storage_class,
+                                      var_user->Variable()->Access(),
+                                      buf_ast_ty, true, nullptr,
+                                      ast::DecorationList{disable_validation}),
+              b.Param("offset", b.ty.u32()),
+              b.Param("value", el_ast_ty),
           };
-          ast::Function* func = nullptr;
+
+          auto name = b.Sym();
+
           if (auto* intrinsic =
                   IntrinsicStoreFor(ctx.dst, storage_class, el_ty)) {
-            func = ctx.dst->create<ast::Function>(
-                ctx.dst->Sym(), params, ctx.dst->ty.void_(), nullptr,
+            auto* func = b.create<ast::Function>(
+                name, params, b.ty.void_(), nullptr,
                 ast::DecorationList{
                     intrinsic,
-                    ctx.dst->ASTNodes()
-                        .Create<ast::DisableValidationDecoration>(
-                            ctx.dst->ID(),
-                            ast::DisabledValidation::kFunctionHasNoBody),
+                    b.ASTNodes().Create<ast::DisableValidationDecoration>(
+                        b.ID(), ast::DisabledValidation::kFunctionHasNoBody),
                 },
                 ast::DecorationList{});
-
+            b.AST().AddFunction(func);
           } else {
             ast::StatementList body;
-            if (auto* mat_ty = el_ty->As<sem::Matrix>()) {
+            if (auto* arr_ty = el_ty->As<sem::Array>()) {
+              // fn store_func(buf : buf_ty, offset : u32, value : el_ty) {
+              //   var array = value; // No dynamic indexing on constant arrays
+              //   for (var i = 0u; i < array_count; i = i + 1) {
+              //     arr[i] = el_store_func(buf, offset + i * array_stride,
+              //                            value[i])
+              //   }
+              //   return arr;
+              // }
+              auto* array =
+                  b.Var(b.Symbols().New("array"), nullptr, b.Expr("value"));
+              auto store =
+                  StoreFunc(buf_ty, arr_ty->ElemType()->UnwrapRef(), var_user);
+              auto* i = b.Var(b.Symbols().New("i"), nullptr, b.Expr(0u));
+              auto* for_init = b.Decl(i);
+              auto* for_cond = b.create<ast::BinaryExpression>(
+                  ast::BinaryOp::kLessThan, b.Expr(i), b.Expr(arr_ty->Count()));
+              auto* for_cont = b.Assign(i, b.Add(i, 1u));
+              auto* arr_el = b.IndexAccessor(array, i);
+              auto* el_offset =
+                  b.Add(b.Expr("offset"), b.Mul(i, arr_ty->Stride()));
+              auto* store_stmt = b.create<ast::CallStatement>(
+                  b.Call(store, "buffer", el_offset, arr_el));
+              auto* for_loop =
+                  b.For(for_init, for_cond, for_cont, b.Block(store_stmt));
+
+              body = {b.Decl(array), for_loop};
+            } else if (auto* mat_ty = el_ty->As<sem::Matrix>()) {
               auto* vec_ty = mat_ty->ColumnType();
-              Symbol store = StoreFunc(ctx, buf_ty, vec_ty, var_user);
+              Symbol store = StoreFunc(buf_ty, vec_ty, var_user);
               for (uint32_t i = 0; i < mat_ty->columns(); i++) {
-                auto* offset =
-                    ctx.dst->Add("offset", i * MatrixColumnStride(mat_ty));
-                auto* access = ctx.dst->IndexAccessor("value", i);
-                auto* call = ctx.dst->Call(store, "buffer", offset, access);
-                body.emplace_back(ctx.dst->create<ast::CallStatement>(call));
+                auto* offset = b.Add("offset", i * MatrixColumnStride(mat_ty));
+                auto* access = b.IndexAccessor("value", i);
+                auto* call = b.Call(store, "buffer", offset, access);
+                body.emplace_back(b.create<ast::CallStatement>(call));
               }
             } else if (auto* str = el_ty->As<sem::Struct>()) {
               for (auto* member : str->Members()) {
-                auto* offset = ctx.dst->Add("offset", member->Offset());
-                auto* access = ctx.dst->MemberAccessor(
+                auto* offset = b.Add("offset", member->Offset());
+                auto* access = b.MemberAccessor(
                     "value", ctx.Clone(member->Declaration()->symbol()));
-                Symbol store = StoreFunc(ctx, buf_ty,
-                                         member->Type()->UnwrapRef(), var_user);
-                auto* call = ctx.dst->Call(store, "buffer", offset, access);
-                body.emplace_back(ctx.dst->create<ast::CallStatement>(call));
-              }
-            } else if (auto* arr = el_ty->As<sem::Array>()) {
-              for (uint32_t i = 0; i < arr->Count(); i++) {
-                auto* offset = ctx.dst->Add("offset", arr->Stride() * i);
-                auto* access =
-                    ctx.dst->IndexAccessor("value", ctx.dst->Expr(i));
-                Symbol store = StoreFunc(
-                    ctx, buf_ty, arr->ElemType()->UnwrapRef(), var_user);
-                auto* call = ctx.dst->Call(store, "buffer", offset, access);
-                body.emplace_back(ctx.dst->create<ast::CallStatement>(call));
+                Symbol store =
+                    StoreFunc(buf_ty, member->Type()->UnwrapRef(), var_user);
+                auto* call = b.Call(store, "buffer", offset, access);
+                body.emplace_back(b.create<ast::CallStatement>(call));
               }
             }
-            func = ctx.dst->create<ast::Function>(
-                ctx.dst->Sym(), params, ctx.dst->ty.void_(),
-                ctx.dst->Block(body), ast::DecorationList{},
-                ast::DecorationList{});
+            b.Func(name, params, b.ty.void_(), body);
           }
 
-          ctx.dst->AST().AddFunction(func);
-          return func->symbol();
+          return name;
         });
   }
 
@@ -616,14 +651,12 @@
   /// atomic operation from a storage buffer of type `buf_ty`. The function has
   /// the signature:
   // `fn atomic_op(buf : buf_ty, offset : u32, ...) -> T`
-  /// @param ctx the CloneContext
   /// @param buf_ty the storage buffer type
   /// @param el_ty the storage buffer element type
   /// @param intrinsic the atomic intrinsic
   /// @param var_user the variable user
   /// @return the name of the function that performs the load
-  Symbol AtomicFunc(CloneContext& ctx,
-                    const sem::Type* buf_ty,
+  Symbol AtomicFunc(const sem::Type* buf_ty,
                     const sem::Type* el_ty,
                     const sem::Intrinsic* intrinsic,
                     const sem::VariableUser* var_user) {
@@ -631,8 +664,8 @@
     return utils::GetOrCreate(atomic_funcs, AtomicKey{buf_ty, el_ty, op}, [&] {
       auto* buf_ast_ty = CreateASTTypeFor(ctx, buf_ty);
       auto* disable_validation =
-          ctx.dst->ASTNodes().Create<ast::DisableValidationDecoration>(
-              ctx.dst->ID(),
+          b.ASTNodes().Create<ast::DisableValidationDecoration>(
+              b.ID(),
               ast::DisabledValidation::kIgnoreConstructibleFunctionParameter);
       // The first parameter to all WGSL atomics is the expression to the
       // atomic. This is replaced with two parameters: the buffer and offset.
@@ -640,38 +673,38 @@
       ast::VariableList params = {
           // Note: The buffer parameter requires the kStorage StorageClass in
           // order for HLSL to emit this as a ByteAddressBuffer.
-          ctx.dst->create<ast::Variable>(
-              ctx.dst->Sym("buffer"), ast::StorageClass::kStorage,
-              var_user->Variable()->Access(), buf_ast_ty, true, nullptr,
-              ast::DecorationList{disable_validation}),
-          ctx.dst->Param("offset", ctx.dst->ty.u32()),
+          b.create<ast::Variable>(b.Sym("buffer"), ast::StorageClass::kStorage,
+                                  var_user->Variable()->Access(), buf_ast_ty,
+                                  true, nullptr,
+                                  ast::DecorationList{disable_validation}),
+          b.Param("offset", b.ty.u32()),
       };
 
       // Other parameters are copied as-is:
       for (size_t i = 1; i < intrinsic->Parameters().size(); i++) {
         auto& param = intrinsic->Parameters()[i];
         auto* ty = CreateASTTypeFor(ctx, param.type);
-        params.emplace_back(ctx.dst->Param("param_" + std::to_string(i), ty));
+        params.emplace_back(b.Param("param_" + std::to_string(i), ty));
       }
 
       auto* atomic = IntrinsicAtomicFor(ctx.dst, op, el_ty);
       if (atomic == nullptr) {
-        TINT_ICE(Transform, ctx.dst->Diagnostics())
+        TINT_ICE(Transform, b.Diagnostics())
             << "IntrinsicAtomicFor() returned nullptr for op " << op
             << " and type " << el_ty->type_name();
       }
 
       auto* ret_ty = CreateASTTypeFor(ctx, intrinsic->ReturnType());
-      auto* func = ctx.dst->create<ast::Function>(
-          ctx.dst->Sym(), params, ret_ty, nullptr,
+      auto* func = b.create<ast::Function>(
+          b.Sym(), params, ret_ty, nullptr,
           ast::DecorationList{
               atomic,
-              ctx.dst->ASTNodes().Create<ast::DisableValidationDecoration>(
-                  ctx.dst->ID(), ast::DisabledValidation::kFunctionHasNoBody),
+              b.ASTNodes().Create<ast::DisableValidationDecoration>(
+                  b.ID(), ast::DisabledValidation::kFunctionHasNoBody),
           },
           ast::DecorationList{});
 
-      ctx.dst->AST().AddFunction(func);
+      b.AST().AddFunction(func);
       return func->symbol();
     });
   }
@@ -777,7 +810,7 @@
 void DecomposeMemoryAccess::Run(CloneContext& ctx, const DataMap&, DataMap&) {
   auto& sem = ctx.src->Sem();
 
-  State state;
+  State state(ctx);
 
   // Scan the AST nodes for storage and uniform buffer accesses. Complex
   // expression chains (e.g. `storage_buffer.foo.bar[20].x`) are handled by
@@ -908,7 +941,7 @@
               auto* buf_ty = access.var->Type()->UnwrapRef();
               auto* el_ty = access.type->UnwrapRef()->As<sem::Atomic>()->Type();
               Symbol func =
-                  state.AtomicFunc(ctx, buf_ty, el_ty, intrinsic,
+                  state.AtomicFunc(buf_ty, el_ty, intrinsic,
                                    access.var->As<sem::VariableUser>());
 
               ast::ExpressionList args{ctx.Clone(buf), offset};
@@ -937,8 +970,8 @@
       auto* offset = access.offset->Build(ctx);
       auto* buf_ty = access.var->Type()->UnwrapRef();
       auto* el_ty = access.type->UnwrapRef();
-      Symbol func = state.LoadFunc(ctx, buf_ty, el_ty,
-                                   access.var->As<sem::VariableUser>());
+      Symbol func =
+          state.LoadFunc(buf_ty, el_ty, access.var->As<sem::VariableUser>());
       return ctx.dst->Call(func, ctx.CloneWithoutTransform(buf), offset);
     });
   }
@@ -951,7 +984,7 @@
       auto* buf_ty = store.target.var->Type()->UnwrapRef();
       auto* el_ty = store.target.type->UnwrapRef();
       auto* value = store.assignment->rhs();
-      Symbol func = state.StoreFunc(ctx, buf_ty, el_ty,
+      Symbol func = state.StoreFunc(buf_ty, el_ty,
                                     store.target.var->As<sem::VariableUser>());
       auto* call = ctx.dst->Call(func, ctx.CloneWithoutTransform(buf), offset,
                                  ctx.Clone(value));
diff --git a/src/transform/decompose_memory_access_test.cc b/src/transform/decompose_memory_access_test.cc
index 5e9e60d..58466c4 100644
--- a/src/transform/decompose_memory_access_test.cc
+++ b/src/transform/decompose_memory_access_test.cc
@@ -181,7 +181,11 @@
 }
 
 fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> array<vec3<f32>, 2> {
-  return array<vec3<f32>, 2>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)));
+  var arr : array<vec3<f32>, 2>;
+  for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
+    arr[i_1] = tint_symbol_8(buffer, (offset + (i_1 * 16u)));
+  }
+  return arr;
 }
 
 [[stage(compute), workgroup_size(1)]]
@@ -375,7 +379,11 @@
 }
 
 fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> array<vec3<f32>, 2> {
-  return array<vec3<f32>, 2>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)));
+  var arr : array<vec3<f32>, 2>;
+  for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
+    arr[i_1] = tint_symbol_8(buffer, (offset + (i_1 * 16u)));
+  }
+  return arr;
 }
 
 [[stage(compute), workgroup_size(1)]]
@@ -587,8 +595,10 @@
 }
 
 fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : array<vec3<f32>, 2>) {
-  tint_symbol_8(buffer, (offset + 0u), value[0u]);
-  tint_symbol_8(buffer, (offset + 16u), value[1u]);
+  var array = value;
+  for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
+    tint_symbol_8(buffer, (offset + (i_1 * 16u)), array[i_1]);
+  }
 }
 
 [[stage(compute), workgroup_size(1)]]
@@ -689,88 +699,92 @@
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
 [[internal(intrinsic_load_storage_i32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> i32
+fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> i32
 
 [[internal(intrinsic_load_storage_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> u32
+fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> u32
 
 [[internal(intrinsic_load_storage_f32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> f32
+fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> f32
 
 [[internal(intrinsic_load_storage_vec2_i32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<i32>
+fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<i32>
 
 [[internal(intrinsic_load_storage_vec2_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<u32>
+fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<u32>
 
 [[internal(intrinsic_load_storage_vec2_f32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<f32>
+fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec2<f32>
 
 [[internal(intrinsic_load_storage_vec3_i32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<i32>
+fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<i32>
 
 [[internal(intrinsic_load_storage_vec3_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<u32>
+fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<u32>
 
 [[internal(intrinsic_load_storage_vec3_f32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<f32>
+fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec3<f32>
 
 [[internal(intrinsic_load_storage_vec4_i32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<i32>
+fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<i32>
 
 [[internal(intrinsic_load_storage_vec4_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<u32>
+fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<u32>
 
 [[internal(intrinsic_load_storage_vec4_f32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<f32>
+fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> vec4<f32>
 
-fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x2<f32> {
-  return mat2x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)));
+fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x2<f32> {
+  return mat2x2<f32>(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u)));
 }
 
-fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x3<f32> {
-  return mat2x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)));
+fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x3<f32> {
+  return mat2x3<f32>(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u)));
 }
 
-fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x4<f32> {
-  return mat2x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)));
+fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat2x4<f32> {
+  return mat2x4<f32>(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u)));
 }
 
-fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x2<f32> {
-  return mat3x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)));
+fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x2<f32> {
+  return mat3x2<f32>(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u)), tint_symbol_6(buffer, (offset + 16u)));
 }
 
-fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x3<f32> {
-  return mat3x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)));
+fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x3<f32> {
+  return mat3x3<f32>(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u)), tint_symbol_9(buffer, (offset + 32u)));
 }
 
-fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x4<f32> {
-  return mat3x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)));
+fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat3x4<f32> {
+  return mat3x4<f32>(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u)), tint_symbol_12(buffer, (offset + 32u)));
 }
 
-fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x2<f32> {
-  return mat4x2<f32>(tint_symbol_5(buffer, (offset + 0u)), tint_symbol_5(buffer, (offset + 8u)), tint_symbol_5(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 24u)));
+fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x2<f32> {
+  return mat4x2<f32>(tint_symbol_6(buffer, (offset + 0u)), tint_symbol_6(buffer, (offset + 8u)), tint_symbol_6(buffer, (offset + 16u)), tint_symbol_6(buffer, (offset + 24u)));
 }
 
-fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x3<f32> {
-  return mat4x3<f32>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)), tint_symbol_8(buffer, (offset + 32u)), tint_symbol_8(buffer, (offset + 48u)));
+fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x3<f32> {
+  return mat4x3<f32>(tint_symbol_9(buffer, (offset + 0u)), tint_symbol_9(buffer, (offset + 16u)), tint_symbol_9(buffer, (offset + 32u)), tint_symbol_9(buffer, (offset + 48u)));
 }
 
-fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x4<f32> {
-  return mat4x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)), tint_symbol_11(buffer, (offset + 48u)));
+fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> mat4x4<f32> {
+  return mat4x4<f32>(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u)), tint_symbol_12(buffer, (offset + 32u)), tint_symbol_12(buffer, (offset + 48u)));
 }
 
-fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> array<vec3<f32>, 2> {
-  return array<vec3<f32>, 2>(tint_symbol_8(buffer, (offset + 0u)), tint_symbol_8(buffer, (offset + 16u)));
+fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> array<vec3<f32>, 2> {
+  var arr : array<vec3<f32>, 2>;
+  for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
+    arr[i_1] = tint_symbol_9(buffer, (offset + (i_1 * 16u)));
+  }
+  return arr;
 }
 
-fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> SB {
-  return SB(tint_symbol(buffer, (offset + 0u)), tint_symbol_1(buffer, (offset + 4u)), tint_symbol_2(buffer, (offset + 8u)), tint_symbol_3(buffer, (offset + 16u)), tint_symbol_4(buffer, (offset + 24u)), tint_symbol_5(buffer, (offset + 32u)), tint_symbol_6(buffer, (offset + 48u)), tint_symbol_7(buffer, (offset + 64u)), tint_symbol_8(buffer, (offset + 80u)), tint_symbol_9(buffer, (offset + 96u)), tint_symbol_10(buffer, (offset + 112u)), tint_symbol_11(buffer, (offset + 128u)), tint_symbol_12(buffer, (offset + 144u)), tint_symbol_13(buffer, (offset + 160u)), tint_symbol_14(buffer, (offset + 192u)), tint_symbol_15(buffer, (offset + 224u)), tint_symbol_16(buffer, (offset + 256u)), tint_symbol_17(buffer, (offset + 304u)), tint_symbol_18(buffer, (offset + 352u)), tint_symbol_19(buffer, (offset + 384u)), tint_symbol_20(buffer, (offset + 448u)), tint_symbol_21(buffer, (offset + 512u)));
+fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> SB {
+  return SB(tint_symbol_1(buffer, (offset + 0u)), tint_symbol_2(buffer, (offset + 4u)), tint_symbol_3(buffer, (offset + 8u)), tint_symbol_4(buffer, (offset + 16u)), tint_symbol_5(buffer, (offset + 24u)), tint_symbol_6(buffer, (offset + 32u)), tint_symbol_7(buffer, (offset + 48u)), tint_symbol_8(buffer, (offset + 64u)), tint_symbol_9(buffer, (offset + 80u)), tint_symbol_10(buffer, (offset + 96u)), tint_symbol_11(buffer, (offset + 112u)), tint_symbol_12(buffer, (offset + 128u)), tint_symbol_13(buffer, (offset + 144u)), tint_symbol_14(buffer, (offset + 160u)), tint_symbol_15(buffer, (offset + 192u)), tint_symbol_16(buffer, (offset + 224u)), tint_symbol_17(buffer, (offset + 256u)), tint_symbol_18(buffer, (offset + 304u)), tint_symbol_19(buffer, (offset + 352u)), tint_symbol_20(buffer, (offset + 384u)), tint_symbol_21(buffer, (offset + 448u)), tint_symbol_22(buffer, (offset + 512u)));
 }
 
 [[stage(compute), workgroup_size(1)]]
 fn main() {
-  var x : SB = tint_symbol_22(sb, 0u);
+  var x : SB = tint_symbol(sb, 0u);
 }
 )";
 
@@ -845,128 +859,130 @@
 [[group(0), binding(0)]] var<storage, read_write> sb : SB;
 
 [[internal(intrinsic_store_storage_i32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : i32)
+fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : i32)
 
 [[internal(intrinsic_store_storage_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_1([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : u32)
+fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : u32)
 
 [[internal(intrinsic_store_storage_f32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_2([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : f32)
+fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : f32)
 
 [[internal(intrinsic_store_storage_vec2_i32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<i32>)
+fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<i32>)
 
 [[internal(intrinsic_store_storage_vec2_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_4([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<u32>)
+fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<u32>)
 
 [[internal(intrinsic_store_storage_vec2_f32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_5([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<f32>)
+fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec2<f32>)
 
 [[internal(intrinsic_store_storage_vec3_i32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<i32>)
+fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<i32>)
 
 [[internal(intrinsic_store_storage_vec3_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_7([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<u32>)
+fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<u32>)
 
 [[internal(intrinsic_store_storage_vec3_f32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_8([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<f32>)
+fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec3<f32>)
 
 [[internal(intrinsic_store_storage_vec4_i32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_9([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<i32>)
+fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<i32>)
 
 [[internal(intrinsic_store_storage_vec4_u32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_10([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<u32>)
+fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<u32>)
 
 [[internal(intrinsic_store_storage_vec4_f32), internal(disable_validation__function_has_no_body)]]
-fn tint_symbol_11([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<f32>)
+fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : vec4<f32>)
 
-fn tint_symbol_12([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x2<f32>) {
-  tint_symbol_5(buffer, (offset + 0u), value[0u]);
-  tint_symbol_5(buffer, (offset + 8u), value[1u]);
+fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x2<f32>) {
+  tint_symbol_6(buffer, (offset + 0u), value[0u]);
+  tint_symbol_6(buffer, (offset + 8u), value[1u]);
 }
 
-fn tint_symbol_13([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x3<f32>) {
-  tint_symbol_8(buffer, (offset + 0u), value[0u]);
-  tint_symbol_8(buffer, (offset + 16u), value[1u]);
+fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x3<f32>) {
+  tint_symbol_9(buffer, (offset + 0u), value[0u]);
+  tint_symbol_9(buffer, (offset + 16u), value[1u]);
 }
 
-fn tint_symbol_14([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x4<f32>) {
-  tint_symbol_11(buffer, (offset + 0u), value[0u]);
-  tint_symbol_11(buffer, (offset + 16u), value[1u]);
+fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat2x4<f32>) {
+  tint_symbol_12(buffer, (offset + 0u), value[0u]);
+  tint_symbol_12(buffer, (offset + 16u), value[1u]);
 }
 
-fn tint_symbol_15([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x2<f32>) {
-  tint_symbol_5(buffer, (offset + 0u), value[0u]);
-  tint_symbol_5(buffer, (offset + 8u), value[1u]);
-  tint_symbol_5(buffer, (offset + 16u), value[2u]);
+fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x2<f32>) {
+  tint_symbol_6(buffer, (offset + 0u), value[0u]);
+  tint_symbol_6(buffer, (offset + 8u), value[1u]);
+  tint_symbol_6(buffer, (offset + 16u), value[2u]);
 }
 
-fn tint_symbol_16([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x3<f32>) {
-  tint_symbol_8(buffer, (offset + 0u), value[0u]);
-  tint_symbol_8(buffer, (offset + 16u), value[1u]);
-  tint_symbol_8(buffer, (offset + 32u), value[2u]);
+fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x3<f32>) {
+  tint_symbol_9(buffer, (offset + 0u), value[0u]);
+  tint_symbol_9(buffer, (offset + 16u), value[1u]);
+  tint_symbol_9(buffer, (offset + 32u), value[2u]);
 }
 
-fn tint_symbol_17([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x4<f32>) {
-  tint_symbol_11(buffer, (offset + 0u), value[0u]);
-  tint_symbol_11(buffer, (offset + 16u), value[1u]);
-  tint_symbol_11(buffer, (offset + 32u), value[2u]);
+fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat3x4<f32>) {
+  tint_symbol_12(buffer, (offset + 0u), value[0u]);
+  tint_symbol_12(buffer, (offset + 16u), value[1u]);
+  tint_symbol_12(buffer, (offset + 32u), value[2u]);
 }
 
-fn tint_symbol_18([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x2<f32>) {
-  tint_symbol_5(buffer, (offset + 0u), value[0u]);
-  tint_symbol_5(buffer, (offset + 8u), value[1u]);
-  tint_symbol_5(buffer, (offset + 16u), value[2u]);
-  tint_symbol_5(buffer, (offset + 24u), value[3u]);
+fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x2<f32>) {
+  tint_symbol_6(buffer, (offset + 0u), value[0u]);
+  tint_symbol_6(buffer, (offset + 8u), value[1u]);
+  tint_symbol_6(buffer, (offset + 16u), value[2u]);
+  tint_symbol_6(buffer, (offset + 24u), value[3u]);
 }
 
-fn tint_symbol_19([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x3<f32>) {
-  tint_symbol_8(buffer, (offset + 0u), value[0u]);
-  tint_symbol_8(buffer, (offset + 16u), value[1u]);
-  tint_symbol_8(buffer, (offset + 32u), value[2u]);
-  tint_symbol_8(buffer, (offset + 48u), value[3u]);
+fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x3<f32>) {
+  tint_symbol_9(buffer, (offset + 0u), value[0u]);
+  tint_symbol_9(buffer, (offset + 16u), value[1u]);
+  tint_symbol_9(buffer, (offset + 32u), value[2u]);
+  tint_symbol_9(buffer, (offset + 48u), value[3u]);
 }
 
-fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x4<f32>) {
-  tint_symbol_11(buffer, (offset + 0u), value[0u]);
-  tint_symbol_11(buffer, (offset + 16u), value[1u]);
-  tint_symbol_11(buffer, (offset + 32u), value[2u]);
-  tint_symbol_11(buffer, (offset + 48u), value[3u]);
+fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : mat4x4<f32>) {
+  tint_symbol_12(buffer, (offset + 0u), value[0u]);
+  tint_symbol_12(buffer, (offset + 16u), value[1u]);
+  tint_symbol_12(buffer, (offset + 32u), value[2u]);
+  tint_symbol_12(buffer, (offset + 48u), value[3u]);
 }
 
-fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : array<vec3<f32>, 2>) {
-  tint_symbol_8(buffer, (offset + 0u), value[0u]);
-  tint_symbol_8(buffer, (offset + 16u), value[1u]);
+fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : array<vec3<f32>, 2>) {
+  var array = value;
+  for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
+    tint_symbol_9(buffer, (offset + (i_1 * 16u)), array[i_1]);
+  }
 }
 
-fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : SB) {
-  tint_symbol(buffer, (offset + 0u), value.a);
-  tint_symbol_1(buffer, (offset + 4u), value.b);
-  tint_symbol_2(buffer, (offset + 8u), value.c);
-  tint_symbol_3(buffer, (offset + 16u), value.d);
-  tint_symbol_4(buffer, (offset + 24u), value.e);
-  tint_symbol_5(buffer, (offset + 32u), value.f);
-  tint_symbol_6(buffer, (offset + 48u), value.g);
-  tint_symbol_7(buffer, (offset + 64u), value.h);
-  tint_symbol_8(buffer, (offset + 80u), value.i);
-  tint_symbol_9(buffer, (offset + 96u), value.j);
-  tint_symbol_10(buffer, (offset + 112u), value.k);
-  tint_symbol_11(buffer, (offset + 128u), value.l);
-  tint_symbol_12(buffer, (offset + 144u), value.m);
-  tint_symbol_13(buffer, (offset + 160u), value.n);
-  tint_symbol_14(buffer, (offset + 192u), value.o);
-  tint_symbol_15(buffer, (offset + 224u), value.p);
-  tint_symbol_16(buffer, (offset + 256u), value.q);
-  tint_symbol_17(buffer, (offset + 304u), value.r);
-  tint_symbol_18(buffer, (offset + 352u), value.s);
-  tint_symbol_19(buffer, (offset + 384u), value.t);
-  tint_symbol_20(buffer, (offset + 448u), value.u);
-  tint_symbol_21(buffer, (offset + 512u), value.v);
+fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : SB) {
+  tint_symbol_1(buffer, (offset + 0u), value.a);
+  tint_symbol_2(buffer, (offset + 4u), value.b);
+  tint_symbol_3(buffer, (offset + 8u), value.c);
+  tint_symbol_4(buffer, (offset + 16u), value.d);
+  tint_symbol_5(buffer, (offset + 24u), value.e);
+  tint_symbol_6(buffer, (offset + 32u), value.f);
+  tint_symbol_7(buffer, (offset + 48u), value.g);
+  tint_symbol_8(buffer, (offset + 64u), value.h);
+  tint_symbol_9(buffer, (offset + 80u), value.i);
+  tint_symbol_10(buffer, (offset + 96u), value.j);
+  tint_symbol_11(buffer, (offset + 112u), value.k);
+  tint_symbol_12(buffer, (offset + 128u), value.l);
+  tint_symbol_13(buffer, (offset + 144u), value.m);
+  tint_symbol_14(buffer, (offset + 160u), value.n);
+  tint_symbol_15(buffer, (offset + 192u), value.o);
+  tint_symbol_16(buffer, (offset + 224u), value.p);
+  tint_symbol_17(buffer, (offset + 256u), value.q);
+  tint_symbol_18(buffer, (offset + 304u), value.r);
+  tint_symbol_19(buffer, (offset + 352u), value.s);
+  tint_symbol_20(buffer, (offset + 384u), value.t);
+  tint_symbol_21(buffer, (offset + 448u), value.u);
+  tint_symbol_22(buffer, (offset + 512u), value.v);
 }
 
 [[stage(compute), workgroup_size(1)]]
 fn main() {
-  tint_symbol_22(sb, 0u, SB());
+  tint_symbol(sb, 0u, SB());
 }
 )";
 
diff --git a/src/writer/hlsl/generator_impl_member_accessor_test.cc b/src/writer/hlsl/generator_impl_member_accessor_test.cc
index 10d7cb9..40b73d4 100644
--- a/src/writer/hlsl/generator_impl_member_accessor_test.cc
+++ b/src/writer/hlsl/generator_impl_member_accessor_test.cc
@@ -343,13 +343,13 @@
   auto* expected =
       R"(RWByteAddressBuffer data : register(u0, space1);
 
-void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, float2x3 value) {
+void tint_symbol(RWByteAddressBuffer buffer, uint offset, float2x3 value) {
   buffer.Store3((offset + 0u), asuint(value[0u]));
   buffer.Store3((offset + 16u), asuint(value[1u]));
 }
 
 void main() {
-  tint_symbol_1(data, 16u, float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
+  tint_symbol(data, 16u, float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
   return;
 }
 )";
diff --git a/test/array/assign_to_function_var.wgsl.expected.hlsl b/test/array/assign_to_function_var.wgsl.expected.hlsl
index 2c7d6be..5b15fc1 100644
--- a/test/array/assign_to_function_var.wgsl.expected.hlsl
+++ b/test/array/assign_to_function_var.wgsl.expected.hlsl
@@ -28,27 +28,34 @@
   return tint_symbol_6;
 }
 
-typedef tint_padded_array_element tint_symbol_2_ret[4];
-tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) {
-  const uint scalar_offset = ((offset + 0u)) / 4;
-  const uint scalar_offset_1 = ((offset + 16u)) / 4;
-  const uint scalar_offset_2 = ((offset + 32u)) / 4;
-  const uint scalar_offset_3 = ((offset + 48u)) / 4;
-  const tint_padded_array_element tint_symbol_7[4] = {{asint(buffer[scalar_offset / 4][scalar_offset % 4])}, {asint(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4])}, {asint(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4])}, {asint(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4])}};
-  return tint_symbol_7;
+typedef tint_padded_array_element tint_symbol_1_ret[4];
+tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
+  tint_padded_array_element arr_1[4] = (tint_padded_array_element[4])0;
+  {
+    for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+      const uint scalar_offset = ((offset + (i * 16u))) / 4;
+      arr_1[i].el = asint(buffer[scalar_offset / 4][scalar_offset % 4]);
+    }
+  }
+  return arr_1;
 }
 
-typedef tint_padded_array_element tint_symbol_4_ret[4];
-tint_symbol_4_ret tint_symbol_4(RWByteAddressBuffer buffer, uint offset) {
-  const tint_padded_array_element tint_symbol_8[4] = {{asint(buffer.Load((offset + 0u)))}, {asint(buffer.Load((offset + 16u)))}, {asint(buffer.Load((offset + 32u)))}, {asint(buffer.Load((offset + 48u)))}};
-  return tint_symbol_8;
+typedef tint_padded_array_element tint_symbol_3_ret[4];
+tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
+  tint_padded_array_element arr_2[4] = (tint_padded_array_element[4])0;
+  {
+    for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+      arr_2[i_1].el = asint(buffer.Load((offset + (i_1 * 16u))));
+    }
+  }
+  return arr_2;
 }
 
 void foo(tint_padded_array_element src_param[4]) {
   tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0;
   tint_padded_array_element tint_symbol[4] = (tint_padded_array_element[4])0;
-  const tint_padded_array_element tint_symbol_9[4] = {{1}, {2}, {3}, {3}};
-  tint_symbol = tint_symbol_9;
+  const tint_padded_array_element tint_symbol_7[4] = {{1}, {2}, {3}, {3}};
+  tint_symbol = tint_symbol_7;
   tint_symbol = src_param;
   tint_symbol = ret_arr();
   const tint_padded_array_element src_let[4] = (tint_padded_array_element[4])0;
@@ -57,8 +64,8 @@
   tint_symbol = src_private;
   tint_symbol = src_workgroup;
   tint_symbol = ret_struct_arr().arr;
-  tint_symbol = tint_symbol_2(src_uniform, 0u);
-  tint_symbol = tint_symbol_4(src_storage, 0u);
+  tint_symbol = tint_symbol_1(src_uniform, 0u);
+  tint_symbol = tint_symbol_3(src_storage, 0u);
   int dst_nested[4][3][2] = (int[4][3][2])0;
   int src_nested[4][3][2] = (int[4][3][2])0;
   dst_nested = src_nested;
diff --git a/test/array/assign_to_private_var.wgsl.expected.hlsl b/test/array/assign_to_private_var.wgsl.expected.hlsl
index 68ea5c3..8cd8ee1 100644
--- a/test/array/assign_to_private_var.wgsl.expected.hlsl
+++ b/test/array/assign_to_private_var.wgsl.expected.hlsl
@@ -30,26 +30,33 @@
   return tint_symbol_6;
 }
 
-typedef tint_padded_array_element tint_symbol_2_ret[4];
-tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) {
-  const uint scalar_offset = ((offset + 0u)) / 4;
-  const uint scalar_offset_1 = ((offset + 16u)) / 4;
-  const uint scalar_offset_2 = ((offset + 32u)) / 4;
-  const uint scalar_offset_3 = ((offset + 48u)) / 4;
-  const tint_padded_array_element tint_symbol_7[4] = {{asint(buffer[scalar_offset / 4][scalar_offset % 4])}, {asint(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4])}, {asint(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4])}, {asint(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4])}};
-  return tint_symbol_7;
+typedef tint_padded_array_element tint_symbol_1_ret[4];
+tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
+  tint_padded_array_element arr_1[4] = (tint_padded_array_element[4])0;
+  {
+    for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+      const uint scalar_offset = ((offset + (i * 16u))) / 4;
+      arr_1[i].el = asint(buffer[scalar_offset / 4][scalar_offset % 4]);
+    }
+  }
+  return arr_1;
 }
 
-typedef tint_padded_array_element tint_symbol_4_ret[4];
-tint_symbol_4_ret tint_symbol_4(RWByteAddressBuffer buffer, uint offset) {
-  const tint_padded_array_element tint_symbol_8[4] = {{asint(buffer.Load((offset + 0u)))}, {asint(buffer.Load((offset + 16u)))}, {asint(buffer.Load((offset + 32u)))}, {asint(buffer.Load((offset + 48u)))}};
-  return tint_symbol_8;
+typedef tint_padded_array_element tint_symbol_3_ret[4];
+tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
+  tint_padded_array_element arr_2[4] = (tint_padded_array_element[4])0;
+  {
+    for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+      arr_2[i_1].el = asint(buffer.Load((offset + (i_1 * 16u))));
+    }
+  }
+  return arr_2;
 }
 
 void foo(tint_padded_array_element src_param[4]) {
   tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0;
-  const tint_padded_array_element tint_symbol_9[4] = {{1}, {2}, {3}, {3}};
-  tint_symbol = tint_symbol_9;
+  const tint_padded_array_element tint_symbol_7[4] = {{1}, {2}, {3}, {3}};
+  tint_symbol = tint_symbol_7;
   tint_symbol = src_param;
   tint_symbol = ret_arr();
   const tint_padded_array_element src_let[4] = (tint_padded_array_element[4])0;
@@ -58,8 +65,8 @@
   tint_symbol = src_private;
   tint_symbol = src_workgroup;
   tint_symbol = ret_struct_arr().arr;
-  tint_symbol = tint_symbol_2(src_uniform, 0u);
-  tint_symbol = tint_symbol_4(src_storage, 0u);
+  tint_symbol = tint_symbol_1(src_uniform, 0u);
+  tint_symbol = tint_symbol_3(src_storage, 0u);
   int src_nested[4][3][2] = (int[4][3][2])0;
   dst_nested = src_nested;
 }
diff --git a/test/array/assign_to_storage_var.wgsl.expected.hlsl b/test/array/assign_to_storage_var.wgsl.expected.hlsl
index ee58066..8825ef3 100644
--- a/test/array/assign_to_storage_var.wgsl.expected.hlsl
+++ b/test/array/assign_to_storage_var.wgsl.expected.hlsl
@@ -30,61 +30,79 @@
   return tint_symbol_12;
 }
 
-void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[4]) {
-  buffer.Store((offset + 0u), asuint(value[0u].el));
-  buffer.Store((offset + 16u), asuint(value[1u].el));
-  buffer.Store((offset + 32u), asuint(value[2u].el));
-  buffer.Store((offset + 48u), asuint(value[3u].el));
+void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[4]) {
+  tint_padded_array_element array[4] = value;
+  {
+    for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+      buffer.Store((offset + (i * 16u)), asuint(array[i].el));
+    }
+  }
 }
 
-typedef tint_padded_array_element tint_symbol_4_ret[4];
-tint_symbol_4_ret tint_symbol_4(uint4 buffer[4], uint offset) {
-  const uint scalar_offset = ((offset + 0u)) / 4;
-  const uint scalar_offset_1 = ((offset + 16u)) / 4;
-  const uint scalar_offset_2 = ((offset + 32u)) / 4;
-  const uint scalar_offset_3 = ((offset + 48u)) / 4;
-  const tint_padded_array_element tint_symbol_13[4] = {{asint(buffer[scalar_offset / 4][scalar_offset % 4])}, {asint(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4])}, {asint(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4])}, {asint(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4])}};
-  return tint_symbol_13;
+typedef tint_padded_array_element tint_symbol_3_ret[4];
+tint_symbol_3_ret tint_symbol_3(uint4 buffer[4], uint offset) {
+  tint_padded_array_element arr_1[4] = (tint_padded_array_element[4])0;
+  {
+    for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+      const uint scalar_offset = ((offset + (i_1 * 16u))) / 4;
+      arr_1[i_1].el = asint(buffer[scalar_offset / 4][scalar_offset % 4]);
+    }
+  }
+  return arr_1;
 }
 
-typedef tint_padded_array_element tint_symbol_6_ret[4];
-tint_symbol_6_ret tint_symbol_6(RWByteAddressBuffer buffer, uint offset) {
-  const tint_padded_array_element tint_symbol_14[4] = {{asint(buffer.Load((offset + 0u)))}, {asint(buffer.Load((offset + 16u)))}, {asint(buffer.Load((offset + 32u)))}, {asint(buffer.Load((offset + 48u)))}};
-  return tint_symbol_14;
+typedef tint_padded_array_element tint_symbol_5_ret[4];
+tint_symbol_5_ret tint_symbol_5(RWByteAddressBuffer buffer, uint offset) {
+  tint_padded_array_element arr_2[4] = (tint_padded_array_element[4])0;
+  {
+    for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) {
+      arr_2[i_2].el = asint(buffer.Load((offset + (i_2 * 16u))));
+    }
+  }
+  return arr_2;
 }
 
-void tint_symbol_8(RWByteAddressBuffer buffer, uint offset, int value[2]) {
-  buffer.Store((offset + 0u), asuint(value[0u]));
-  buffer.Store((offset + 4u), asuint(value[1u]));
+void tint_symbol_9(RWByteAddressBuffer buffer, uint offset, int value[2]) {
+  int array_3[2] = value;
+  {
+    for(uint i_3 = 0u; (i_3 < 2u); i_3 = (i_3 + 1u)) {
+      buffer.Store((offset + (i_3 * 4u)), asuint(array_3[i_3]));
+    }
+  }
 }
 
-void tint_symbol_9(RWByteAddressBuffer buffer, uint offset, int value[3][2]) {
-  tint_symbol_8(buffer, (offset + 0u), value[0u]);
-  tint_symbol_8(buffer, (offset + 8u), value[1u]);
-  tint_symbol_8(buffer, (offset + 16u), value[2u]);
+void tint_symbol_8(RWByteAddressBuffer buffer, uint offset, int value[3][2]) {
+  int array_2[3][2] = value;
+  {
+    for(uint i_4 = 0u; (i_4 < 3u); i_4 = (i_4 + 1u)) {
+      tint_symbol_9(buffer, (offset + (i_4 * 8u)), array_2[i_4]);
+    }
+  }
 }
 
-void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, int value[4][3][2]) {
-  tint_symbol_9(buffer, (offset + 0u), value[0u]);
-  tint_symbol_9(buffer, (offset + 24u), value[1u]);
-  tint_symbol_9(buffer, (offset + 48u), value[2u]);
-  tint_symbol_9(buffer, (offset + 72u), value[3u]);
+void tint_symbol_7(RWByteAddressBuffer buffer, uint offset, int value[4][3][2]) {
+  int array_1[4][3][2] = value;
+  {
+    for(uint i_5 = 0u; (i_5 < 4u); i_5 = (i_5 + 1u)) {
+      tint_symbol_8(buffer, (offset + (i_5 * 24u)), array_1[i_5]);
+    }
+  }
 }
 
 void foo(tint_padded_array_element src_param[4]) {
   tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0;
-  const tint_padded_array_element tint_symbol_15[4] = {{1}, {2}, {3}, {3}};
-  tint_symbol_2(tint_symbol, 0u, tint_symbol_15);
-  tint_symbol_2(tint_symbol, 0u, src_param);
-  tint_symbol_2(tint_symbol, 0u, ret_arr());
+  const tint_padded_array_element tint_symbol_13[4] = {{1}, {2}, {3}, {3}};
+  tint_symbol_1(tint_symbol, 0u, tint_symbol_13);
+  tint_symbol_1(tint_symbol, 0u, src_param);
+  tint_symbol_1(tint_symbol, 0u, ret_arr());
   const tint_padded_array_element src_let[4] = (tint_padded_array_element[4])0;
-  tint_symbol_2(tint_symbol, 0u, src_let);
-  tint_symbol_2(tint_symbol, 0u, src_function);
-  tint_symbol_2(tint_symbol, 0u, src_private);
-  tint_symbol_2(tint_symbol, 0u, src_workgroup);
-  tint_symbol_2(tint_symbol, 0u, ret_struct_arr().arr);
-  tint_symbol_2(tint_symbol, 0u, tint_symbol_4(src_uniform, 0u));
-  tint_symbol_2(tint_symbol, 0u, tint_symbol_6(src_storage, 0u));
+  tint_symbol_1(tint_symbol, 0u, src_let);
+  tint_symbol_1(tint_symbol, 0u, src_function);
+  tint_symbol_1(tint_symbol, 0u, src_private);
+  tint_symbol_1(tint_symbol, 0u, src_workgroup);
+  tint_symbol_1(tint_symbol, 0u, ret_struct_arr().arr);
+  tint_symbol_1(tint_symbol, 0u, tint_symbol_3(src_uniform, 0u));
+  tint_symbol_1(tint_symbol, 0u, tint_symbol_5(src_storage, 0u));
   int src_nested[4][3][2] = (int[4][3][2])0;
-  tint_symbol_10(dst_nested, 0u, src_nested);
+  tint_symbol_7(dst_nested, 0u, src_nested);
 }
diff --git a/test/array/assign_to_workgroup_var.wgsl.expected.hlsl b/test/array/assign_to_workgroup_var.wgsl.expected.hlsl
index 48c5fae..43a8a1c 100644
--- a/test/array/assign_to_workgroup_var.wgsl.expected.hlsl
+++ b/test/array/assign_to_workgroup_var.wgsl.expected.hlsl
@@ -30,26 +30,33 @@
   return tint_symbol_6;
 }
 
-typedef tint_padded_array_element tint_symbol_2_ret[4];
-tint_symbol_2_ret tint_symbol_2(uint4 buffer[4], uint offset) {
-  const uint scalar_offset = ((offset + 0u)) / 4;
-  const uint scalar_offset_1 = ((offset + 16u)) / 4;
-  const uint scalar_offset_2 = ((offset + 32u)) / 4;
-  const uint scalar_offset_3 = ((offset + 48u)) / 4;
-  const tint_padded_array_element tint_symbol_7[4] = {{asint(buffer[scalar_offset / 4][scalar_offset % 4])}, {asint(buffer[scalar_offset_1 / 4][scalar_offset_1 % 4])}, {asint(buffer[scalar_offset_2 / 4][scalar_offset_2 % 4])}, {asint(buffer[scalar_offset_3 / 4][scalar_offset_3 % 4])}};
-  return tint_symbol_7;
+typedef tint_padded_array_element tint_symbol_1_ret[4];
+tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
+  tint_padded_array_element arr_1[4] = (tint_padded_array_element[4])0;
+  {
+    for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+      const uint scalar_offset = ((offset + (i * 16u))) / 4;
+      arr_1[i].el = asint(buffer[scalar_offset / 4][scalar_offset % 4]);
+    }
+  }
+  return arr_1;
 }
 
-typedef tint_padded_array_element tint_symbol_4_ret[4];
-tint_symbol_4_ret tint_symbol_4(RWByteAddressBuffer buffer, uint offset) {
-  const tint_padded_array_element tint_symbol_8[4] = {{asint(buffer.Load((offset + 0u)))}, {asint(buffer.Load((offset + 16u)))}, {asint(buffer.Load((offset + 32u)))}, {asint(buffer.Load((offset + 48u)))}};
-  return tint_symbol_8;
+typedef tint_padded_array_element tint_symbol_3_ret[4];
+tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
+  tint_padded_array_element arr_2[4] = (tint_padded_array_element[4])0;
+  {
+    for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+      arr_2[i_1].el = asint(buffer.Load((offset + (i_1 * 16u))));
+    }
+  }
+  return arr_2;
 }
 
 void foo(tint_padded_array_element src_param[4]) {
   tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0;
-  const tint_padded_array_element tint_symbol_9[4] = {{1}, {2}, {3}, {3}};
-  tint_symbol = tint_symbol_9;
+  const tint_padded_array_element tint_symbol_7[4] = {{1}, {2}, {3}, {3}};
+  tint_symbol = tint_symbol_7;
   tint_symbol = src_param;
   tint_symbol = ret_arr();
   const tint_padded_array_element src_let[4] = (tint_padded_array_element[4])0;
@@ -58,8 +65,8 @@
   tint_symbol = src_private;
   tint_symbol = src_workgroup;
   tint_symbol = ret_struct_arr().arr;
-  tint_symbol = tint_symbol_2(src_uniform, 0u);
-  tint_symbol = tint_symbol_4(src_storage, 0u);
+  tint_symbol = tint_symbol_1(src_uniform, 0u);
+  tint_symbol = tint_symbol_3(src_storage, 0u);
   int src_nested[4][3][2] = (int[4][3][2])0;
   dst_nested = src_nested;
 }
diff --git a/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl b/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl
index 9b9390a..24312d1 100644
--- a/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl
+++ b/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl
@@ -8,14 +8,19 @@
   return float2x3(asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load3((offset + 16u))));
 }
 
-float3x2 tint_symbol_10(ByteAddressBuffer buffer, uint offset) {
+float3x2 tint_symbol_9(ByteAddressBuffer buffer, uint offset) {
   return float3x2(asfloat(buffer.Load2((offset + 0u))), asfloat(buffer.Load2((offset + 8u))), asfloat(buffer.Load2((offset + 16u))));
 }
 
-typedef int4 tint_symbol_12_ret[4];
-tint_symbol_12_ret tint_symbol_12(ByteAddressBuffer buffer, uint offset) {
-  const int4 tint_symbol_13[4] = {asint(buffer.Load4((offset + 0u))), asint(buffer.Load4((offset + 16u))), asint(buffer.Load4((offset + 32u))), asint(buffer.Load4((offset + 48u)))};
-  return tint_symbol_13;
+typedef int4 tint_symbol_11_ret[4];
+tint_symbol_11_ret tint_symbol_11(ByteAddressBuffer buffer, uint offset) {
+  int4 arr_1[4] = (int4[4])0;
+  {
+    for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+      arr_1[i_1] = asint(buffer.Load4((offset + (i_1 * 16u))));
+    }
+  }
+  return arr_1;
 }
 
 [numthreads(1, 1, 1)]
@@ -28,7 +33,7 @@
   const float3 e = asfloat(s.Load3(((176u * idx) + 32u)));
   const float f = asfloat(s.Load(((176u * idx) + 44u)));
   const float2x3 g = tint_symbol_8(s, ((176u * idx) + 48u));
-  const float3x2 h = tint_symbol_10(s, ((176u * idx) + 80u));
-  const int4 i[4] = tint_symbol_12(s, ((176u * idx) + 112u));
+  const float3x2 h = tint_symbol_9(s, ((176u * idx) + 80u));
+  const int4 i[4] = tint_symbol_11(s, ((176u * idx) + 112u));
   return;
 }
diff --git a/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl b/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl
index 74aeaf2..c18b7c6 100644
--- a/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl
+++ b/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl
@@ -9,17 +9,19 @@
   buffer.Store3((offset + 16u), asuint(value[1u]));
 }
 
-void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, float3x2 value) {
+void tint_symbol_9(RWByteAddressBuffer buffer, uint offset, float3x2 value) {
   buffer.Store2((offset + 0u), asuint(value[0u]));
   buffer.Store2((offset + 8u), asuint(value[1u]));
   buffer.Store2((offset + 16u), asuint(value[2u]));
 }
 
-void tint_symbol_12(RWByteAddressBuffer buffer, uint offset, int4 value[4]) {
-  buffer.Store4((offset + 0u), asuint(value[0u]));
-  buffer.Store4((offset + 16u), asuint(value[1u]));
-  buffer.Store4((offset + 32u), asuint(value[2u]));
-  buffer.Store4((offset + 48u), asuint(value[3u]));
+void tint_symbol_11(RWByteAddressBuffer buffer, uint offset, int4 value[4]) {
+  int4 array[4] = value;
+  {
+    for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+      buffer.Store4((offset + (i_1 * 16u)), asuint(array[i_1]));
+    }
+  }
 }
 
 [numthreads(1, 1, 1)]
@@ -32,8 +34,8 @@
   s.Store3(((176u * idx) + 32u), asuint(float3(0.0f, 0.0f, 0.0f)));
   s.Store(((176u * idx) + 44u), asuint(0.0f));
   tint_symbol_8(s, ((176u * idx) + 48u), float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
-  tint_symbol_10(s, ((176u * idx) + 80u), float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
+  tint_symbol_9(s, ((176u * idx) + 80u), float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
   const int4 tint_symbol_13[4] = (int4[4])0;
-  tint_symbol_12(s, ((176u * idx) + 112u), tint_symbol_13);
+  tint_symbol_11(s, ((176u * idx) + 112u), tint_symbol_13);
   return;
 }
diff --git a/test/buffer/storage/static_index/read.wgsl.expected.hlsl b/test/buffer/storage/static_index/read.wgsl.expected.hlsl
index c14caff..42a6a71 100644
--- a/test/buffer/storage/static_index/read.wgsl.expected.hlsl
+++ b/test/buffer/storage/static_index/read.wgsl.expected.hlsl
@@ -11,7 +11,7 @@
   return float2x3(asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load3((offset + 16u))));
 }
 
-float3x2 tint_symbol_8(ByteAddressBuffer buffer, uint offset) {
+float3x2 tint_symbol_7(ByteAddressBuffer buffer, uint offset) {
   return float3x2(asfloat(buffer.Load2((offset + 0u))), asfloat(buffer.Load2((offset + 8u))), asfloat(buffer.Load2((offset + 16u))));
 }
 
@@ -22,8 +22,13 @@
 
 typedef tint_padded_array_element tint_symbol_10_ret[4];
 tint_symbol_10_ret tint_symbol_10(ByteAddressBuffer buffer, uint offset) {
-  const tint_padded_array_element tint_symbol_12[4] = {{tint_symbol_9(buffer, (offset + 0u))}, {tint_symbol_9(buffer, (offset + 16u))}, {tint_symbol_9(buffer, (offset + 32u))}, {tint_symbol_9(buffer, (offset + 48u))}};
-  return tint_symbol_12;
+  tint_padded_array_element arr[4] = (tint_padded_array_element[4])0;
+  {
+    for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+      arr[i_1].el = tint_symbol_9(buffer, (offset + (i_1 * 16u)));
+    }
+  }
+  return arr;
 }
 
 [numthreads(1, 1, 1)]
@@ -35,7 +40,7 @@
   const float3 e = asfloat(s.Load3(32u));
   const float f = asfloat(s.Load(44u));
   const float2x3 g = tint_symbol_6(s, 48u);
-  const float3x2 h = tint_symbol_8(s, 80u);
+  const float3x2 h = tint_symbol_7(s, 80u);
   const Inner i = tint_symbol_9(s, 104u);
   const tint_padded_array_element j[4] = tint_symbol_10(s, 108u);
   return;
diff --git a/test/buffer/storage/static_index/write.wgsl.expected.hlsl b/test/buffer/storage/static_index/write.wgsl.expected.hlsl
index 7e569cd..8c9b7ea 100644
--- a/test/buffer/storage/static_index/write.wgsl.expected.hlsl
+++ b/test/buffer/storage/static_index/write.wgsl.expected.hlsl
@@ -12,7 +12,7 @@
   buffer.Store3((offset + 16u), asuint(value[1u]));
 }
 
-void tint_symbol_8(RWByteAddressBuffer buffer, uint offset, float3x2 value) {
+void tint_symbol_7(RWByteAddressBuffer buffer, uint offset, float3x2 value) {
   buffer.Store2((offset + 0u), asuint(value[0u]));
   buffer.Store2((offset + 8u), asuint(value[1u]));
   buffer.Store2((offset + 16u), asuint(value[2u]));
@@ -23,10 +23,12 @@
 }
 
 void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[4]) {
-  tint_symbol_9(buffer, (offset + 0u), value[0u].el);
-  tint_symbol_9(buffer, (offset + 16u), value[1u].el);
-  tint_symbol_9(buffer, (offset + 32u), value[2u].el);
-  tint_symbol_9(buffer, (offset + 48u), value[3u].el);
+  tint_padded_array_element array[4] = value;
+  {
+    for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+      tint_symbol_9(buffer, (offset + (i_1 * 16u)), array[i_1].el);
+    }
+  }
 }
 
 [numthreads(1, 1, 1)]
@@ -38,7 +40,7 @@
   s.Store3(32u, asuint(float3(0.0f, 0.0f, 0.0f)));
   s.Store(44u, asuint(0.0f));
   tint_symbol_6(s, 48u, float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
-  tint_symbol_8(s, 80u, float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
+  tint_symbol_7(s, 80u, float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
   const Inner tint_symbol_11 = (Inner)0;
   tint_symbol_9(s, 104u, tint_symbol_11);
   const tint_padded_array_element tint_symbol_12[4] = (tint_padded_array_element[4])0;
diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.hlsl b/test/buffer/uniform/static_index/read.wgsl.expected.hlsl
index fa6d133..f66c81b 100644
--- a/test/buffer/uniform/static_index/read.wgsl.expected.hlsl
+++ b/test/buffer/uniform/static_index/read.wgsl.expected.hlsl
@@ -15,7 +15,7 @@
   return float2x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz));
 }
 
-float3x2 tint_symbol_9(uint4 buffer[13], uint offset) {
+float3x2 tint_symbol_8(uint4 buffer[13], uint offset) {
   const uint scalar_offset_2 = ((offset + 0u)) / 4;
   uint4 ubo_load = buffer[scalar_offset_2 / 4];
   const uint scalar_offset_3 = ((offset + 8u)) / 4;
@@ -33,8 +33,13 @@
 
 typedef tint_padded_array_element tint_symbol_11_ret[4];
 tint_symbol_11_ret tint_symbol_11(uint4 buffer[13], uint offset) {
-  const tint_padded_array_element tint_symbol_13[4] = {{tint_symbol_10(buffer, (offset + 0u))}, {tint_symbol_10(buffer, (offset + 16u))}, {tint_symbol_10(buffer, (offset + 32u))}, {tint_symbol_10(buffer, (offset + 48u))}};
-  return tint_symbol_13;
+  tint_padded_array_element arr[4] = (tint_padded_array_element[4])0;
+  {
+    for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+      arr[i_1].el = tint_symbol_10(buffer, (offset + (i_1 * 16u)));
+    }
+  }
+  return arr;
 }
 
 [numthreads(1, 1, 1)]
@@ -48,7 +53,7 @@
   const int2 g = asint(s[3].xy);
   const int2 h = asint(s[3].zw);
   const float2x3 i = tint_symbol_7(s, 64u);
-  const float3x2 j = tint_symbol_9(s, 96u);
+  const float3x2 j = tint_symbol_8(s, 96u);
   const Inner k = tint_symbol_10(s, 128u);
   const tint_padded_array_element l[4] = tint_symbol_11(s, 144u);
   return;
diff --git a/test/bug/tint/403.wgsl.expected.hlsl b/test/bug/tint/403.wgsl.expected.hlsl
index ef0848c..f798195 100644
--- a/test/bug/tint/403.wgsl.expected.hlsl
+++ b/test/bug/tint/403.wgsl.expected.hlsl
@@ -12,7 +12,7 @@
   float4 value : SV_Position;
 };
 
-float2x2 tint_symbol_4(uint4 buffer[1], uint offset) {
+float2x2 tint_symbol_3(uint4 buffer[1], uint offset) {
   const uint scalar_offset = ((offset + 0u)) / 4;
   uint4 ubo_load = buffer[scalar_offset / 4];
   const uint scalar_offset_1 = ((offset + 8u)) / 4;
@@ -20,7 +20,7 @@
   return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
 }
 
-float2x2 tint_symbol_6(uint4 buffer[1], uint offset) {
+float2x2 tint_symbol_5(uint4 buffer[1], uint offset) {
   const uint scalar_offset_2 = ((offset + 0u)) / 4;
   uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
   const uint scalar_offset_3 = ((offset + 8u)) / 4;
@@ -31,8 +31,8 @@
 tint_symbol_2 main(tint_symbol_1 tint_symbol) {
   const uint gl_VertexIndex = tint_symbol.gl_VertexIndex;
   float2 indexable[3] = (float2[3])0;
-  const float2x2 x_23 = tint_symbol_4(x_20, 0u);
-  const float2x2 x_28 = tint_symbol_6(x_26, 0u);
+  const float2x2 x_23 = tint_symbol_3(x_20, 0u);
+  const float2x2 x_28 = tint_symbol_5(x_26, 0u);
   const uint x_46 = gl_VertexIndex;
   const float2 tint_symbol_7[3] = {float2(-1.0f, 1.0f), float2(1.0f, 1.0f), float2(-1.0f, -1.0f)};
   indexable = tint_symbol_7;
diff --git a/test/bug/tint/870.spvasm.expected.hlsl b/test/bug/tint/870.spvasm.expected.hlsl
index b3b5071..db19a18 100644
--- a/test/bug/tint/870.spvasm.expected.hlsl
+++ b/test/bug/tint/870.spvasm.expected.hlsl
@@ -1,14 +1,19 @@
 ByteAddressBuffer sspp962805860buildInformation : register(t2, space0);
 
-typedef int tint_symbol_1_ret[6];
-tint_symbol_1_ret tint_symbol_1(ByteAddressBuffer buffer, uint offset) {
-  const int tint_symbol_2[6] = {asint(buffer.Load((offset + 0u))), asint(buffer.Load((offset + 4u))), asint(buffer.Load((offset + 8u))), asint(buffer.Load((offset + 12u))), asint(buffer.Load((offset + 16u))), asint(buffer.Load((offset + 20u)))};
-  return tint_symbol_2;
+typedef int tint_symbol_ret[6];
+tint_symbol_ret tint_symbol(ByteAddressBuffer buffer, uint offset) {
+  int arr[6] = (int[6])0;
+  {
+    for(uint i = 0u; (i < 6u); i = (i + 1u)) {
+      arr[i] = asint(buffer.Load((offset + (i * 4u))));
+    }
+  }
+  return arr;
 }
 
 void main_1() {
   int orientation[6] = (int[6])0;
-  const int x_23[6] = tint_symbol_1(sspp962805860buildInformation, 36u);
+  const int x_23[6] = tint_symbol(sspp962805860buildInformation, 36u);
   orientation[0] = x_23[0u];
   orientation[1] = x_23[1u];
   orientation[2] = x_23[2u];
diff --git a/test/bug/tint/922.wgsl.expected.hlsl b/test/bug/tint/922.wgsl.expected.hlsl
index e0b2049..c3b7973 100644
--- a/test/bug/tint/922.wgsl.expected.hlsl
+++ b/test/bug/tint/922.wgsl.expected.hlsl
@@ -154,7 +154,7 @@
   return o4;
 }
 
-Mat4x3_ tint_symbol_5(uint4 buffer[96], uint offset) {
+Mat4x3_ tint_symbol_4(uint4 buffer[96], uint offset) {
   const uint scalar_offset = ((offset + 0u)) / 4;
   const uint scalar_offset_1 = ((offset + 16u)) / 4;
   const uint scalar_offset_2 = ((offset + 32u)) / 4;
@@ -162,7 +162,7 @@
   return tint_symbol_10;
 }
 
-Mat4x4_ tint_symbol_7(uint4 buffer[4], uint offset) {
+Mat4x4_ tint_symbol_6(uint4 buffer[4], uint offset) {
   const uint scalar_offset_3 = ((offset + 0u)) / 4;
   const uint scalar_offset_4 = ((offset + 16u)) / 4;
   const uint scalar_offset_5 = ((offset + 32u)) / 4;
@@ -181,13 +181,13 @@
 void main1() {
   Mat4x3_ t_PosMtx = (Mat4x3_)0;
   float2 t_TexSpaceCoord = float2(0.0f, 0.0f);
-  const Mat4x3_ _e18 = tint_symbol_5(global2, (48u * uint(int(a_PosMtxIdx1))));
+  const Mat4x3_ _e18 = tint_symbol_4(global2, (48u * uint(int(a_PosMtxIdx1))));
   t_PosMtx = _e18;
   const Mat4x4_ _e24 = _Mat4x4_1(t_PosMtx);
   const float3 _e25 = a_Position1;
   const Mat4x4_ _e30 = _Mat4x4_1(t_PosMtx);
   const float4 _e34 = Mul(_e30, float4(a_Position1, 1.0f));
-  const Mat4x4_ _e35 = tint_symbol_7(global, 0u);
+  const Mat4x4_ _e35 = tint_symbol_6(global, 0u);
   const Mat4x4_ _e38 = _Mat4x4_1(t_PosMtx);
   const float3 _e39 = a_Position1;
   const Mat4x4_ _e44 = _Mat4x4_1(t_PosMtx);
diff --git a/test/bug/tint/998.wgsl.expected.hlsl b/test/bug/tint/998.wgsl.expected.hlsl
index 1acfe64..88f4d37 100644
--- a/test/bug/tint/998.wgsl.expected.hlsl
+++ b/test/bug/tint/998.wgsl.expected.hlsl
@@ -1,5 +1,3 @@
-SKIP: FAILED
-
 cbuffer cbuffer_constants : register(b0, space1) {
   uint4 constants[1];
 };
@@ -17,5 +15,3 @@
   s.data[constants[0].x] = 0u;
   return;
 }
-C:\src\tint\test\Shader@0x0000015D0E1BAC50(15,3-24): error X3500: array reference cannot be used as an l-value; not natively addressable
-
diff --git a/test/samples/cube.wgsl.expected.hlsl b/test/samples/cube.wgsl.expected.hlsl
index 9791b42..f83fd1f 100644
--- a/test/samples/cube.wgsl.expected.hlsl
+++ b/test/samples/cube.wgsl.expected.hlsl
@@ -19,7 +19,7 @@
   float4 Position : SV_Position;
 };
 
-float4x4 tint_symbol_7(uint4 buffer[4], uint offset) {
+float4x4 tint_symbol_6(uint4 buffer[4], uint offset) {
   const uint scalar_offset = ((offset + 0u)) / 4;
   const uint scalar_offset_1 = ((offset + 16u)) / 4;
   const uint scalar_offset_2 = ((offset + 32u)) / 4;
@@ -30,7 +30,7 @@
 tint_symbol_2 vtx_main(tint_symbol_1 tint_symbol) {
   const VertexInput input = {tint_symbol.cur_position, tint_symbol.color};
   VertexOutput output = (VertexOutput)0;
-  output.Position = mul(input.cur_position, tint_symbol_7(uniforms, 0u));
+  output.Position = mul(input.cur_position, tint_symbol_6(uniforms, 0u));
   output.vtxFragColor = input.color;
   const tint_symbol_2 tint_symbol_8 = {output.vtxFragColor, output.Position};
   return tint_symbol_8;
diff --git a/test/shader_io/shared_struct_storage_buffer.wgsl.expected.hlsl b/test/shader_io/shared_struct_storage_buffer.wgsl.expected.hlsl
index af4440e..c06b40f 100644
--- a/test/shader_io/shared_struct_storage_buffer.wgsl.expected.hlsl
+++ b/test/shader_io/shared_struct_storage_buffer.wgsl.expected.hlsl
@@ -12,7 +12,7 @@
   float4 v : SV_Position;
 };
 
-void tint_symbol_5(RWByteAddressBuffer buffer, uint offset, S value) {
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, S value) {
   buffer.Store((offset + 0u), asuint(value.f));
   buffer.Store((offset + 4u), asuint(value.u));
   buffer.Store4((offset + 128u), asuint(value.v));
@@ -23,6 +23,6 @@
   const float f = input.f;
   const uint u = input.u;
   const float4 v = input.v;
-  tint_symbol_5(output, 0u, input);
+  tint_symbol_2(output, 0u, input);
   return;
 }