diff --git a/src/transform/hlsl.cc b/src/transform/hlsl.cc
index 823f2b5..baf8897 100644
--- a/src/transform/hlsl.cc
+++ b/src/transform/hlsl.cc
@@ -66,6 +66,10 @@
   manager.Add<ExternalTextureTransform>();
   manager.Add<PromoteInitializersToConstVar>();
   manager.Add<PadArrayElements>();
+
+  ZeroInitWorkgroupMemory::Config zero_init_cfg;
+  zero_init_cfg.init_arrays_with_loop_size_threshold = 32;  // 8 scalars
+  data.Add<ZeroInitWorkgroupMemory::Config>(zero_init_cfg);
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::BuiltinStyle::kStructMember);
   auto out = manager.Run(in, data);
diff --git a/src/transform/zero_init_workgroup_memory.cc b/src/transform/zero_init_workgroup_memory.cc
index 80ba53a..600a02f 100644
--- a/src/transform/zero_init_workgroup_memory.cc
+++ b/src/transform/zero_init_workgroup_memory.cc
@@ -24,6 +24,7 @@
 #include "src/utils/get_or_create.h"
 
 TINT_INSTANTIATE_TYPEINFO(tint::transform::ZeroInitWorkgroupMemory);
+TINT_INSTANTIATE_TYPEINFO(tint::transform::ZeroInitWorkgroupMemory::Config);
 
 namespace tint {
 namespace transform {
@@ -32,14 +33,16 @@
 struct ZeroInitWorkgroupMemory::State {
   /// The clone context
   CloneContext& ctx;
-  /// The built statements
-  ast::StatementList& stmts;
+  /// The config
+  Config cfg;
 
   /// Zero() generates the statements required to zero initialize the workgroup
   /// storage expression of type `ty`.
   /// @param ty the expression type
+  /// @param stmts the built statements
   /// @param get_expr a function that builds the AST nodes for the expression
   void Zero(const sem::Type* ty,
+            ast::StatementList& stmts,
             const std::function<ast::Expression*()>& get_expr) {
     if (CanZero(ty)) {
       auto* var = get_expr();
@@ -61,21 +64,32 @@
     if (auto* str = ty->As<sem::Struct>()) {
       for (auto* member : str->Members()) {
         auto name = ctx.Clone(member->Declaration()->symbol());
-        Zero(member->Type(),
+        Zero(member->Type(), stmts,
              [&] { return ctx.dst->MemberAccessor(get_expr(), name); });
       }
       return;
     }
 
     if (auto* arr = ty->As<sem::Array>()) {
-      // TODO(bclayton): If array sizes become pipeline-overridable then this
-      // will need to emit code for a loop.
-      // See https://github.com/gpuweb/gpuweb/pull/1792
-      for (size_t i = 0; i < arr->Count(); i++) {
-        Zero(arr->ElemType(), [&] {
-          return ctx.dst->IndexAccessor(get_expr(),
-                                        static_cast<ProgramBuilder::u32>(i));
-        });
+      if (ShouldEmitForLoop(arr)) {
+        auto i = ctx.dst->Symbols().New("i");
+        auto* i_decl = ctx.dst->Decl(ctx.dst->Var(i, ctx.dst->ty.i32()));
+        auto* cond = ctx.dst->create<ast::BinaryExpression>(
+            ast::BinaryOp::kLessThan, ctx.dst->Expr(i),
+            ctx.dst->Expr(static_cast<int>(arr->Count())));
+        auto* inc = ctx.dst->Assign(i, ctx.dst->Add(i, 1));
+        ast::StatementList for_stmts;
+        Zero(arr->ElemType(), for_stmts,
+             [&] { return ctx.dst->IndexAccessor(get_expr(), i); });
+        auto* body = ctx.dst->Block(for_stmts);
+        stmts.emplace_back(ctx.dst->For(i_decl, cond, inc, body));
+      } else {
+        for (size_t i = 0; i < arr->Count(); i++) {
+          Zero(arr->ElemType(), stmts, [&] {
+            return ctx.dst->IndexAccessor(get_expr(),
+                                          static_cast<ProgramBuilder::u32>(i));
+          });
+        }
       }
       return;
     }
@@ -89,7 +103,7 @@
   /// CanZero() returns false, then the type needs to be initialized by
   /// decomposing the initialization into multiple sub-initializations.
   /// @param ty the type to inspect
-  static bool CanZero(const sem::Type* ty) {
+  bool CanZero(const sem::Type* ty) {
     if (ty->Is<sem::Atomic>()) {
       return false;
     }
@@ -101,21 +115,39 @@
       }
     }
     if (auto* arr = ty->As<sem::Array>()) {
-      if (!CanZero(arr->ElemType())) {
+      if (ShouldEmitForLoop(arr) || !CanZero(arr->ElemType())) {
         return false;
       }
     }
     return true;
   }
+
+  /// @returns true if the array should be emitted as a for-loop instead of
+  /// using zero-initializer statements.
+  /// @param array the array
+  bool ShouldEmitForLoop(const sem::Array* array) {
+    // TODO(bclayton): If array sizes become pipeline-overridable then this
+    // we need to return true for these arrays.
+    // See https://github.com/gpuweb/gpuweb/pull/1792
+    return (cfg.init_arrays_with_loop_size_threshold != 0) &&
+           (array->SizeInBytes() >= cfg.init_arrays_with_loop_size_threshold);
+  }
 };
 
 ZeroInitWorkgroupMemory::ZeroInitWorkgroupMemory() = default;
 
 ZeroInitWorkgroupMemory::~ZeroInitWorkgroupMemory() = default;
 
-void ZeroInitWorkgroupMemory::Run(CloneContext& ctx, const DataMap&, DataMap&) {
+void ZeroInitWorkgroupMemory::Run(CloneContext& ctx,
+                                  const DataMap& inputs,
+                                  DataMap&) {
   auto& sem = ctx.src->Sem();
 
+  Config cfg;
+  if (auto* c = inputs.Get<Config>()) {
+    cfg = *c;
+  }
+
   for (auto* ast_func : ctx.src->AST().Functions()) {
     if (!ast_func->IsEntryPoint()) {
       continue;
@@ -129,7 +161,7 @@
       if (var->StorageClass() != ast::StorageClass::kWorkgroup) {
         continue;
       }
-      State{ctx, stmts}.Zero(var->Type()->UnwrapRef(), [&] {
+      State{ctx, cfg}.Zero(var->Type()->UnwrapRef(), stmts, [&] {
         auto var_name = ctx.Clone(var->Declaration()->symbol());
         return ctx.dst->Expr(var_name);
       });
@@ -193,5 +225,11 @@
   ctx.Clone();
 }
 
+ZeroInitWorkgroupMemory::Config::Config() = default;
+ZeroInitWorkgroupMemory::Config::Config(const Config&) = default;
+ZeroInitWorkgroupMemory::Config::~Config() = default;
+ZeroInitWorkgroupMemory::Config& ZeroInitWorkgroupMemory::Config::operator=(
+    const Config&) = default;
+
 }  // namespace transform
 }  // namespace tint
diff --git a/src/transform/zero_init_workgroup_memory.h b/src/transform/zero_init_workgroup_memory.h
index 1644b4d..e40a458 100644
--- a/src/transform/zero_init_workgroup_memory.h
+++ b/src/transform/zero_init_workgroup_memory.h
@@ -26,6 +26,27 @@
 class ZeroInitWorkgroupMemory
     : public Castable<ZeroInitWorkgroupMemory, Transform> {
  public:
+  /// Configuration options for the transform
+  struct Config : public Castable<Config, Data> {
+    /// Constructor
+    Config();
+
+    /// Copy constructor
+    Config(const Config&);
+
+    /// Destructor
+    ~Config() override;
+
+    /// Assignment operator
+    /// @returns this Config
+    Config& operator=(const Config&);
+
+    /// If greater than 0, then arrays of at least this size in bytes will be
+    /// zero initialized using a for loop. If 0, then the array is assigned a
+    /// zero initialized array with a single statement.
+    uint32_t init_arrays_with_loop_size_threshold = 0;
+  };
+
   /// Constructor
   ZeroInitWorkgroupMemory();
 
diff --git a/src/transform/zero_init_workgroup_memory_test.cc b/src/transform/zero_init_workgroup_memory_test.cc
index ab1b305..6c318ab 100644
--- a/src/transform/zero_init_workgroup_memory_test.cc
+++ b/src/transform/zero_init_workgroup_memory_test.cc
@@ -558,6 +558,56 @@
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupArray_InitWithLoop) {
+  auto* src = R"(
+struct S {
+  a : array<i32, 3>; // size: 12, less than the loop threshold
+  b : array<i32, 4>; // size: 16, equal to the loop threshold
+  c : array<i32, 5>; // size: 20, greater than the loop threshold
+};
+
+var<workgroup> w : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn f() {
+  ignore(w); // Initialization should be inserted above this statement
+}
+)";
+  auto* expect = R"(
+struct S {
+  a : array<i32, 3>;
+  b : array<i32, 4>;
+  c : array<i32, 5>;
+};
+
+var<workgroup> w : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn f([[builtin(local_invocation_index)]] local_invocation_index : u32) {
+  if ((local_invocation_index == 0u)) {
+    w.a = array<i32, 3>();
+    for(var i : i32; (i < 4); i = (i + 1)) {
+      w.b[i] = i32();
+    }
+    for(var i_1 : i32; (i_1 < 5); i_1 = (i_1 + 1)) {
+      w.c[i_1] = i32();
+    }
+  }
+  workgroupBarrier();
+  ignore(w);
+}
+)";
+
+  ZeroInitWorkgroupMemory::Config cfg;
+  cfg.init_arrays_with_loop_size_threshold = 16;
+
+  DataMap data;
+  data.Add<ZeroInitWorkgroupMemory::Config>(cfg);
+  auto got = Run<ZeroInitWorkgroupMemory>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 }  // namespace
 }  // namespace transform
 }  // namespace tint
diff --git a/test/bug/tint/914.wgsl.expected.hlsl b/test/bug/tint/914.wgsl.expected.hlsl
index 4030d77..3de9408 100644
--- a/test/bug/tint/914.wgsl.expected.hlsl
+++ b/test/bug/tint/914.wgsl.expected.hlsl
@@ -69,10 +69,16 @@
   const uint3 global_id = tint_symbol.global_id;
   const uint local_invocation_index = tint_symbol.local_invocation_index;
   if ((local_invocation_index == 0u)) {
-    const float tint_symbol_5[64][64] = (float[64][64])0;
-    mm_Asub = tint_symbol_5;
-    const float tint_symbol_6[64][64] = (float[64][64])0;
-    mm_Bsub = tint_symbol_6;
+    for(int i = 0; (i < 64); i = (i + 1)) {
+      for(int i_1 = 0; (i_1 < 64); i_1 = (i_1 + 1)) {
+        mm_Asub[i][i_1] = 0.0f;
+      }
+    }
+    for(int i_2 = 0; (i_2 < 64); i_2 = (i_2 + 1)) {
+      for(int i_3 = 0; (i_3 < 64); i_3 = (i_3 + 1)) {
+        mm_Bsub[i_2][i_3] = 0.0f;
+      }
+    }
   }
   GroupMemoryBarrierWithGroupSync();
   const uint tileRow = (local_id.y * RowPerThread);
diff --git a/test/bug/tint/942.wgsl b/test/bug/tint/942.wgsl
new file mode 100644
index 0000000..aaf7549
--- /dev/null
+++ b/test/bug/tint/942.wgsl
@@ -0,0 +1,80 @@
+[[block]] struct Params {
+  filterDim : u32;
+  blockDim : u32;
+};
+
+[[group(0), binding(0)]] var samp : sampler;
+[[group(0), binding(1)]] var<uniform> params : Params;
+[[group(1), binding(1)]] var inputTex : texture_2d<f32>;
+[[group(1), binding(2)]] var outputTex : texture_storage_2d<rgba8unorm, write>;
+
+[[block]] struct Flip {
+  value : u32;
+};
+[[group(1), binding(3)]] var<uniform> flip : Flip;
+
+// This shader blurs the input texture in one direction, depending on whether
+// |flip.value| is 0 or 1.
+// It does so by running (256 / 4) threads per workgroup to load 256
+// texels into 4 rows of shared memory. Each thread loads a
+// 4 x 4 block of texels to take advantage of the texture sampling
+// hardware.
+// Then, each thread computes the blur result by averaging the adjacent texel values
+// in shared memory.
+// Because we're operating on a subset of the texture, we cannot compute all of the
+// results since not all of the neighbors are available in shared memory.
+// Specifically, with 256 x 256 tiles, we can only compute and write out
+// square blocks of size 256 - (filterSize - 1). We compute the number of blocks
+// needed in Javascript and dispatch that amount.
+
+var<workgroup> tile : array<array<vec3<f32>, 256>, 4>;
+
+[[stage(compute), workgroup_size(64, 1, 1)]]
+fn main(
+  [[builtin(workgroup_id)]] WorkGroupID : vec3<u32>,
+  [[builtin(local_invocation_id)]] LocalInvocationID : vec3<u32>
+) {
+  let filterOffset : u32 = (params.filterDim - 1u) / 2u;
+  let dims : vec2<i32> = textureDimensions(inputTex, 0);
+
+  let baseIndex = vec2<i32>(
+    WorkGroupID.xy * vec2<u32>(params.blockDim, 4u) +
+    LocalInvocationID.xy * vec2<u32>(4u, 1u)
+  ) - vec2<i32>(i32(filterOffset), 0);
+
+  for (var r : u32 = 0u; r < 4u; r = r + 1u) {
+    for (var c : u32 = 0u; c < 4u; c = c + 1u) {
+      var loadIndex = baseIndex + vec2<i32>(i32(c), i32(r));
+      if (flip.value != 0u) {
+        loadIndex = loadIndex.yx;
+      }
+
+      tile[r][4u * LocalInvocationID.x + c] =
+        textureSampleLevel(inputTex, samp,
+          (vec2<f32>(loadIndex) + vec2<f32>(0.25, 0.25)) / vec2<f32>(dims), 0.0).rgb;
+    }
+  }
+
+  workgroupBarrier();
+
+  for (var r : u32 = 0u; r < 4u; r = r + 1u) {
+    for (var c : u32 = 0u; c < 4u; c = c + 1u) {
+      var writeIndex = baseIndex + vec2<i32>(i32(c), i32(r));
+      if (flip.value != 0u) {
+        writeIndex = writeIndex.yx;
+      }
+
+      let center : u32 = 4u * LocalInvocationID.x + c;
+      if (center >= filterOffset &&
+          center < 256u - filterOffset &&
+          all(writeIndex < dims)) {
+        var acc : vec3<f32> = vec3<f32>(0.0, 0.0, 0.0);
+        for (var f : u32 = 0u; f < params.filterDim; f = f + 1u) {
+          var i : u32 = center + f - filterOffset;
+          acc = acc + (1.0 / f32(params.filterDim)) * tile[r][i];
+        }
+        textureStore(outputTex, writeIndex, vec4<f32>(acc, 1.0));
+      }
+    }
+  }
+}
diff --git a/test/bug/tint/942.wgsl.expected.hlsl b/test/bug/tint/942.wgsl.expected.hlsl
new file mode 100644
index 0000000..ff2739f
--- /dev/null
+++ b/test/bug/tint/942.wgsl.expected.hlsl
@@ -0,0 +1,96 @@
+SamplerState samp : register(s0, space0);
+cbuffer cbuffer_params : register(b1, space0) {
+  uint4 params[1];
+};
+Texture2D<float4> inputTex : register(t1, space1);
+RWTexture2D<float4> outputTex : register(u2, space1);
+
+cbuffer cbuffer_flip : register(b3, space1) {
+  uint4 flip[1];
+};
+groupshared float3 tile[4][256];
+
+struct tint_symbol_1 {
+  uint3 LocalInvocationID : SV_GroupThreadID;
+  uint local_invocation_index : SV_GroupIndex;
+  uint3 WorkGroupID : SV_GroupID;
+};
+
+[numthreads(64, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  const uint3 WorkGroupID = tint_symbol.WorkGroupID;
+  const uint3 LocalInvocationID = tint_symbol.LocalInvocationID;
+  const uint local_invocation_index = tint_symbol.local_invocation_index;
+  if ((local_invocation_index == 0u)) {
+    for(int i_1 = 0; (i_1 < 4); i_1 = (i_1 + 1)) {
+      for(int i_2 = 0; (i_2 < 256); i_2 = (i_2 + 1)) {
+        tile[i_1][i_2] = float3(0.0f, 0.0f, 0.0f);
+      }
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  const uint scalar_offset = (0u) / 4;
+  const uint filterOffset = ((params[scalar_offset / 4][scalar_offset % 4] - 1u) / 2u);
+  int3 tint_tmp;
+  inputTex.GetDimensions(0, tint_tmp.x, tint_tmp.y, tint_tmp.z);
+  const int2 dims = tint_tmp.xy;
+  const uint scalar_offset_1 = (4u) / 4;
+  const int2 baseIndex = (int2(((WorkGroupID.xy * uint2(params[scalar_offset_1 / 4][scalar_offset_1 % 4], 4u)) + (LocalInvocationID.xy * uint2(4u, 1u)))) - int2(int(filterOffset), 0));
+  {
+    uint r = 0u;
+    for(; !(!((r < 4u))); r = (r + 1u)) {
+      {
+        uint c = 0u;
+        for(; !(!((c < 4u))); c = (c + 1u)) {
+          int2 loadIndex = (baseIndex + int2(int(c), int(r)));
+          const uint scalar_offset_2 = (0u) / 4;
+          if ((flip[scalar_offset_2 / 4][scalar_offset_2 % 4] != 0u)) {
+            loadIndex = loadIndex.yx;
+          }
+          tile[r][((4u * LocalInvocationID.x) + c)] = inputTex.SampleLevel(samp, ((float2(loadIndex) + float2(0.25f, 0.25f)) / float2(dims)), 0.0f).rgb;
+        }
+      }
+    }
+  }
+  GroupMemoryBarrierWithGroupSync();
+  {
+    uint r = 0u;
+    for(; !(!((r < 4u))); r = (r + 1u)) {
+      {
+        uint c = 0u;
+        for(; !(!((c < 4u))); c = (c + 1u)) {
+          int2 writeIndex = (baseIndex + int2(int(c), int(r)));
+          const uint scalar_offset_3 = (0u) / 4;
+          if ((flip[scalar_offset_3 / 4][scalar_offset_3 % 4] != 0u)) {
+            writeIndex = writeIndex.yx;
+          }
+          const uint center = ((4u * LocalInvocationID.x) + c);
+          bool tint_tmp_2 = (center >= filterOffset);
+          if (tint_tmp_2) {
+            tint_tmp_2 = (center < (256u - filterOffset));
+          }
+          bool tint_tmp_1 = (tint_tmp_2);
+          if (tint_tmp_1) {
+            tint_tmp_1 = all((writeIndex < dims));
+          }
+          if ((tint_tmp_1)) {
+            float3 acc = float3(0.0f, 0.0f, 0.0f);
+            {
+              uint f = 0u;
+              while (true) {
+                const uint scalar_offset_4 = (0u) / 4;
+                if (!(!(!((f < params[scalar_offset_4 / 4][scalar_offset_4 % 4]))))) { break; }
+                uint i = ((center + f) - filterOffset);
+                const uint scalar_offset_5 = (0u) / 4;
+                acc = (acc + ((1.0f / float(params[scalar_offset_5 / 4][scalar_offset_5 % 4])) * tile[r][i]));
+                f = (f + 1u);
+              }
+            }
+            outputTex[writeIndex] = float4(acc, 1.0f);
+          }
+        }
+      }
+    }
+  }
+  return;
+}
diff --git a/test/bug/tint/942.wgsl.expected.msl b/test/bug/tint/942.wgsl.expected.msl
new file mode 100644
index 0000000..16e440d
--- /dev/null
+++ b/test/bug/tint/942.wgsl.expected.msl
@@ -0,0 +1,102 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Params {
+  /* 0x0000 */ uint filterDim;
+  /* 0x0004 */ uint blockDim;
+};
+struct Flip {
+  /* 0x0000 */ uint value;
+};
+struct tint_array_wrapper_1 {
+  float3 arr[256];
+};
+struct tint_array_wrapper {
+  tint_array_wrapper_1 arr[4];
+};
+
+kernel void tint_symbol(texture2d<float, access::sample> tint_symbol_4 [[texture(1)]], sampler tint_symbol_5 [[sampler(0)]], texture2d<float, access::write> tint_symbol_6 [[texture(2)]], uint3 WorkGroupID [[threadgroup_position_in_grid]], uint3 LocalInvocationID [[thread_position_in_threadgroup]], uint local_invocation_index [[thread_index_in_threadgroup]], constant Params& params [[buffer(1)]], constant Flip& flip [[buffer(3)]]) {
+  threadgroup tint_array_wrapper tint_symbol_3;
+  if ((local_invocation_index == 0u)) {
+    tint_array_wrapper const tint_symbol_2 = {.arr={}};
+    tint_symbol_3 = tint_symbol_2;
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  uint const filterOffset = ((params.filterDim - 1u) / 2u);
+  int2 const dims = int2(tint_symbol_4.get_width(0), tint_symbol_4.get_height(0));
+  int2 const baseIndex = (int2(((WorkGroupID.xy * uint2(params.blockDim, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u)))) - int2(int(filterOffset), 0));
+  {
+    uint r = 0u;
+    while (true) {
+      if (!((r < 4u))) {
+        break;
+      }
+      {
+        uint c = 0u;
+        while (true) {
+          if (!((c < 4u))) {
+            break;
+          }
+          int2 loadIndex = (baseIndex + int2(int(c), int(r)));
+          if ((flip.value != 0u)) {
+            loadIndex = loadIndex.yx;
+          }
+          tint_symbol_3.arr[r].arr[((4u * LocalInvocationID.x) + c)] = tint_symbol_4.sample(tint_symbol_5, ((float2(loadIndex) + float2(0.25f, 0.25f)) / float2(dims)), level(0.0f)).rgb;
+          {
+            c = (c + 1u);
+          }
+        }
+      }
+      {
+        r = (r + 1u);
+      }
+    }
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  {
+    uint r = 0u;
+    while (true) {
+      if (!((r < 4u))) {
+        break;
+      }
+      {
+        uint c = 0u;
+        while (true) {
+          if (!((c < 4u))) {
+            break;
+          }
+          int2 writeIndex = (baseIndex + int2(int(c), int(r)));
+          if ((flip.value != 0u)) {
+            writeIndex = writeIndex.yx;
+          }
+          uint const center = ((4u * LocalInvocationID.x) + c);
+          if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
+            float3 acc = float3(0.0f, 0.0f, 0.0f);
+            {
+              uint f = 0u;
+              while (true) {
+                if (!((f < params.filterDim))) {
+                  break;
+                }
+                uint i = ((center + f) - filterOffset);
+                acc = (acc + ((1.0f / float(params.filterDim)) * tint_symbol_3.arr[r].arr[i]));
+                {
+                  f = (f + 1u);
+                }
+              }
+            }
+            tint_symbol_6.write(float4(acc, 1.0f), uint2(writeIndex));
+          }
+          {
+            c = (c + 1u);
+          }
+        }
+      }
+      {
+        r = (r + 1u);
+      }
+    }
+  }
+  return;
+}
+
diff --git a/test/bug/tint/942.wgsl.expected.spvasm b/test/bug/tint/942.wgsl.expected.spvasm
new file mode 100644
index 0000000..ef1f0a9
--- /dev/null
+++ b/test/bug/tint/942.wgsl.expected.spvasm
@@ -0,0 +1,374 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 239
+; Schema: 0
+               OpCapability Shader
+               OpCapability ImageQuery
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %tint_symbol_2 %tint_symbol %tint_symbol_1
+               OpExecutionMode %main LocalSize 64 1 1
+               OpName %samp "samp"
+               OpName %Params "Params"
+               OpMemberName %Params 0 "filterDim"
+               OpMemberName %Params 1 "blockDim"
+               OpName %params "params"
+               OpName %inputTex "inputTex"
+               OpName %outputTex "outputTex"
+               OpName %Flip "Flip"
+               OpMemberName %Flip 0 "value"
+               OpName %flip "flip"
+               OpName %tile "tile"
+               OpName %tint_symbol "tint_symbol"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %main "main"
+               OpName %r "r"
+               OpName %c "c"
+               OpName %loadIndex "loadIndex"
+               OpName %r_0 "r"
+               OpName %c_0 "c"
+               OpName %writeIndex "writeIndex"
+               OpName %acc "acc"
+               OpName %f "f"
+               OpName %i "i"
+               OpDecorate %samp DescriptorSet 0
+               OpDecorate %samp Binding 0
+               OpDecorate %Params Block
+               OpMemberDecorate %Params 0 Offset 0
+               OpMemberDecorate %Params 1 Offset 4
+               OpDecorate %params NonWritable
+               OpDecorate %params DescriptorSet 0
+               OpDecorate %params Binding 1
+               OpDecorate %inputTex DescriptorSet 1
+               OpDecorate %inputTex Binding 1
+               OpDecorate %outputTex NonReadable
+               OpDecorate %outputTex DescriptorSet 1
+               OpDecorate %outputTex Binding 2
+               OpDecorate %Flip Block
+               OpMemberDecorate %Flip 0 Offset 0
+               OpDecorate %flip NonWritable
+               OpDecorate %flip DescriptorSet 1
+               OpDecorate %flip Binding 3
+               OpDecorate %_arr_v3float_uint_256 ArrayStride 16
+               OpDecorate %_arr__arr_v3float_uint_256_uint_4 ArrayStride 4096
+               OpDecorate %tint_symbol BuiltIn WorkgroupId
+               OpDecorate %tint_symbol_1 BuiltIn LocalInvocationId
+               OpDecorate %tint_symbol_2 BuiltIn LocalInvocationIndex
+          %3 = OpTypeSampler
+%_ptr_UniformConstant_3 = OpTypePointer UniformConstant %3
+       %samp = OpVariable %_ptr_UniformConstant_3 UniformConstant
+       %uint = OpTypeInt 32 0
+     %Params = OpTypeStruct %uint %uint
+%_ptr_Uniform_Params = OpTypePointer Uniform %Params
+     %params = OpVariable %_ptr_Uniform_Params Uniform
+      %float = OpTypeFloat 32
+         %10 = OpTypeImage %float 2D 0 0 0 1 Unknown
+%_ptr_UniformConstant_10 = OpTypePointer UniformConstant %10
+   %inputTex = OpVariable %_ptr_UniformConstant_10 UniformConstant
+         %14 = OpTypeImage %float 2D 0 0 0 2 Rgba8
+%_ptr_UniformConstant_14 = OpTypePointer UniformConstant %14
+  %outputTex = OpVariable %_ptr_UniformConstant_14 UniformConstant
+       %Flip = OpTypeStruct %uint
+%_ptr_Uniform_Flip = OpTypePointer Uniform %Flip
+       %flip = OpVariable %_ptr_Uniform_Flip Uniform
+    %v3float = OpTypeVector %float 3
+   %uint_256 = OpConstant %uint 256
+%_arr_v3float_uint_256 = OpTypeArray %v3float %uint_256
+     %uint_4 = OpConstant %uint 4
+%_arr__arr_v3float_uint_256_uint_4 = OpTypeArray %_arr_v3float_uint_256 %uint_4
+%_ptr_Workgroup__arr__arr_v3float_uint_256_uint_4 = OpTypePointer Workgroup %_arr__arr_v3float_uint_256_uint_4
+       %tile = OpVariable %_ptr_Workgroup__arr__arr_v3float_uint_256_uint_4 Workgroup
+     %v3uint = OpTypeVector %uint 3
+%_ptr_Input_v3uint = OpTypePointer Input %v3uint
+%tint_symbol = OpVariable %_ptr_Input_v3uint Input
+%tint_symbol_1 = OpVariable %_ptr_Input_v3uint Input
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol_2 = OpVariable %_ptr_Input_uint Input
+       %void = OpTypeVoid
+         %31 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+       %bool = OpTypeBool
+         %41 = OpConstantNull %_arr__arr_v3float_uint_256_uint_4
+     %uint_2 = OpConstant %uint 2
+   %uint_264 = OpConstant %uint 264
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+     %uint_1 = OpConstant %uint 1
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+      %int_0 = OpConstant %int 0
+     %v2uint = OpTypeVector %uint 2
+         %66 = OpConstantComposite %v2uint %uint_4 %uint_1
+%_ptr_Function_uint = OpTypePointer Function %uint
+         %74 = OpConstantNull %uint
+%_ptr_Function_v2int = OpTypePointer Function %v2int
+        %102 = OpConstantNull %v2int
+%_ptr_Workgroup_v3float = OpTypePointer Workgroup %v3float
+    %v4float = OpTypeVector %float 4
+        %122 = OpTypeSampledImage %10
+    %v2float = OpTypeVector %float 2
+ %float_0_25 = OpConstant %float 0.25
+        %128 = OpConstantComposite %v2float %float_0_25 %float_0_25
+    %float_0 = OpConstant %float 0
+     %v2bool = OpTypeVector %bool 2
+        %193 = OpConstantComposite %v3float %float_0 %float_0 %float_0
+%_ptr_Function_v3float = OpTypePointer Function %v3float
+        %196 = OpConstantNull %v3float
+    %float_1 = OpConstant %float 1
+       %main = OpFunction %void None %31
+         %34 = OpLabel
+          %r = OpVariable %_ptr_Function_uint Function %74
+          %c = OpVariable %_ptr_Function_uint Function %74
+  %loadIndex = OpVariable %_ptr_Function_v2int Function %102
+        %r_0 = OpVariable %_ptr_Function_uint Function %74
+        %c_0 = OpVariable %_ptr_Function_uint Function %74
+ %writeIndex = OpVariable %_ptr_Function_v2int Function %102
+        %acc = OpVariable %_ptr_Function_v3float Function %196
+          %f = OpVariable %_ptr_Function_uint Function %74
+          %i = OpVariable %_ptr_Function_uint Function %74
+         %35 = OpLoad %uint %tint_symbol_2
+         %37 = OpIEqual %bool %35 %uint_0
+               OpSelectionMerge %39 None
+               OpBranchConditional %37 %40 %39
+         %40 = OpLabel
+               OpStore %tile %41
+               OpBranch %39
+         %39 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %46 = OpAccessChain %_ptr_Uniform_uint %params %uint_0
+         %47 = OpLoad %uint %46
+         %49 = OpISub %uint %47 %uint_1
+         %50 = OpUDiv %uint %49 %uint_2
+         %54 = OpLoad %10 %inputTex
+         %51 = OpImageQuerySizeLod %v2int %54 %int_0
+         %58 = OpLoad %v3uint %tint_symbol
+         %59 = OpVectorShuffle %v2uint %58 %58 0 1
+         %60 = OpAccessChain %_ptr_Uniform_uint %params %uint_1
+         %61 = OpLoad %uint %60
+         %62 = OpCompositeConstruct %v2uint %61 %uint_4
+         %63 = OpIMul %v2uint %59 %62
+         %64 = OpLoad %v3uint %tint_symbol_1
+         %65 = OpVectorShuffle %v2uint %64 %64 0 1
+         %67 = OpIMul %v2uint %65 %66
+         %68 = OpIAdd %v2uint %63 %67
+         %56 = OpBitcast %v2int %68
+         %69 = OpBitcast %int %50
+         %70 = OpCompositeConstruct %v2int %69 %int_0
+         %71 = OpISub %v2int %56 %70
+               OpStore %r %uint_0
+               OpBranch %75
+         %75 = OpLabel
+               OpLoopMerge %76 %77 None
+               OpBranch %78
+         %78 = OpLabel
+         %80 = OpLoad %uint %r
+         %81 = OpULessThan %bool %80 %uint_4
+         %79 = OpLogicalNot %bool %81
+               OpSelectionMerge %82 None
+               OpBranchConditional %79 %83 %82
+         %83 = OpLabel
+               OpBranch %76
+         %82 = OpLabel
+               OpStore %c %uint_0
+               OpBranch %85
+         %85 = OpLabel
+               OpLoopMerge %86 %87 None
+               OpBranch %88
+         %88 = OpLabel
+         %90 = OpLoad %uint %c
+         %91 = OpULessThan %bool %90 %uint_4
+         %89 = OpLogicalNot %bool %91
+               OpSelectionMerge %92 None
+               OpBranchConditional %89 %93 %92
+         %93 = OpLabel
+               OpBranch %86
+         %92 = OpLabel
+         %95 = OpLoad %uint %c
+         %94 = OpBitcast %int %95
+         %97 = OpLoad %uint %r
+         %96 = OpBitcast %int %97
+         %98 = OpCompositeConstruct %v2int %94 %96
+         %99 = OpIAdd %v2int %71 %98
+               OpStore %loadIndex %99
+        %103 = OpAccessChain %_ptr_Uniform_uint %flip %uint_0
+        %104 = OpLoad %uint %103
+        %105 = OpINotEqual %bool %104 %uint_0
+               OpSelectionMerge %106 None
+               OpBranchConditional %105 %107 %106
+        %107 = OpLabel
+        %108 = OpLoad %v2int %loadIndex
+        %109 = OpVectorShuffle %v2int %108 %108 1 0
+               OpStore %loadIndex %109
+               OpBranch %106
+        %106 = OpLabel
+        %110 = OpLoad %uint %r
+        %111 = OpAccessChain %_ptr_Input_uint %tint_symbol_1 %uint_0
+        %112 = OpLoad %uint %111
+        %113 = OpIMul %uint %uint_4 %112
+        %114 = OpLoad %uint %c
+        %115 = OpIAdd %uint %113 %114
+        %117 = OpAccessChain %_ptr_Workgroup_v3float %tile %110 %115
+        %120 = OpLoad %3 %samp
+        %121 = OpLoad %10 %inputTex
+        %123 = OpSampledImage %122 %121 %120
+        %126 = OpLoad %v2int %loadIndex
+        %124 = OpConvertSToF %v2float %126
+        %129 = OpFAdd %v2float %124 %128
+        %130 = OpConvertSToF %v2float %51
+        %131 = OpFDiv %v2float %129 %130
+        %118 = OpImageSampleExplicitLod %v4float %123 %131 Lod %float_0
+        %133 = OpVectorShuffle %v3float %118 %118 0 1 2
+               OpStore %117 %133
+               OpBranch %87
+         %87 = OpLabel
+        %134 = OpLoad %uint %c
+        %135 = OpIAdd %uint %134 %uint_1
+               OpStore %c %135
+               OpBranch %85
+         %86 = OpLabel
+               OpBranch %77
+         %77 = OpLabel
+        %136 = OpLoad %uint %r
+        %137 = OpIAdd %uint %136 %uint_1
+               OpStore %r %137
+               OpBranch %75
+         %76 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+               OpStore %r_0 %uint_0
+               OpBranch %140
+        %140 = OpLabel
+               OpLoopMerge %141 %142 None
+               OpBranch %143
+        %143 = OpLabel
+        %145 = OpLoad %uint %r_0
+        %146 = OpULessThan %bool %145 %uint_4
+        %144 = OpLogicalNot %bool %146
+               OpSelectionMerge %147 None
+               OpBranchConditional %144 %148 %147
+        %148 = OpLabel
+               OpBranch %141
+        %147 = OpLabel
+               OpStore %c_0 %uint_0
+               OpBranch %150
+        %150 = OpLabel
+               OpLoopMerge %151 %152 None
+               OpBranch %153
+        %153 = OpLabel
+        %155 = OpLoad %uint %c_0
+        %156 = OpULessThan %bool %155 %uint_4
+        %154 = OpLogicalNot %bool %156
+               OpSelectionMerge %157 None
+               OpBranchConditional %154 %158 %157
+        %158 = OpLabel
+               OpBranch %151
+        %157 = OpLabel
+        %160 = OpLoad %uint %c_0
+        %159 = OpBitcast %int %160
+        %162 = OpLoad %uint %r_0
+        %161 = OpBitcast %int %162
+        %163 = OpCompositeConstruct %v2int %159 %161
+        %164 = OpIAdd %v2int %71 %163
+               OpStore %writeIndex %164
+        %166 = OpAccessChain %_ptr_Uniform_uint %flip %uint_0
+        %167 = OpLoad %uint %166
+        %168 = OpINotEqual %bool %167 %uint_0
+               OpSelectionMerge %169 None
+               OpBranchConditional %168 %170 %169
+        %170 = OpLabel
+        %171 = OpLoad %v2int %writeIndex
+        %172 = OpVectorShuffle %v2int %171 %171 1 0
+               OpStore %writeIndex %172
+               OpBranch %169
+        %169 = OpLabel
+        %173 = OpAccessChain %_ptr_Input_uint %tint_symbol_1 %uint_0
+        %174 = OpLoad %uint %173
+        %175 = OpIMul %uint %uint_4 %174
+        %176 = OpLoad %uint %c_0
+        %177 = OpIAdd %uint %175 %176
+        %178 = OpUGreaterThanEqual %bool %177 %50
+               OpSelectionMerge %179 None
+               OpBranchConditional %178 %180 %179
+        %180 = OpLabel
+        %181 = OpISub %uint %uint_256 %50
+        %182 = OpULessThan %bool %177 %181
+               OpBranch %179
+        %179 = OpLabel
+        %183 = OpPhi %bool %178 %169 %182 %180
+               OpSelectionMerge %184 None
+               OpBranchConditional %183 %185 %184
+        %185 = OpLabel
+        %187 = OpLoad %v2int %writeIndex
+        %188 = OpSLessThan %v2bool %187 %51
+        %186 = OpAll %bool %188
+               OpBranch %184
+        %184 = OpLabel
+        %190 = OpPhi %bool %183 %179 %186 %185
+               OpSelectionMerge %191 None
+               OpBranchConditional %190 %192 %191
+        %192 = OpLabel
+               OpStore %acc %193
+               OpStore %f %uint_0
+               OpBranch %198
+        %198 = OpLabel
+               OpLoopMerge %199 %200 None
+               OpBranch %201
+        %201 = OpLabel
+        %203 = OpLoad %uint %f
+        %204 = OpAccessChain %_ptr_Uniform_uint %params %uint_0
+        %205 = OpLoad %uint %204
+        %206 = OpULessThan %bool %203 %205
+        %202 = OpLogicalNot %bool %206
+               OpSelectionMerge %207 None
+               OpBranchConditional %202 %208 %207
+        %208 = OpLabel
+               OpBranch %199
+        %207 = OpLabel
+        %209 = OpLoad %uint %f
+        %210 = OpIAdd %uint %177 %209
+        %211 = OpISub %uint %210 %50
+               OpStore %i %211
+        %213 = OpLoad %v3float %acc
+        %216 = OpAccessChain %_ptr_Uniform_uint %params %uint_0
+        %217 = OpLoad %uint %216
+        %215 = OpConvertUToF %float %217
+        %218 = OpFDiv %float %float_1 %215
+        %219 = OpLoad %uint %r_0
+        %220 = OpLoad %uint %i
+        %221 = OpAccessChain %_ptr_Workgroup_v3float %tile %219 %220
+        %222 = OpLoad %v3float %221
+        %223 = OpVectorTimesScalar %v3float %222 %218
+        %224 = OpFAdd %v3float %213 %223
+               OpStore %acc %224
+               OpBranch %200
+        %200 = OpLabel
+        %225 = OpLoad %uint %f
+        %226 = OpIAdd %uint %225 %uint_1
+               OpStore %f %226
+               OpBranch %198
+        %199 = OpLabel
+        %228 = OpLoad %14 %outputTex
+        %229 = OpLoad %v2int %writeIndex
+        %230 = OpLoad %v3float %acc
+        %231 = OpCompositeExtract %float %230 0
+        %232 = OpCompositeExtract %float %230 1
+        %233 = OpCompositeExtract %float %230 2
+        %234 = OpCompositeConstruct %v4float %231 %232 %233 %float_1
+               OpImageWrite %228 %229 %234
+               OpBranch %191
+        %191 = OpLabel
+               OpBranch %152
+        %152 = OpLabel
+        %235 = OpLoad %uint %c_0
+        %236 = OpIAdd %uint %235 %uint_1
+               OpStore %c_0 %236
+               OpBranch %150
+        %151 = OpLabel
+               OpBranch %142
+        %142 = OpLabel
+        %237 = OpLoad %uint %r_0
+        %238 = OpIAdd %uint %237 %uint_1
+               OpStore %r_0 %238
+               OpBranch %140
+        %141 = OpLabel
+               OpReturn
+               OpFunctionEnd
diff --git a/test/bug/tint/942.wgsl.expected.wgsl b/test/bug/tint/942.wgsl.expected.wgsl
new file mode 100644
index 0000000..d261497
--- /dev/null
+++ b/test/bug/tint/942.wgsl.expected.wgsl
@@ -0,0 +1,106 @@
+[[block]]
+struct Params {
+  filterDim : u32;
+  blockDim : u32;
+};
+
+[[group(0), binding(0)]] var samp : sampler;
+
+[[group(0), binding(1)]] var<uniform> params : Params;
+
+[[group(1), binding(1)]] var inputTex : texture_2d<f32>;
+
+[[group(1), binding(2)]] var outputTex : texture_storage_2d<rgba8unorm, write>;
+
+[[block]]
+struct Flip {
+  value : u32;
+};
+
+[[group(1), binding(3)]] var<uniform> flip : Flip;
+
+var<workgroup> tile : array<array<vec3<f32>, 256>, 4>;
+
+[[stage(compute), workgroup_size(64, 1, 1)]]
+fn main([[builtin(workgroup_id)]] WorkGroupID : vec3<u32>, [[builtin(local_invocation_id)]] LocalInvocationID : vec3<u32>) {
+  let filterOffset : u32 = ((params.filterDim - 1u) / 2u);
+  let dims : vec2<i32> = textureDimensions(inputTex, 0);
+  let baseIndex = (vec2<i32>(((WorkGroupID.xy * vec2<u32>(params.blockDim, 4u)) + (LocalInvocationID.xy * vec2<u32>(4u, 1u)))) - vec2<i32>(i32(filterOffset), 0));
+  {
+    var r : u32 = 0u;
+    loop {
+      if (!((r < 4u))) {
+        break;
+      }
+      {
+        var c : u32 = 0u;
+        loop {
+          if (!((c < 4u))) {
+            break;
+          }
+          var loadIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
+          if ((flip.value != 0u)) {
+            loadIndex = loadIndex.yx;
+          }
+          tile[r][((4u * LocalInvocationID.x) + c)] = textureSampleLevel(inputTex, samp, ((vec2<f32>(loadIndex) + vec2<f32>(0.25, 0.25)) / vec2<f32>(dims)), 0.0).rgb;
+
+          continuing {
+            c = (c + 1u);
+          }
+        }
+      }
+
+      continuing {
+        r = (r + 1u);
+      }
+    }
+  }
+  workgroupBarrier();
+  {
+    var r : u32 = 0u;
+    loop {
+      if (!((r < 4u))) {
+        break;
+      }
+      {
+        var c : u32 = 0u;
+        loop {
+          if (!((c < 4u))) {
+            break;
+          }
+          var writeIndex = (baseIndex + vec2<i32>(i32(c), i32(r)));
+          if ((flip.value != 0u)) {
+            writeIndex = writeIndex.yx;
+          }
+          let center : u32 = ((4u * LocalInvocationID.x) + c);
+          if ((((center >= filterOffset) && (center < (256u - filterOffset))) && all((writeIndex < dims)))) {
+            var acc : vec3<f32> = vec3<f32>(0.0, 0.0, 0.0);
+            {
+              var f : u32 = 0u;
+              loop {
+                if (!((f < params.filterDim))) {
+                  break;
+                }
+                var i : u32 = ((center + f) - filterOffset);
+                acc = (acc + ((1.0 / f32(params.filterDim)) * tile[r][i]));
+
+                continuing {
+                  f = (f + 1u);
+                }
+              }
+            }
+            textureStore(outputTex, writeIndex, vec4<f32>(acc, 1.0));
+          }
+
+          continuing {
+            c = (c + 1u);
+          }
+        }
+      }
+
+      continuing {
+        r = (r + 1u);
+      }
+    }
+  }
+}
diff --git a/test/bug/tint/943.spvasm.expected.hlsl b/test/bug/tint/943.spvasm.expected.hlsl
index a2192d2..47f46cd 100644
--- a/test/bug/tint/943.spvasm.expected.hlsl
+++ b/test/bug/tint/943.spvasm.expected.hlsl
@@ -336,10 +336,15 @@
   const uint3 gl_GlobalInvocationID_param = tint_symbol.gl_GlobalInvocationID_param;
   const uint local_invocation_index = tint_symbol.local_invocation_index;
   if ((local_invocation_index == 0u)) {
-    const float tint_symbol_6[64][64] = (float[64][64])0;
-    mm_Asub = tint_symbol_6;
-    const float tint_symbol_7[64][1] = (float[64][1])0;
-    mm_Bsub = tint_symbol_7;
+    for(int i = 0; (i < 64); i = (i + 1)) {
+      for(int i_1 = 0; (i_1 < 64); i_1 = (i_1 + 1)) {
+        mm_Asub[i][i_1] = 0.0f;
+      }
+    }
+    for(int i_2 = 0; (i_2 < 64); i_2 = (i_2 + 1)) {
+      const float tint_symbol_6[1] = (float[1])0;
+      mm_Bsub[i_2] = tint_symbol_6;
+    }
   }
   GroupMemoryBarrierWithGroupSync();
   gl_LocalInvocationID = gl_LocalInvocationID_param;