diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 0f5fad6..372e9de 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -624,11 +624,27 @@
     ast::CallExpression* expr,
     const transform::DecomposeMemoryAccess::Intrinsic* intrinsic) {
   const auto& params = expr->params();
+  auto* offset_arg = builder_.Sem().Get(params[1]);
 
-  std::string scalar_offset = UniqueIdentifier("scalar_offset");
-  {
+  uint32_t scalar_offset_value = 0;
+  std::string scalar_offset_expr;
+
+  // If true, use scalar_offset_value, otherwise use scalar_offset_expr
+  bool scalar_offset_constant = false;
+
+  if (auto val = offset_arg->ConstantValue()) {
+    TINT_ASSERT(Writer, val.Type()->Is<sem::U32>());
+    scalar_offset_value = val.Elements()[0].u32;
+    scalar_offset_value /= 4;  // bytes -> scalar index
+    scalar_offset_constant = true;
+  }
+
+  if (!scalar_offset_constant) {
+    // UBO offset not compile-time known.
+    // Calculate the scalar offset into a temporary.
+    scalar_offset_expr = UniqueIdentifier("scalar_offset");
     auto pre = line();
-    pre << "const uint " << scalar_offset << " = (";
+    pre << "const uint " << scalar_offset_expr << " = (";
     if (!EmitExpression(pre, params[1])) {  // offset
       return false;
     }
@@ -649,32 +665,37 @@
         if (!EmitExpression(out, params[0])) {  // buffer
           return false;
         }
-        out << "[" << scalar_offset << " / 4][" << scalar_offset << " % 4]";
+        if (scalar_offset_constant) {
+          char swizzle[] = {'x', 'y', 'z', 'w'};
+          out << "[" << (scalar_offset_value / 4) << "]."
+              << swizzle[scalar_offset_value & 3];
+        } else {
+          out << "[" << scalar_offset_expr << " / 4][" << scalar_offset_expr
+              << " % 4]";
+        }
         return true;
       };
       // Has a minimum alignment of 8 bytes, so is either .xy or .zw
       auto load_vec2 = [&] {
-        std::string ubo_load = UniqueIdentifier("ubo_load");
-
-        {
-          auto pre = line();
-          pre << "uint4 " << ubo_load << " = ";
-          if (!EmitExpression(pre, params[0])) {  // buffer
+        if (scalar_offset_constant) {
+          if (!EmitExpression(out, params[0])) {  // buffer
             return false;
           }
-          pre << "[" << scalar_offset << " / 4];";
+          out << "[" << (scalar_offset_value / 4) << "]";
+          out << ((scalar_offset_value & 2) == 0 ? ".xy" : ".zw");
+        } else {
+          std::string ubo_load = UniqueIdentifier("ubo_load");
+          {
+            auto pre = line();
+            pre << "uint4 " << ubo_load << " = ";
+            if (!EmitExpression(pre, params[0])) {  // buffer
+              return false;
+            }
+            pre << "[" << scalar_offset_expr << " / 4];";
+          }
+          out << "((" << scalar_offset_expr << " & 2) ? " << ubo_load
+              << ".zw : " << ubo_load << ".xy)";
         }
-
-        out << "((" << scalar_offset << " & 2) ? " << ubo_load
-            << ".zw : " << ubo_load << ".xy)";
-        return true;
-      };
-      // vec3 has a minimum alignment of 16 bytes, so is just a .xyz swizzle
-      auto load_vec3 = [&] {
-        if (!EmitExpression(out, params[0])) {  // buffer
-          return false;
-        }
-        out << "[" << scalar_offset << " / 4].xyz";
         return true;
       };
       // vec4 has a minimum alignment of 16 bytes, easiest case
@@ -682,7 +703,19 @@
         if (!EmitExpression(out, params[0])) {  // buffer
           return false;
         }
-        out << "[" << scalar_offset << " / 4]";
+        if (scalar_offset_constant) {
+          out << "[" << (scalar_offset_value / 4) << "]";
+        } else {
+          out << "[" << scalar_offset_expr << " / 4]";
+        }
+        return true;
+      };
+      // vec3 has a minimum alignment of 16 bytes, so is just a .xyz swizzle
+      auto load_vec3 = [&] {
+        if (!load_vec4()) {
+          return false;
+        }
+        out << ".xyz";
         return true;
       };
       switch (intrinsic->type) {
diff --git a/src/writer/hlsl/generator_impl_function_test.cc b/src/writer/hlsl/generator_impl_function_test.cc
index 57de37c..59ea1e7 100644
--- a/src/writer/hlsl/generator_impl_function_test.cc
+++ b/src/writer/hlsl/generator_impl_function_test.cc
@@ -335,7 +335,7 @@
            Stage(ast::PipelineStage::kFragment),
        });
 
-  GeneratorImpl& gen = Build();
+  GeneratorImpl& gen = SanitizeAndBuild();
 
   ASSERT_TRUE(gen.Generate()) << gen.error();
   EXPECT_EQ(gen.result(), R"(cbuffer cbuffer_ubo : register(b0, space1) {
@@ -343,7 +343,7 @@
 };
 
 float sub_func(float param) {
-  return ubo.coord.x;
+  return asfloat(ubo[0].x);
 }
 
 void frag_main() {
diff --git a/test/buffer/storage/dynamic_index/read.wgsl b/test/buffer/storage/dynamic_index/read.wgsl
new file mode 100644
index 0000000..8295c4d
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/read.wgsl
@@ -0,0 +1,31 @@
+struct Inner {
+    a : vec3<i32>;
+    b : i32;
+    c : vec3<u32>;
+    d : u32;
+    e : vec3<f32>;
+    f : f32;
+    g : mat2x3<f32>;
+    h : mat3x2<f32>;
+    i : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+struct S {
+    arr : array<Inner>;
+};
+
+[[binding(0), group(0)]] var<storage, read> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[builtin(local_invocation_index)]] idx : u32) {
+    let a = s.arr[idx].a;
+    let b = s.arr[idx].b;
+    let c = s.arr[idx].c;
+    let d = s.arr[idx].d;
+    let e = s.arr[idx].e;
+    let f = s.arr[idx].f;
+    let g = s.arr[idx].g;
+    let h = s.arr[idx].h;
+    let i = s.arr[idx].i;
+}
diff --git a/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl b/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl
new file mode 100644
index 0000000..89fc6b9
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/read.wgsl.expected.hlsl
@@ -0,0 +1,34 @@
+float2x3 tint_symbol_8(ByteAddressBuffer buffer, uint offset) {
+  return float2x3(asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load3((offset + 16u))));
+}
+
+float3x2 tint_symbol_10(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;
+}
+
+ByteAddressBuffer s : register(t0, space0);
+
+struct tint_symbol_1 {
+  uint idx : SV_GroupIndex;
+};
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  const uint idx = tint_symbol.idx;
+  const int3 a = asint(s.Load3((176u * idx)));
+  const int b = asint(s.Load(((176u * idx) + 12u)));
+  const uint3 c = s.Load3(((176u * idx) + 16u));
+  const uint d = s.Load(((176u * idx) + 28u));
+  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));
+  return;
+}
diff --git a/test/buffer/storage/dynamic_index/read.wgsl.expected.msl b/test/buffer/storage/dynamic_index/read.wgsl.expected.msl
new file mode 100644
index 0000000..b9aa14f
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/read.wgsl.expected.msl
@@ -0,0 +1,35 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  /* 0x0000 */ int4 arr[4];
+};
+struct Inner {
+  /* 0x0000 */ packed_int3 a;
+  /* 0x000c */ int b;
+  /* 0x0010 */ packed_uint3 c;
+  /* 0x001c */ uint d;
+  /* 0x0020 */ packed_float3 e;
+  /* 0x002c */ float f;
+  /* 0x0030 */ float2x3 g;
+  /* 0x0050 */ float3x2 h;
+  /* 0x0068 */ int8_t tint_pad[8];
+  /* 0x0070 */ tint_array_wrapper i;
+};
+struct S {
+  /* 0x0000 */ Inner arr[1];
+};
+
+kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], const device S& s [[buffer(0)]]) {
+  int3 const a = s.arr[idx].a;
+  int const b = s.arr[idx].b;
+  uint3 const c = s.arr[idx].c;
+  uint const d = s.arr[idx].d;
+  float3 const e = s.arr[idx].e;
+  float const f = s.arr[idx].f;
+  float2x3 const g = s.arr[idx].g;
+  float3x2 const h = s.arr[idx].h;
+  tint_array_wrapper const i = s.arr[idx].i;
+  return;
+}
+
diff --git a/test/buffer/storage/dynamic_index/read.wgsl.expected.spvasm b/test/buffer/storage/dynamic_index/read.wgsl.expected.spvasm
new file mode 100644
index 0000000..e5df892
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/read.wgsl.expected.spvasm
@@ -0,0 +1,114 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 68
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %tint_symbol
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "arr"
+               OpName %Inner "Inner"
+               OpMemberName %Inner 0 "a"
+               OpMemberName %Inner 1 "b"
+               OpMemberName %Inner 2 "c"
+               OpMemberName %Inner 3 "d"
+               OpMemberName %Inner 4 "e"
+               OpMemberName %Inner 5 "f"
+               OpMemberName %Inner 6 "g"
+               OpMemberName %Inner 7 "h"
+               OpMemberName %Inner 8 "i"
+               OpName %s "s"
+               OpName %tint_symbol "tint_symbol"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %Inner 0 Offset 0
+               OpMemberDecorate %Inner 1 Offset 12
+               OpMemberDecorate %Inner 2 Offset 16
+               OpMemberDecorate %Inner 3 Offset 28
+               OpMemberDecorate %Inner 4 Offset 32
+               OpMemberDecorate %Inner 5 Offset 44
+               OpMemberDecorate %Inner 6 Offset 48
+               OpMemberDecorate %Inner 6 ColMajor
+               OpMemberDecorate %Inner 6 MatrixStride 16
+               OpMemberDecorate %Inner 7 Offset 80
+               OpMemberDecorate %Inner 7 ColMajor
+               OpMemberDecorate %Inner 7 MatrixStride 8
+               OpMemberDecorate %Inner 8 Offset 112
+               OpDecorate %_arr_v4int_uint_4 ArrayStride 16
+               OpDecorate %_runtimearr_Inner ArrayStride 176
+               OpDecorate %s NonWritable
+               OpDecorate %s Binding 0
+               OpDecorate %s DescriptorSet 0
+               OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+      %v4int = OpTypeVector %int 4
+     %uint_4 = OpConstant %uint 4
+%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
+      %Inner = OpTypeStruct %v3int %int %v3uint %uint %v3float %float %mat2v3float %mat3v2float %_arr_v4int_uint_4
+%_runtimearr_Inner = OpTypeRuntimeArray %Inner
+          %S = OpTypeStruct %_runtimearr_Inner
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+          %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol = OpVariable %_ptr_Input_uint Input
+       %void = OpTypeVoid
+         %20 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3int = OpTypePointer StorageBuffer %v3int
+     %uint_1 = OpConstant %uint 1
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+     %uint_2 = OpConstant %uint 2
+%_ptr_StorageBuffer_v3uint = OpTypePointer StorageBuffer %v3uint
+     %uint_3 = OpConstant %uint 3
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+     %uint_5 = OpConstant %uint 5
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+     %uint_6 = OpConstant %uint 6
+%_ptr_StorageBuffer_mat2v3float = OpTypePointer StorageBuffer %mat2v3float
+     %uint_7 = OpConstant %uint 7
+%_ptr_StorageBuffer_mat3v2float = OpTypePointer StorageBuffer %mat3v2float
+     %uint_8 = OpConstant %uint 8
+%_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
+       %main = OpFunction %void None %20
+         %23 = OpLabel
+         %25 = OpLoad %uint %tint_symbol
+         %27 = OpAccessChain %_ptr_StorageBuffer_v3int %s %uint_0 %25 %uint_0
+         %28 = OpLoad %v3int %27
+         %29 = OpLoad %uint %tint_symbol
+         %32 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0 %29 %uint_1
+         %33 = OpLoad %int %32
+         %34 = OpLoad %uint %tint_symbol
+         %37 = OpAccessChain %_ptr_StorageBuffer_v3uint %s %uint_0 %34 %uint_2
+         %38 = OpLoad %v3uint %37
+         %39 = OpLoad %uint %tint_symbol
+         %42 = OpAccessChain %_ptr_StorageBuffer_uint %s %uint_0 %39 %uint_3
+         %43 = OpLoad %uint %42
+         %44 = OpLoad %uint %tint_symbol
+         %46 = OpAccessChain %_ptr_StorageBuffer_v3float %s %uint_0 %44 %uint_4
+         %47 = OpLoad %v3float %46
+         %48 = OpLoad %uint %tint_symbol
+         %51 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %48 %uint_5
+         %52 = OpLoad %float %51
+         %53 = OpLoad %uint %tint_symbol
+         %56 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %s %uint_0 %53 %uint_6
+         %57 = OpLoad %mat2v3float %56
+         %58 = OpLoad %uint %tint_symbol
+         %61 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %s %uint_0 %58 %uint_7
+         %62 = OpLoad %mat3v2float %61
+         %63 = OpLoad %uint %tint_symbol
+         %66 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %s %uint_0 %63 %uint_8
+         %67 = OpLoad %_arr_v4int_uint_4 %66
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/storage/dynamic_index/read.wgsl.expected.wgsl b/test/buffer/storage/dynamic_index/read.wgsl.expected.wgsl
new file mode 100644
index 0000000..3af1fa8
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/read.wgsl.expected.wgsl
@@ -0,0 +1,31 @@
+struct Inner {
+  a : vec3<i32>;
+  b : i32;
+  c : vec3<u32>;
+  d : u32;
+  e : vec3<f32>;
+  f : f32;
+  g : mat2x3<f32>;
+  h : mat3x2<f32>;
+  i : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+struct S {
+  arr : array<Inner>;
+};
+
+[[binding(0), group(0)]] var<storage, read> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[builtin(local_invocation_index)]] idx : u32) {
+  let a = s.arr[idx].a;
+  let b = s.arr[idx].b;
+  let c = s.arr[idx].c;
+  let d = s.arr[idx].d;
+  let e = s.arr[idx].e;
+  let f = s.arr[idx].f;
+  let g = s.arr[idx].g;
+  let h = s.arr[idx].h;
+  let i = s.arr[idx].i;
+}
diff --git a/test/buffer/storage/dynamic_index/write.wgsl b/test/buffer/storage/dynamic_index/write.wgsl
new file mode 100644
index 0000000..0e4b480
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/write.wgsl
@@ -0,0 +1,31 @@
+struct Inner {
+    a : vec3<i32>;
+    b : i32;
+    c : vec3<u32>;
+    d : u32;
+    e : vec3<f32>;
+    f : f32;
+    g : mat2x3<f32>;
+    h : mat3x2<f32>;
+    i : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+struct S {
+    arr : array<Inner>;
+};
+
+[[binding(0), group(0)]] var<storage, read_write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[builtin(local_invocation_index)]] idx : u32) {
+    s.arr[idx].a = vec3<i32>();
+    s.arr[idx].b = i32();
+    s.arr[idx].c = vec3<u32>();
+    s.arr[idx].d = u32();
+    s.arr[idx].e = vec3<f32>();
+    s.arr[idx].f = f32();
+    s.arr[idx].g = mat2x3<f32>();
+    s.arr[idx].h = mat3x2<f32>();
+    s.arr[idx].i = [[stride(16)]] array<vec4<i32>, 4>();
+}
diff --git a/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl b/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl
new file mode 100644
index 0000000..9d03558
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/write.wgsl.expected.hlsl
@@ -0,0 +1,39 @@
+void tint_symbol_8(RWByteAddressBuffer buffer, uint offset, float2x3 value) {
+  buffer.Store3((offset + 0u), asuint(value[0u]));
+  buffer.Store3((offset + 16u), asuint(value[1u]));
+}
+
+void tint_symbol_10(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]));
+}
+
+RWByteAddressBuffer s : register(u0, space0);
+
+struct tint_symbol_1 {
+  uint idx : SV_GroupIndex;
+};
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  const uint idx = tint_symbol.idx;
+  s.Store3((176u * idx), asuint(int3(0, 0, 0)));
+  s.Store(((176u * idx) + 12u), asuint(0));
+  s.Store3(((176u * idx) + 16u), asuint(uint3(0u, 0u, 0u)));
+  s.Store(((176u * idx) + 28u), asuint(0u));
+  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));
+  const int4 tint_symbol_13[4] = (int4[4])0;
+  tint_symbol_12(s, ((176u * idx) + 112u), tint_symbol_13);
+  return;
+}
diff --git a/test/buffer/storage/dynamic_index/write.wgsl.expected.msl b/test/buffer/storage/dynamic_index/write.wgsl.expected.msl
new file mode 100644
index 0000000..be0d4db
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/write.wgsl.expected.msl
@@ -0,0 +1,36 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  /* 0x0000 */ int4 arr[4];
+};
+struct Inner {
+  /* 0x0000 */ packed_int3 a;
+  /* 0x000c */ int b;
+  /* 0x0010 */ packed_uint3 c;
+  /* 0x001c */ uint d;
+  /* 0x0020 */ packed_float3 e;
+  /* 0x002c */ float f;
+  /* 0x0030 */ float2x3 g;
+  /* 0x0050 */ float3x2 h;
+  /* 0x0068 */ int8_t tint_pad[8];
+  /* 0x0070 */ tint_array_wrapper i;
+};
+struct S {
+  /* 0x0000 */ Inner arr[1];
+};
+
+kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], device S& s [[buffer(0)]]) {
+  s.arr[idx].a = int3();
+  s.arr[idx].b = int();
+  s.arr[idx].c = uint3();
+  s.arr[idx].d = uint();
+  s.arr[idx].e = float3();
+  s.arr[idx].f = float();
+  s.arr[idx].g = float2x3();
+  s.arr[idx].h = float3x2();
+  tint_array_wrapper const tint_symbol_2 = {.arr={}};
+  s.arr[idx].i = tint_symbol_2;
+  return;
+}
+
diff --git a/test/buffer/storage/dynamic_index/write.wgsl.expected.spvasm b/test/buffer/storage/dynamic_index/write.wgsl.expected.spvasm
new file mode 100644
index 0000000..ca9c2b3
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/write.wgsl.expected.spvasm
@@ -0,0 +1,122 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 68
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %tint_symbol
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "arr"
+               OpName %Inner "Inner"
+               OpMemberName %Inner 0 "a"
+               OpMemberName %Inner 1 "b"
+               OpMemberName %Inner 2 "c"
+               OpMemberName %Inner 3 "d"
+               OpMemberName %Inner 4 "e"
+               OpMemberName %Inner 5 "f"
+               OpMemberName %Inner 6 "g"
+               OpMemberName %Inner 7 "h"
+               OpMemberName %Inner 8 "i"
+               OpName %s "s"
+               OpName %tint_symbol "tint_symbol"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %Inner 0 Offset 0
+               OpMemberDecorate %Inner 1 Offset 12
+               OpMemberDecorate %Inner 2 Offset 16
+               OpMemberDecorate %Inner 3 Offset 28
+               OpMemberDecorate %Inner 4 Offset 32
+               OpMemberDecorate %Inner 5 Offset 44
+               OpMemberDecorate %Inner 6 Offset 48
+               OpMemberDecorate %Inner 6 ColMajor
+               OpMemberDecorate %Inner 6 MatrixStride 16
+               OpMemberDecorate %Inner 7 Offset 80
+               OpMemberDecorate %Inner 7 ColMajor
+               OpMemberDecorate %Inner 7 MatrixStride 8
+               OpMemberDecorate %Inner 8 Offset 112
+               OpDecorate %_arr_v4int_uint_4 ArrayStride 16
+               OpDecorate %_runtimearr_Inner ArrayStride 176
+               OpDecorate %s Binding 0
+               OpDecorate %s DescriptorSet 0
+               OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+      %v4int = OpTypeVector %int 4
+     %uint_4 = OpConstant %uint 4
+%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
+      %Inner = OpTypeStruct %v3int %int %v3uint %uint %v3float %float %mat2v3float %mat3v2float %_arr_v4int_uint_4
+%_runtimearr_Inner = OpTypeRuntimeArray %Inner
+          %S = OpTypeStruct %_runtimearr_Inner
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+          %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol = OpVariable %_ptr_Input_uint Input
+       %void = OpTypeVoid
+         %20 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3int = OpTypePointer StorageBuffer %v3int
+         %28 = OpConstantNull %v3int
+     %uint_1 = OpConstant %uint 1
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+         %33 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+%_ptr_StorageBuffer_v3uint = OpTypePointer StorageBuffer %v3uint
+         %38 = OpConstantNull %v3uint
+     %uint_3 = OpConstant %uint 3
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+         %43 = OpConstantNull %uint
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+         %47 = OpConstantNull %v3float
+     %uint_5 = OpConstant %uint 5
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+         %52 = OpConstantNull %float
+     %uint_6 = OpConstant %uint 6
+%_ptr_StorageBuffer_mat2v3float = OpTypePointer StorageBuffer %mat2v3float
+         %57 = OpConstantNull %mat2v3float
+     %uint_7 = OpConstant %uint 7
+%_ptr_StorageBuffer_mat3v2float = OpTypePointer StorageBuffer %mat3v2float
+         %62 = OpConstantNull %mat3v2float
+     %uint_8 = OpConstant %uint 8
+%_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
+         %67 = OpConstantNull %_arr_v4int_uint_4
+       %main = OpFunction %void None %20
+         %23 = OpLabel
+         %25 = OpLoad %uint %tint_symbol
+         %27 = OpAccessChain %_ptr_StorageBuffer_v3int %s %uint_0 %25 %uint_0
+               OpStore %27 %28
+         %29 = OpLoad %uint %tint_symbol
+         %32 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_0 %29 %uint_1
+               OpStore %32 %33
+         %34 = OpLoad %uint %tint_symbol
+         %37 = OpAccessChain %_ptr_StorageBuffer_v3uint %s %uint_0 %34 %uint_2
+               OpStore %37 %38
+         %39 = OpLoad %uint %tint_symbol
+         %42 = OpAccessChain %_ptr_StorageBuffer_uint %s %uint_0 %39 %uint_3
+               OpStore %42 %43
+         %44 = OpLoad %uint %tint_symbol
+         %46 = OpAccessChain %_ptr_StorageBuffer_v3float %s %uint_0 %44 %uint_4
+               OpStore %46 %47
+         %48 = OpLoad %uint %tint_symbol
+         %51 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %48 %uint_5
+               OpStore %51 %52
+         %53 = OpLoad %uint %tint_symbol
+         %56 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %s %uint_0 %53 %uint_6
+               OpStore %56 %57
+         %58 = OpLoad %uint %tint_symbol
+         %61 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %s %uint_0 %58 %uint_7
+               OpStore %61 %62
+         %63 = OpLoad %uint %tint_symbol
+         %66 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %s %uint_0 %63 %uint_8
+               OpStore %66 %67
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/storage/dynamic_index/write.wgsl.expected.wgsl b/test/buffer/storage/dynamic_index/write.wgsl.expected.wgsl
new file mode 100644
index 0000000..d6aea47
--- /dev/null
+++ b/test/buffer/storage/dynamic_index/write.wgsl.expected.wgsl
@@ -0,0 +1,31 @@
+struct Inner {
+  a : vec3<i32>;
+  b : i32;
+  c : vec3<u32>;
+  d : u32;
+  e : vec3<f32>;
+  f : f32;
+  g : mat2x3<f32>;
+  h : mat3x2<f32>;
+  i : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+struct S {
+  arr : array<Inner>;
+};
+
+[[binding(0), group(0)]] var<storage, read_write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[builtin(local_invocation_index)]] idx : u32) {
+  s.arr[idx].a = vec3<i32>();
+  s.arr[idx].b = i32();
+  s.arr[idx].c = vec3<u32>();
+  s.arr[idx].d = u32();
+  s.arr[idx].e = vec3<f32>();
+  s.arr[idx].f = f32();
+  s.arr[idx].g = mat2x3<f32>();
+  s.arr[idx].h = mat3x2<f32>();
+  s.arr[idx].i = [[stride(16)]] array<vec4<i32>, 4>();
+}
diff --git a/test/buffer/storage/static_index/read.wgsl b/test/buffer/storage/static_index/read.wgsl
new file mode 100644
index 0000000..096492d
--- /dev/null
+++ b/test/buffer/storage/static_index/read.wgsl
@@ -0,0 +1,33 @@
+struct Inner {
+    x : i32;
+};
+
+[[block]]
+struct S {
+    a : vec3<i32>;
+    b : i32;
+    c : vec3<u32>;
+    d : u32;
+    e : vec3<f32>;
+    f : f32;
+    g : mat2x3<f32>;
+    h : mat3x2<f32>;
+    i : Inner;
+    j : [[stride(16)]] array<Inner, 4>;
+};
+
+[[binding(0), group(0)]] var<storage, read> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    let a = s.a;
+    let b = s.b;
+    let c = s.c;
+    let d = s.d;
+    let e = s.e;
+    let f = s.f;
+    let g = s.g;
+    let h = s.h;
+    let i = s.i;
+    let j = s.j;
+}
diff --git a/test/buffer/storage/static_index/read.wgsl.expected.hlsl b/test/buffer/storage/static_index/read.wgsl.expected.hlsl
new file mode 100644
index 0000000..95ed272
--- /dev/null
+++ b/test/buffer/storage/static_index/read.wgsl.expected.hlsl
@@ -0,0 +1,42 @@
+struct Inner {
+  int x;
+};
+struct tint_padded_array_element {
+  Inner el;
+};
+
+float2x3 tint_symbol_6(ByteAddressBuffer buffer, uint offset) {
+  return float2x3(asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load3((offset + 16u))));
+}
+
+float3x2 tint_symbol_8(ByteAddressBuffer buffer, uint offset) {
+  return float3x2(asfloat(buffer.Load2((offset + 0u))), asfloat(buffer.Load2((offset + 8u))), asfloat(buffer.Load2((offset + 16u))));
+}
+
+Inner tint_symbol_9(ByteAddressBuffer buffer, uint offset) {
+  const Inner tint_symbol_11 = {asint(buffer.Load((offset + 0u)))};
+  return tint_symbol_11;
+}
+
+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;
+}
+
+ByteAddressBuffer s : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  const int3 a = asint(s.Load3(0u));
+  const int b = asint(s.Load(12u));
+  const uint3 c = s.Load3(16u);
+  const uint d = s.Load(28u);
+  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 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/read.wgsl.expected.msl b/test/buffer/storage/static_index/read.wgsl.expected.msl
new file mode 100644
index 0000000..316ed31
--- /dev/null
+++ b/test/buffer/storage/static_index/read.wgsl.expected.msl
@@ -0,0 +1,41 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Inner {
+  /* 0x0000 */ int x;
+};
+struct tint_padded_array_element {
+  /* 0x0000 */ Inner el;
+  /* 0x0004 */ int8_t tint_pad[12];
+};
+struct tint_array_wrapper {
+  /* 0x0000 */ tint_padded_array_element arr[4];
+};
+struct S {
+  /* 0x0000 */ packed_int3 a;
+  /* 0x000c */ int b;
+  /* 0x0010 */ packed_uint3 c;
+  /* 0x001c */ uint d;
+  /* 0x0020 */ packed_float3 e;
+  /* 0x002c */ float f;
+  /* 0x0030 */ float2x3 g;
+  /* 0x0050 */ float3x2 h;
+  /* 0x0068 */ Inner i;
+  /* 0x006c */ tint_array_wrapper j;
+  /* 0x00ac */ int8_t tint_pad_1[4];
+};
+
+kernel void tint_symbol(const device S& s [[buffer(0)]]) {
+  int3 const a = s.a;
+  int const b = s.b;
+  uint3 const c = s.c;
+  uint const d = s.d;
+  float3 const e = s.e;
+  float const f = s.f;
+  float2x3 const g = s.g;
+  float3x2 const h = s.h;
+  Inner const i = s.i;
+  tint_array_wrapper const j = s.j;
+  return;
+}
+
diff --git a/test/buffer/storage/static_index/read.wgsl.expected.spvasm b/test/buffer/storage/static_index/read.wgsl.expected.spvasm
new file mode 100644
index 0000000..98bf962
--- /dev/null
+++ b/test/buffer/storage/static_index/read.wgsl.expected.spvasm
@@ -0,0 +1,104 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 59
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "a"
+               OpMemberName %S 1 "b"
+               OpMemberName %S 2 "c"
+               OpMemberName %S 3 "d"
+               OpMemberName %S 4 "e"
+               OpMemberName %S 5 "f"
+               OpMemberName %S 6 "g"
+               OpMemberName %S 7 "h"
+               OpMemberName %S 8 "i"
+               OpName %Inner "Inner"
+               OpMemberName %Inner 0 "x"
+               OpMemberName %S 9 "j"
+               OpName %s "s"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 12
+               OpMemberDecorate %S 2 Offset 16
+               OpMemberDecorate %S 3 Offset 28
+               OpMemberDecorate %S 4 Offset 32
+               OpMemberDecorate %S 5 Offset 44
+               OpMemberDecorate %S 6 Offset 48
+               OpMemberDecorate %S 6 ColMajor
+               OpMemberDecorate %S 6 MatrixStride 16
+               OpMemberDecorate %S 7 Offset 80
+               OpMemberDecorate %S 7 ColMajor
+               OpMemberDecorate %S 7 MatrixStride 8
+               OpMemberDecorate %S 8 Offset 104
+               OpMemberDecorate %Inner 0 Offset 0
+               OpMemberDecorate %S 9 Offset 108
+               OpDecorate %_arr_Inner_uint_4 ArrayStride 16
+               OpDecorate %s NonWritable
+               OpDecorate %s Binding 0
+               OpDecorate %s DescriptorSet 0
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+      %Inner = OpTypeStruct %int
+     %uint_4 = OpConstant %uint 4
+%_arr_Inner_uint_4 = OpTypeArray %Inner %uint_4
+          %S = OpTypeStruct %v3int %int %v3uint %uint %v3float %float %mat2v3float %mat3v2float %Inner %_arr_Inner_uint_4
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+          %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3int = OpTypePointer StorageBuffer %v3int
+     %uint_1 = OpConstant %uint 1
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+     %uint_2 = OpConstant %uint 2
+%_ptr_StorageBuffer_v3uint = OpTypePointer StorageBuffer %v3uint
+     %uint_3 = OpConstant %uint 3
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+     %uint_5 = OpConstant %uint 5
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+     %uint_6 = OpConstant %uint 6
+%_ptr_StorageBuffer_mat2v3float = OpTypePointer StorageBuffer %mat2v3float
+     %uint_7 = OpConstant %uint 7
+%_ptr_StorageBuffer_mat3v2float = OpTypePointer StorageBuffer %mat3v2float
+     %uint_8 = OpConstant %uint 8
+%_ptr_StorageBuffer_Inner = OpTypePointer StorageBuffer %Inner
+     %uint_9 = OpConstant %uint 9
+%_ptr_StorageBuffer__arr_Inner_uint_4 = OpTypePointer StorageBuffer %_arr_Inner_uint_4
+       %main = OpFunction %void None %16
+         %19 = OpLabel
+         %22 = OpAccessChain %_ptr_StorageBuffer_v3int %s %uint_0
+         %23 = OpLoad %v3int %22
+         %26 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_1
+         %27 = OpLoad %int %26
+         %30 = OpAccessChain %_ptr_StorageBuffer_v3uint %s %uint_2
+         %31 = OpLoad %v3uint %30
+         %34 = OpAccessChain %_ptr_StorageBuffer_uint %s %uint_3
+         %35 = OpLoad %uint %34
+         %37 = OpAccessChain %_ptr_StorageBuffer_v3float %s %uint_4
+         %38 = OpLoad %v3float %37
+         %41 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_5
+         %42 = OpLoad %float %41
+         %45 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %s %uint_6
+         %46 = OpLoad %mat2v3float %45
+         %49 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %s %uint_7
+         %50 = OpLoad %mat3v2float %49
+         %53 = OpAccessChain %_ptr_StorageBuffer_Inner %s %uint_8
+         %54 = OpLoad %Inner %53
+         %57 = OpAccessChain %_ptr_StorageBuffer__arr_Inner_uint_4 %s %uint_9
+         %58 = OpLoad %_arr_Inner_uint_4 %57
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/storage/static_index/read.wgsl.expected.wgsl b/test/buffer/storage/static_index/read.wgsl.expected.wgsl
new file mode 100644
index 0000000..5dcc611
--- /dev/null
+++ b/test/buffer/storage/static_index/read.wgsl.expected.wgsl
@@ -0,0 +1,33 @@
+struct Inner {
+  x : i32;
+};
+
+[[block]]
+struct S {
+  a : vec3<i32>;
+  b : i32;
+  c : vec3<u32>;
+  d : u32;
+  e : vec3<f32>;
+  f : f32;
+  g : mat2x3<f32>;
+  h : mat3x2<f32>;
+  i : Inner;
+  j : [[stride(16)]] array<Inner, 4>;
+};
+
+[[binding(0), group(0)]] var<storage, read> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let a = s.a;
+  let b = s.b;
+  let c = s.c;
+  let d = s.d;
+  let e = s.e;
+  let f = s.f;
+  let g = s.g;
+  let h = s.h;
+  let i = s.i;
+  let j = s.j;
+}
diff --git a/test/buffer/storage/static_index/write.wgsl b/test/buffer/storage/static_index/write.wgsl
new file mode 100644
index 0000000..6b6affb
--- /dev/null
+++ b/test/buffer/storage/static_index/write.wgsl
@@ -0,0 +1,33 @@
+struct Inner {
+    x : i32;
+};
+
+[[block]]
+struct S {
+    a : vec3<i32>;
+    b : i32;
+    c : vec3<u32>;
+    d : u32;
+    e : vec3<f32>;
+    f : f32;
+    g : mat2x3<f32>;
+    h : mat3x2<f32>;
+    i : Inner;
+    j : [[stride(16)]] array<Inner, 4>;
+};
+
+[[binding(0), group(0)]] var<storage, write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    s.a = vec3<i32>();
+    s.b = i32();
+    s.c = vec3<u32>();
+    s.d = u32();
+    s.e = vec3<f32>();
+    s.f = f32();
+    s.g = mat2x3<f32>();
+    s.h = mat3x2<f32>();
+    s.i = Inner();
+    s.j = [[stride(16)]] array<Inner, 4>();
+}
diff --git a/test/buffer/storage/static_index/write.wgsl.expected.hlsl b/test/buffer/storage/static_index/write.wgsl.expected.hlsl
new file mode 100644
index 0000000..52c36c6
--- /dev/null
+++ b/test/buffer/storage/static_index/write.wgsl.expected.hlsl
@@ -0,0 +1,47 @@
+struct Inner {
+  int x;
+};
+struct tint_padded_array_element {
+  Inner el;
+};
+
+void tint_symbol_6(RWByteAddressBuffer buffer, uint offset, float2x3 value) {
+  buffer.Store3((offset + 0u), asuint(value[0u]));
+  buffer.Store3((offset + 16u), asuint(value[1u]));
+}
+
+void tint_symbol_8(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_9(RWByteAddressBuffer buffer, uint offset, Inner value) {
+  buffer.Store((offset + 0u), asuint(value.x));
+}
+
+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);
+}
+
+RWByteAddressBuffer s : register(u0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  s.Store3(0u, asuint(int3(0, 0, 0)));
+  s.Store(12u, asuint(0));
+  s.Store3(16u, asuint(uint3(0u, 0u, 0u)));
+  s.Store(28u, asuint(0u));
+  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));
+  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;
+  tint_symbol_10(s, 108u, tint_symbol_12);
+  return;
+}
diff --git a/test/buffer/storage/static_index/write.wgsl.expected.msl b/test/buffer/storage/static_index/write.wgsl.expected.msl
new file mode 100644
index 0000000..eab1f4f
--- /dev/null
+++ b/test/buffer/storage/static_index/write.wgsl.expected.msl
@@ -0,0 +1,43 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Inner {
+  /* 0x0000 */ int x;
+};
+struct tint_padded_array_element {
+  /* 0x0000 */ Inner el;
+  /* 0x0004 */ int8_t tint_pad[12];
+};
+struct tint_array_wrapper {
+  /* 0x0000 */ tint_padded_array_element arr[4];
+};
+struct S {
+  /* 0x0000 */ packed_int3 a;
+  /* 0x000c */ int b;
+  /* 0x0010 */ packed_uint3 c;
+  /* 0x001c */ uint d;
+  /* 0x0020 */ packed_float3 e;
+  /* 0x002c */ float f;
+  /* 0x0030 */ float2x3 g;
+  /* 0x0050 */ float3x2 h;
+  /* 0x0068 */ Inner i;
+  /* 0x006c */ tint_array_wrapper j;
+  /* 0x00ac */ int8_t tint_pad_1[4];
+};
+
+kernel void tint_symbol(device S& s [[buffer(0)]]) {
+  s.a = int3();
+  s.b = int();
+  s.c = uint3();
+  s.d = uint();
+  s.e = float3();
+  s.f = float();
+  s.g = float2x3();
+  s.h = float3x2();
+  Inner const tint_symbol_1 = {};
+  s.i = tint_symbol_1;
+  tint_array_wrapper const tint_symbol_2 = {.arr={}};
+  s.j = tint_symbol_2;
+  return;
+}
+
diff --git a/test/buffer/storage/static_index/write.wgsl.expected.spvasm b/test/buffer/storage/static_index/write.wgsl.expected.spvasm
new file mode 100644
index 0000000..faa44c5
--- /dev/null
+++ b/test/buffer/storage/static_index/write.wgsl.expected.spvasm
@@ -0,0 +1,114 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 59
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "a"
+               OpMemberName %S 1 "b"
+               OpMemberName %S 2 "c"
+               OpMemberName %S 3 "d"
+               OpMemberName %S 4 "e"
+               OpMemberName %S 5 "f"
+               OpMemberName %S 6 "g"
+               OpMemberName %S 7 "h"
+               OpMemberName %S 8 "i"
+               OpName %Inner "Inner"
+               OpMemberName %Inner 0 "x"
+               OpMemberName %S 9 "j"
+               OpName %s "s"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 12
+               OpMemberDecorate %S 2 Offset 16
+               OpMemberDecorate %S 3 Offset 28
+               OpMemberDecorate %S 4 Offset 32
+               OpMemberDecorate %S 5 Offset 44
+               OpMemberDecorate %S 6 Offset 48
+               OpMemberDecorate %S 6 ColMajor
+               OpMemberDecorate %S 6 MatrixStride 16
+               OpMemberDecorate %S 7 Offset 80
+               OpMemberDecorate %S 7 ColMajor
+               OpMemberDecorate %S 7 MatrixStride 8
+               OpMemberDecorate %S 8 Offset 104
+               OpMemberDecorate %Inner 0 Offset 0
+               OpMemberDecorate %S 9 Offset 108
+               OpDecorate %_arr_Inner_uint_4 ArrayStride 16
+               OpDecorate %s NonReadable
+               OpDecorate %s Binding 0
+               OpDecorate %s DescriptorSet 0
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+      %Inner = OpTypeStruct %int
+     %uint_4 = OpConstant %uint 4
+%_arr_Inner_uint_4 = OpTypeArray %Inner %uint_4
+          %S = OpTypeStruct %v3int %int %v3uint %uint %v3float %float %mat2v3float %mat3v2float %Inner %_arr_Inner_uint_4
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+          %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+       %void = OpTypeVoid
+         %16 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3int = OpTypePointer StorageBuffer %v3int
+         %23 = OpConstantNull %v3int
+     %uint_1 = OpConstant %uint 1
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+         %27 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+%_ptr_StorageBuffer_v3uint = OpTypePointer StorageBuffer %v3uint
+         %31 = OpConstantNull %v3uint
+     %uint_3 = OpConstant %uint 3
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+         %35 = OpConstantNull %uint
+%_ptr_StorageBuffer_v3float = OpTypePointer StorageBuffer %v3float
+         %38 = OpConstantNull %v3float
+     %uint_5 = OpConstant %uint 5
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+         %42 = OpConstantNull %float
+     %uint_6 = OpConstant %uint 6
+%_ptr_StorageBuffer_mat2v3float = OpTypePointer StorageBuffer %mat2v3float
+         %46 = OpConstantNull %mat2v3float
+     %uint_7 = OpConstant %uint 7
+%_ptr_StorageBuffer_mat3v2float = OpTypePointer StorageBuffer %mat3v2float
+         %50 = OpConstantNull %mat3v2float
+     %uint_8 = OpConstant %uint 8
+%_ptr_StorageBuffer_Inner = OpTypePointer StorageBuffer %Inner
+         %54 = OpConstantNull %Inner
+     %uint_9 = OpConstant %uint 9
+%_ptr_StorageBuffer__arr_Inner_uint_4 = OpTypePointer StorageBuffer %_arr_Inner_uint_4
+         %58 = OpConstantNull %_arr_Inner_uint_4
+       %main = OpFunction %void None %16
+         %19 = OpLabel
+         %22 = OpAccessChain %_ptr_StorageBuffer_v3int %s %uint_0
+               OpStore %22 %23
+         %26 = OpAccessChain %_ptr_StorageBuffer_int %s %uint_1
+               OpStore %26 %27
+         %30 = OpAccessChain %_ptr_StorageBuffer_v3uint %s %uint_2
+               OpStore %30 %31
+         %34 = OpAccessChain %_ptr_StorageBuffer_uint %s %uint_3
+               OpStore %34 %35
+         %37 = OpAccessChain %_ptr_StorageBuffer_v3float %s %uint_4
+               OpStore %37 %38
+         %41 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_5
+               OpStore %41 %42
+         %45 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %s %uint_6
+               OpStore %45 %46
+         %49 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %s %uint_7
+               OpStore %49 %50
+         %53 = OpAccessChain %_ptr_StorageBuffer_Inner %s %uint_8
+               OpStore %53 %54
+         %57 = OpAccessChain %_ptr_StorageBuffer__arr_Inner_uint_4 %s %uint_9
+               OpStore %57 %58
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/storage/static_index/write.wgsl.expected.wgsl b/test/buffer/storage/static_index/write.wgsl.expected.wgsl
new file mode 100644
index 0000000..16725de
--- /dev/null
+++ b/test/buffer/storage/static_index/write.wgsl.expected.wgsl
@@ -0,0 +1,33 @@
+struct Inner {
+  x : i32;
+};
+
+[[block]]
+struct S {
+  a : vec3<i32>;
+  b : i32;
+  c : vec3<u32>;
+  d : u32;
+  e : vec3<f32>;
+  f : f32;
+  g : mat2x3<f32>;
+  h : mat3x2<f32>;
+  i : Inner;
+  j : [[stride(16)]] array<Inner, 4>;
+};
+
+[[binding(0), group(0)]] var<storage, write> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  s.a = vec3<i32>();
+  s.b = i32();
+  s.c = vec3<u32>();
+  s.d = u32();
+  s.e = vec3<f32>();
+  s.f = f32();
+  s.g = mat2x3<f32>();
+  s.h = mat3x2<f32>();
+  s.i = Inner();
+  s.j = [[stride(16)]] array<Inner, 4>();
+}
diff --git a/test/buffer/uniform/dynamic_index/read.wgsl b/test/buffer/uniform/dynamic_index/read.wgsl
new file mode 100644
index 0000000..12cc0c0
--- /dev/null
+++ b/test/buffer/uniform/dynamic_index/read.wgsl
@@ -0,0 +1,33 @@
+struct Inner {
+    a : vec3<i32>;
+    b : i32;
+    c : vec3<u32>;
+    d : u32;
+    e : vec3<f32>;
+    f : f32;
+    g : vec2<i32>;
+    h : vec2<i32>;
+    i : mat2x3<f32>;
+    [[align(16)]] j : mat3x2<f32>;
+    [[align(16)]] k : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+struct S {
+    arr : array<Inner, 8>;
+};
+
+[[binding(0), group(0)]] var<uniform> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[builtin(local_invocation_index)]] idx : u32) {
+    let a = s.arr[idx].a;
+    let b = s.arr[idx].b;
+    let c = s.arr[idx].c;
+    let d = s.arr[idx].d;
+    let e = s.arr[idx].e;
+    let f = s.arr[idx].f;
+    let g = s.arr[idx].g;
+    let h = s.arr[idx].h;
+    let i = s.arr[idx].i;
+}
diff --git a/test/buffer/uniform/dynamic_index/read.wgsl.expected.hlsl b/test/buffer/uniform/dynamic_index/read.wgsl.expected.hlsl
new file mode 100644
index 0000000..378e877
--- /dev/null
+++ b/test/buffer/uniform/dynamic_index/read.wgsl.expected.hlsl
@@ -0,0 +1,38 @@
+float2x3 tint_symbol_9(uint4 buffer[96], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const uint scalar_offset_1 = ((offset + 16u)) / 4;
+  return float2x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz));
+}
+
+cbuffer cbuffer_s : register(b0, space0) {
+  uint4 s[96];
+};
+
+struct tint_symbol_1 {
+  uint idx : SV_GroupIndex;
+};
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  const uint idx = tint_symbol.idx;
+  const uint scalar_offset_2 = ((192u * idx)) / 4;
+  const int3 a = asint(s[scalar_offset_2 / 4].xyz);
+  const uint scalar_offset_3 = (((192u * idx) + 12u)) / 4;
+  const int b = asint(s[scalar_offset_3 / 4][scalar_offset_3 % 4]);
+  const uint scalar_offset_4 = (((192u * idx) + 16u)) / 4;
+  const uint3 c = s[scalar_offset_4 / 4].xyz;
+  const uint scalar_offset_5 = (((192u * idx) + 28u)) / 4;
+  const uint d = s[scalar_offset_5 / 4][scalar_offset_5 % 4];
+  const uint scalar_offset_6 = (((192u * idx) + 32u)) / 4;
+  const float3 e = asfloat(s[scalar_offset_6 / 4].xyz);
+  const uint scalar_offset_7 = (((192u * idx) + 44u)) / 4;
+  const float f = asfloat(s[scalar_offset_7 / 4][scalar_offset_7 % 4]);
+  const uint scalar_offset_8 = (((192u * idx) + 48u)) / 4;
+  uint4 ubo_load = s[scalar_offset_8 / 4];
+  const int2 g = asint(((scalar_offset_8 & 2) ? ubo_load.zw : ubo_load.xy));
+  const uint scalar_offset_9 = (((192u * idx) + 56u)) / 4;
+  uint4 ubo_load_1 = s[scalar_offset_9 / 4];
+  const int2 h = asint(((scalar_offset_9 & 2) ? ubo_load_1.zw : ubo_load_1.xy));
+  const float2x3 i = tint_symbol_9(s, ((192u * idx) + 64u));
+  return;
+}
diff --git a/test/buffer/uniform/dynamic_index/read.wgsl.expected.msl b/test/buffer/uniform/dynamic_index/read.wgsl.expected.msl
new file mode 100644
index 0000000..08c4e7d
--- /dev/null
+++ b/test/buffer/uniform/dynamic_index/read.wgsl.expected.msl
@@ -0,0 +1,40 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  /* 0x0000 */ int4 arr[4];
+};
+struct Inner {
+  /* 0x0000 */ packed_int3 a;
+  /* 0x000c */ int b;
+  /* 0x0010 */ packed_uint3 c;
+  /* 0x001c */ uint d;
+  /* 0x0020 */ packed_float3 e;
+  /* 0x002c */ float f;
+  /* 0x0030 */ packed_int2 g;
+  /* 0x0038 */ packed_int2 h;
+  /* 0x0040 */ float2x3 i;
+  /* 0x0060 */ float3x2 j;
+  /* 0x0078 */ int8_t tint_pad[8];
+  /* 0x0080 */ tint_array_wrapper k;
+};
+struct tint_array_wrapper_1 {
+  /* 0x0000 */ Inner arr[8];
+};
+struct S {
+  /* 0x0000 */ tint_array_wrapper_1 arr;
+};
+
+kernel void tint_symbol(uint idx [[thread_index_in_threadgroup]], constant S& s [[buffer(0)]]) {
+  int3 const a = s.arr.arr[idx].a;
+  int const b = s.arr.arr[idx].b;
+  uint3 const c = s.arr.arr[idx].c;
+  uint const d = s.arr.arr[idx].d;
+  float3 const e = s.arr.arr[idx].e;
+  float const f = s.arr.arr[idx].f;
+  int2 const g = s.arr.arr[idx].g;
+  int2 const h = s.arr.arr[idx].h;
+  float2x3 const i = s.arr.arr[idx].i;
+  return;
+}
+
diff --git a/test/buffer/uniform/dynamic_index/read.wgsl.expected.spvasm b/test/buffer/uniform/dynamic_index/read.wgsl.expected.spvasm
new file mode 100644
index 0000000..4ea6207
--- /dev/null
+++ b/test/buffer/uniform/dynamic_index/read.wgsl.expected.spvasm
@@ -0,0 +1,118 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 68
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %tint_symbol
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "arr"
+               OpName %Inner "Inner"
+               OpMemberName %Inner 0 "a"
+               OpMemberName %Inner 1 "b"
+               OpMemberName %Inner 2 "c"
+               OpMemberName %Inner 3 "d"
+               OpMemberName %Inner 4 "e"
+               OpMemberName %Inner 5 "f"
+               OpMemberName %Inner 6 "g"
+               OpMemberName %Inner 7 "h"
+               OpMemberName %Inner 8 "i"
+               OpMemberName %Inner 9 "j"
+               OpMemberName %Inner 10 "k"
+               OpName %s "s"
+               OpName %tint_symbol "tint_symbol"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %Inner 0 Offset 0
+               OpMemberDecorate %Inner 1 Offset 12
+               OpMemberDecorate %Inner 2 Offset 16
+               OpMemberDecorate %Inner 3 Offset 28
+               OpMemberDecorate %Inner 4 Offset 32
+               OpMemberDecorate %Inner 5 Offset 44
+               OpMemberDecorate %Inner 6 Offset 48
+               OpMemberDecorate %Inner 7 Offset 56
+               OpMemberDecorate %Inner 8 Offset 64
+               OpMemberDecorate %Inner 8 ColMajor
+               OpMemberDecorate %Inner 8 MatrixStride 16
+               OpMemberDecorate %Inner 9 Offset 96
+               OpMemberDecorate %Inner 9 ColMajor
+               OpMemberDecorate %Inner 9 MatrixStride 8
+               OpMemberDecorate %Inner 10 Offset 128
+               OpDecorate %_arr_v4int_uint_4 ArrayStride 16
+               OpDecorate %_arr_Inner_uint_8 ArrayStride 192
+               OpDecorate %s NonWritable
+               OpDecorate %s Binding 0
+               OpDecorate %s DescriptorSet 0
+               OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+      %v2int = OpTypeVector %int 2
+%mat2v3float = OpTypeMatrix %v3float 2
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+      %v4int = OpTypeVector %int 4
+     %uint_4 = OpConstant %uint 4
+%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
+      %Inner = OpTypeStruct %v3int %int %v3uint %uint %v3float %float %v2int %v2int %mat2v3float %mat3v2float %_arr_v4int_uint_4
+     %uint_8 = OpConstant %uint 8
+%_arr_Inner_uint_8 = OpTypeArray %Inner %uint_8
+          %S = OpTypeStruct %_arr_Inner_uint_8
+%_ptr_Uniform_S = OpTypePointer Uniform %S
+          %s = OpVariable %_ptr_Uniform_S Uniform
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol = OpVariable %_ptr_Input_uint Input
+       %void = OpTypeVoid
+         %22 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_v3int = OpTypePointer Uniform %v3int
+     %uint_1 = OpConstant %uint 1
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+     %uint_2 = OpConstant %uint 2
+%_ptr_Uniform_v3uint = OpTypePointer Uniform %v3uint
+     %uint_3 = OpConstant %uint 3
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
+     %uint_5 = OpConstant %uint 5
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+     %uint_6 = OpConstant %uint 6
+%_ptr_Uniform_v2int = OpTypePointer Uniform %v2int
+     %uint_7 = OpConstant %uint 7
+%_ptr_Uniform_mat2v3float = OpTypePointer Uniform %mat2v3float
+       %main = OpFunction %void None %22
+         %25 = OpLabel
+         %27 = OpLoad %uint %tint_symbol
+         %29 = OpAccessChain %_ptr_Uniform_v3int %s %uint_0 %27 %uint_0
+         %30 = OpLoad %v3int %29
+         %31 = OpLoad %uint %tint_symbol
+         %34 = OpAccessChain %_ptr_Uniform_int %s %uint_0 %31 %uint_1
+         %35 = OpLoad %int %34
+         %36 = OpLoad %uint %tint_symbol
+         %39 = OpAccessChain %_ptr_Uniform_v3uint %s %uint_0 %36 %uint_2
+         %40 = OpLoad %v3uint %39
+         %41 = OpLoad %uint %tint_symbol
+         %44 = OpAccessChain %_ptr_Uniform_uint %s %uint_0 %41 %uint_3
+         %45 = OpLoad %uint %44
+         %46 = OpLoad %uint %tint_symbol
+         %48 = OpAccessChain %_ptr_Uniform_v3float %s %uint_0 %46 %uint_4
+         %49 = OpLoad %v3float %48
+         %50 = OpLoad %uint %tint_symbol
+         %53 = OpAccessChain %_ptr_Uniform_float %s %uint_0 %50 %uint_5
+         %54 = OpLoad %float %53
+         %55 = OpLoad %uint %tint_symbol
+         %58 = OpAccessChain %_ptr_Uniform_v2int %s %uint_0 %55 %uint_6
+         %59 = OpLoad %v2int %58
+         %60 = OpLoad %uint %tint_symbol
+         %62 = OpAccessChain %_ptr_Uniform_v2int %s %uint_0 %60 %uint_7
+         %63 = OpLoad %v2int %62
+         %64 = OpLoad %uint %tint_symbol
+         %66 = OpAccessChain %_ptr_Uniform_mat2v3float %s %uint_0 %64 %uint_8
+         %67 = OpLoad %mat2v3float %66
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/uniform/dynamic_index/read.wgsl.expected.wgsl b/test/buffer/uniform/dynamic_index/read.wgsl.expected.wgsl
new file mode 100644
index 0000000..7377f9e
--- /dev/null
+++ b/test/buffer/uniform/dynamic_index/read.wgsl.expected.wgsl
@@ -0,0 +1,35 @@
+struct Inner {
+  a : vec3<i32>;
+  b : i32;
+  c : vec3<u32>;
+  d : u32;
+  e : vec3<f32>;
+  f : f32;
+  g : vec2<i32>;
+  h : vec2<i32>;
+  i : mat2x3<f32>;
+  [[align(16)]]
+  j : mat3x2<f32>;
+  [[align(16)]]
+  k : [[stride(16)]] array<vec4<i32>, 4>;
+};
+
+[[block]]
+struct S {
+  arr : array<Inner, 8>;
+};
+
+[[binding(0), group(0)]] var<uniform> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[builtin(local_invocation_index)]] idx : u32) {
+  let a = s.arr[idx].a;
+  let b = s.arr[idx].b;
+  let c = s.arr[idx].c;
+  let d = s.arr[idx].d;
+  let e = s.arr[idx].e;
+  let f = s.arr[idx].f;
+  let g = s.arr[idx].g;
+  let h = s.arr[idx].h;
+  let i = s.arr[idx].i;
+}
diff --git a/test/buffer/uniform/static_index/read.wgsl b/test/buffer/uniform/static_index/read.wgsl
new file mode 100644
index 0000000..c3ad9ad
--- /dev/null
+++ b/test/buffer/uniform/static_index/read.wgsl
@@ -0,0 +1,37 @@
+struct Inner {
+    x : i32;
+};
+
+[[block]]
+struct S {
+    a : vec3<i32>;
+    b : i32;
+    c : vec3<u32>;
+    d : u32;
+    e : vec3<f32>;
+    f : f32;
+    g : vec2<i32>;
+    h : vec2<i32>;
+    i : mat2x3<f32>;
+    j : mat3x2<f32>;
+    [[align(16)]] k : Inner;
+    [[align(16)]] l : [[stride(16)]] array<Inner, 4>;
+};
+
+[[binding(0), group(0)]] var<uniform> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    let a = s.a;
+    let b = s.b;
+    let c = s.c;
+    let d = s.d;
+    let e = s.e;
+    let f = s.f;
+    let g = s.g;
+    let h = s.h;
+    let i = s.i;
+    let j = s.j;
+    let k = s.k;
+    let l = s.l;
+}
diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.hlsl b/test/buffer/uniform/static_index/read.wgsl.expected.hlsl
new file mode 100644
index 0000000..7acef2b
--- /dev/null
+++ b/test/buffer/uniform/static_index/read.wgsl.expected.hlsl
@@ -0,0 +1,55 @@
+struct Inner {
+  int x;
+};
+struct tint_padded_array_element {
+  Inner el;
+};
+
+float2x3 tint_symbol_7(uint4 buffer[13], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const uint scalar_offset_1 = ((offset + 16u)) / 4;
+  return float2x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz));
+}
+
+float3x2 tint_symbol_9(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;
+  uint4 ubo_load_1 = buffer[scalar_offset_3 / 4];
+  const uint scalar_offset_4 = ((offset + 16u)) / 4;
+  uint4 ubo_load_2 = buffer[scalar_offset_4 / 4];
+  return float3x2(asfloat(((scalar_offset_2 & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_3 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_4 & 2) ? ubo_load_2.zw : ubo_load_2.xy)));
+}
+
+Inner tint_symbol_10(uint4 buffer[13], uint offset) {
+  const uint scalar_offset_5 = ((offset + 0u)) / 4;
+  const Inner tint_symbol_12 = {asint(buffer[scalar_offset_5 / 4][scalar_offset_5 % 4])};
+  return tint_symbol_12;
+}
+
+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;
+}
+
+cbuffer cbuffer_s : register(b0, space0) {
+  uint4 s[13];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+  const int3 a = asint(s[0].xyz);
+  const int b = asint(s[0].w);
+  const uint3 c = s[1].xyz;
+  const uint d = s[1].w;
+  const float3 e = asfloat(s[2].xyz);
+  const float f = asfloat(s[2].w);
+  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 Inner k = tint_symbol_10(s, 128u);
+  const tint_padded_array_element l[4] = tint_symbol_11(s, 144u);
+  return;
+}
diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.msl b/test/buffer/uniform/static_index/read.wgsl.expected.msl
new file mode 100644
index 0000000..bc3ce89
--- /dev/null
+++ b/test/buffer/uniform/static_index/read.wgsl.expected.msl
@@ -0,0 +1,46 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Inner {
+  /* 0x0000 */ int x;
+};
+struct tint_padded_array_element {
+  /* 0x0000 */ Inner el;
+  /* 0x0004 */ int8_t tint_pad[12];
+};
+struct tint_array_wrapper {
+  /* 0x0000 */ tint_padded_array_element arr[4];
+};
+struct S {
+  /* 0x0000 */ packed_int3 a;
+  /* 0x000c */ int b;
+  /* 0x0010 */ packed_uint3 c;
+  /* 0x001c */ uint d;
+  /* 0x0020 */ packed_float3 e;
+  /* 0x002c */ float f;
+  /* 0x0030 */ packed_int2 g;
+  /* 0x0038 */ packed_int2 h;
+  /* 0x0040 */ float2x3 i;
+  /* 0x0060 */ float3x2 j;
+  /* 0x0078 */ int8_t tint_pad_1[8];
+  /* 0x0080 */ Inner k;
+  /* 0x0084 */ int8_t tint_pad_2[12];
+  /* 0x0090 */ tint_array_wrapper l;
+};
+
+kernel void tint_symbol(constant S& s [[buffer(0)]]) {
+  int3 const a = s.a;
+  int const b = s.b;
+  uint3 const c = s.c;
+  uint const d = s.d;
+  float3 const e = s.e;
+  float const f = s.f;
+  int2 const g = s.g;
+  int2 const h = s.h;
+  float2x3 const i = s.i;
+  float3x2 const j = s.j;
+  Inner const k = s.k;
+  tint_array_wrapper const l = s.l;
+  return;
+}
+
diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.spvasm b/test/buffer/uniform/static_index/read.wgsl.expected.spvasm
new file mode 100644
index 0000000..25854c6
--- /dev/null
+++ b/test/buffer/uniform/static_index/read.wgsl.expected.spvasm
@@ -0,0 +1,116 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 67
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "a"
+               OpMemberName %S 1 "b"
+               OpMemberName %S 2 "c"
+               OpMemberName %S 3 "d"
+               OpMemberName %S 4 "e"
+               OpMemberName %S 5 "f"
+               OpMemberName %S 6 "g"
+               OpMemberName %S 7 "h"
+               OpMemberName %S 8 "i"
+               OpMemberName %S 9 "j"
+               OpMemberName %S 10 "k"
+               OpName %Inner "Inner"
+               OpMemberName %Inner 0 "x"
+               OpMemberName %S 11 "l"
+               OpName %s "s"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 12
+               OpMemberDecorate %S 2 Offset 16
+               OpMemberDecorate %S 3 Offset 28
+               OpMemberDecorate %S 4 Offset 32
+               OpMemberDecorate %S 5 Offset 44
+               OpMemberDecorate %S 6 Offset 48
+               OpMemberDecorate %S 7 Offset 56
+               OpMemberDecorate %S 8 Offset 64
+               OpMemberDecorate %S 8 ColMajor
+               OpMemberDecorate %S 8 MatrixStride 16
+               OpMemberDecorate %S 9 Offset 96
+               OpMemberDecorate %S 9 ColMajor
+               OpMemberDecorate %S 9 MatrixStride 8
+               OpMemberDecorate %S 10 Offset 128
+               OpMemberDecorate %Inner 0 Offset 0
+               OpMemberDecorate %S 11 Offset 144
+               OpDecorate %_arr_Inner_uint_4 ArrayStride 16
+               OpDecorate %s NonWritable
+               OpDecorate %s Binding 0
+               OpDecorate %s DescriptorSet 0
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+      %v2int = OpTypeVector %int 2
+%mat2v3float = OpTypeMatrix %v3float 2
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+      %Inner = OpTypeStruct %int
+     %uint_4 = OpConstant %uint 4
+%_arr_Inner_uint_4 = OpTypeArray %Inner %uint_4
+          %S = OpTypeStruct %v3int %int %v3uint %uint %v3float %float %v2int %v2int %mat2v3float %mat3v2float %Inner %_arr_Inner_uint_4
+%_ptr_Uniform_S = OpTypePointer Uniform %S
+          %s = OpVariable %_ptr_Uniform_S Uniform
+       %void = OpTypeVoid
+         %17 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_v3int = OpTypePointer Uniform %v3int
+     %uint_1 = OpConstant %uint 1
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+     %uint_2 = OpConstant %uint 2
+%_ptr_Uniform_v3uint = OpTypePointer Uniform %v3uint
+     %uint_3 = OpConstant %uint 3
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+%_ptr_Uniform_v3float = OpTypePointer Uniform %v3float
+     %uint_5 = OpConstant %uint 5
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+     %uint_6 = OpConstant %uint 6
+%_ptr_Uniform_v2int = OpTypePointer Uniform %v2int
+     %uint_7 = OpConstant %uint 7
+     %uint_8 = OpConstant %uint 8
+%_ptr_Uniform_mat2v3float = OpTypePointer Uniform %mat2v3float
+     %uint_9 = OpConstant %uint 9
+%_ptr_Uniform_mat3v2float = OpTypePointer Uniform %mat3v2float
+    %uint_10 = OpConstant %uint 10
+%_ptr_Uniform_Inner = OpTypePointer Uniform %Inner
+    %uint_11 = OpConstant %uint 11
+%_ptr_Uniform__arr_Inner_uint_4 = OpTypePointer Uniform %_arr_Inner_uint_4
+       %main = OpFunction %void None %17
+         %20 = OpLabel
+         %23 = OpAccessChain %_ptr_Uniform_v3int %s %uint_0
+         %24 = OpLoad %v3int %23
+         %27 = OpAccessChain %_ptr_Uniform_int %s %uint_1
+         %28 = OpLoad %int %27
+         %31 = OpAccessChain %_ptr_Uniform_v3uint %s %uint_2
+         %32 = OpLoad %v3uint %31
+         %35 = OpAccessChain %_ptr_Uniform_uint %s %uint_3
+         %36 = OpLoad %uint %35
+         %38 = OpAccessChain %_ptr_Uniform_v3float %s %uint_4
+         %39 = OpLoad %v3float %38
+         %42 = OpAccessChain %_ptr_Uniform_float %s %uint_5
+         %43 = OpLoad %float %42
+         %46 = OpAccessChain %_ptr_Uniform_v2int %s %uint_6
+         %47 = OpLoad %v2int %46
+         %49 = OpAccessChain %_ptr_Uniform_v2int %s %uint_7
+         %50 = OpLoad %v2int %49
+         %53 = OpAccessChain %_ptr_Uniform_mat2v3float %s %uint_8
+         %54 = OpLoad %mat2v3float %53
+         %57 = OpAccessChain %_ptr_Uniform_mat3v2float %s %uint_9
+         %58 = OpLoad %mat3v2float %57
+         %61 = OpAccessChain %_ptr_Uniform_Inner %s %uint_10
+         %62 = OpLoad %Inner %61
+         %65 = OpAccessChain %_ptr_Uniform__arr_Inner_uint_4 %s %uint_11
+         %66 = OpLoad %_arr_Inner_uint_4 %65
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.wgsl b/test/buffer/uniform/static_index/read.wgsl.expected.wgsl
new file mode 100644
index 0000000..d856b8c
--- /dev/null
+++ b/test/buffer/uniform/static_index/read.wgsl.expected.wgsl
@@ -0,0 +1,39 @@
+struct Inner {
+  x : i32;
+};
+
+[[block]]
+struct S {
+  a : vec3<i32>;
+  b : i32;
+  c : vec3<u32>;
+  d : u32;
+  e : vec3<f32>;
+  f : f32;
+  g : vec2<i32>;
+  h : vec2<i32>;
+  i : mat2x3<f32>;
+  j : mat3x2<f32>;
+  [[align(16)]]
+  k : Inner;
+  [[align(16)]]
+  l : [[stride(16)]] array<Inner, 4>;
+};
+
+[[binding(0), group(0)]] var<uniform> s : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let a = s.a;
+  let b = s.b;
+  let c = s.c;
+  let d = s.d;
+  let e = s.e;
+  let f = s.f;
+  let g = s.g;
+  let h = s.h;
+  let i = s.i;
+  let j = s.j;
+  let k = s.k;
+  let l = s.l;
+}
diff --git a/test/bug/dawn/947.wgsl.expected.hlsl b/test/bug/dawn/947.wgsl.expected.hlsl
index a65f3b1..56a4250 100644
--- a/test/bug/dawn/947.wgsl.expected.hlsl
+++ b/test/bug/dawn/947.wgsl.expected.hlsl
@@ -19,20 +19,11 @@
   float2 texcoord[3] = {float2(-0.5f, 0.0f), float2(1.5f, 0.0f), float2(0.5f, 2.0f)};
   VertexOutputs output = (VertexOutputs)0;
   output.position = float4(((texcoord[VertexIndex] * 2.0f) - float2(1.0f, 1.0f)), 0.0f, 1.0f);
-  const uint scalar_offset = (4u) / 4;
-  bool flipY = (asfloat(uniforms[scalar_offset / 4][scalar_offset % 4]) < 0.0f);
+  bool flipY = (asfloat(uniforms[0].y) < 0.0f);
   if (flipY) {
-    const uint scalar_offset_1 = (0u) / 4;
-    uint4 ubo_load = uniforms[scalar_offset_1 / 4];
-    const uint scalar_offset_2 = (8u) / 4;
-    uint4 ubo_load_1 = uniforms[scalar_offset_2 / 4];
-    output.texcoords = ((((texcoord[VertexIndex] * asfloat(((scalar_offset_1 & 2) ? ubo_load.zw : ubo_load.xy))) + asfloat(((scalar_offset_2 & 2) ? ubo_load_1.zw : ubo_load_1.xy))) * float2(1.0f, -1.0f)) + float2(0.0f, 1.0f));
+    output.texcoords = ((((texcoord[VertexIndex] * asfloat(uniforms[0].xy)) + asfloat(uniforms[0].zw)) * float2(1.0f, -1.0f)) + float2(0.0f, 1.0f));
   } else {
-    const uint scalar_offset_3 = (0u) / 4;
-    uint4 ubo_load_2 = uniforms[scalar_offset_3 / 4];
-    const uint scalar_offset_4 = (8u) / 4;
-    uint4 ubo_load_3 = uniforms[scalar_offset_4 / 4];
-    output.texcoords = ((((texcoord[VertexIndex] * float2(1.0f, -1.0f)) + float2(0.0f, 1.0f)) * asfloat(((scalar_offset_3 & 2) ? ubo_load_2.zw : ubo_load_2.xy))) + asfloat(((scalar_offset_4 & 2) ? ubo_load_3.zw : ubo_load_3.xy)));
+    output.texcoords = ((((texcoord[VertexIndex] * float2(1.0f, -1.0f)) + float2(0.0f, 1.0f)) * asfloat(uniforms[0].xy)) + asfloat(uniforms[0].zw));
   }
   const tint_symbol_2 tint_symbol_8 = {output.texcoords, output.position};
   return tint_symbol_8;
diff --git a/test/bug/tint/534.wgsl.expected.hlsl b/test/bug/tint/534.wgsl.expected.hlsl
index e7a735c..da61ba4 100644
--- a/test/bug/tint/534.wgsl.expected.hlsl
+++ b/test/bug/tint/534.wgsl.expected.hlsl
@@ -30,8 +30,7 @@
   int2 size = tint_tmp;
   int2 dstTexCoord = int2(GlobalInvocationID.xy);
   int2 srcTexCoord = dstTexCoord;
-  const uint scalar_offset = (0u) / 4;
-  if ((uniforms[scalar_offset / 4][scalar_offset % 4] == 1u)) {
+  if ((uniforms[0].x == 1u)) {
     srcTexCoord.y = ((size.y - dstTexCoord.y) - 1);
   }
   float4 srcColor = src.Load(int3(srcTexCoord, 0));
@@ -40,17 +39,13 @@
   uint4 srcColorBits = uint4(0u, 0u, 0u, 0u);
   uint4 dstColorBits = uint4(dstColor);
   {
-    uint i = 0u;
-    while (true) {
-      const uint scalar_offset_1 = (12u) / 4;
-      if (!((i < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]))) { break; }
+    for(uint i = 0u; (i < uniforms[0].w); i = (i + 1u)) {
       Set_uint4(srcColorBits, i, ConvertToFp16FloatValue(srcColor[i]));
       bool tint_tmp_1 = success;
       if (tint_tmp_1) {
         tint_tmp_1 = (srcColorBits[i] == dstColorBits[i]);
       }
       success = (tint_tmp_1);
-      i = (i + 1u);
     }
   }
   uint outputIndex = ((GlobalInvocationID.y * uint(size.x)) + GlobalInvocationID.x);
diff --git a/test/bug/tint/744.wgsl.expected.hlsl b/test/bug/tint/744.wgsl.expected.hlsl
index 47f47a7..c7bab90 100644
--- a/test/bug/tint/744.wgsl.expected.hlsl
+++ b/test/bug/tint/744.wgsl.expected.hlsl
@@ -13,10 +13,8 @@
 void main(tint_symbol_1 tint_symbol) {
   const uint3 global_id = tint_symbol.global_id;
   const uint2 resultCell = uint2(global_id.y, global_id.x);
-  const uint scalar_offset = (4u) / 4;
-  const uint dimInner = uniforms[scalar_offset / 4][scalar_offset % 4];
-  const uint scalar_offset_1 = (20u) / 4;
-  const uint dimOutter = uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4];
+  const uint dimInner = uniforms[0].y;
+  const uint dimOutter = uniforms[1].y;
   uint result = 0u;
   {
     for(uint i = 0u; (i < dimInner); i = (i + 1u)) {
diff --git a/test/bug/tint/749.spvasm.expected.hlsl b/test/bug/tint/749.spvasm.expected.hlsl
index 5d8b15b..c0db54c 100644
--- a/test/bug/tint/749.spvasm.expected.hlsl
+++ b/test/bug/tint/749.spvasm.expected.hlsl
@@ -861,9 +861,7 @@
   const float2 x_762 = uv;
   uv = float2(0.0f, 0.0f);
   uv = x_762;
-  const uint scalar_offset = (0u) / 4;
-  uint4 ubo_load = x_188[scalar_offset / 4];
-  const float2 x_191 = asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy));
+  const float2 x_191 = asfloat(x_188[0].xy);
   const QuicksortObject x_763 = obj;
   const int tint_symbol_51[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0};
   const QuicksortObject tint_symbol_52 = {tint_symbol_51};
diff --git a/test/bug/tint/913.wgsl.expected.hlsl b/test/bug/tint/913.wgsl.expected.hlsl
index 18dca3f..2e66d2a 100644
--- a/test/bug/tint/913.wgsl.expected.hlsl
+++ b/test/bug/tint/913.wgsl.expected.hlsl
@@ -25,23 +25,17 @@
   const uint2 dstTexCoord = uint2(GlobalInvocationID.xy);
   const float4 nonCoveredColor = float4(0.0f, 1.0f, 0.0f, 1.0f);
   bool success = true;
-  const uint scalar_offset = (16u) / 4;
-  bool tint_tmp_4 = (dstTexCoord.x < uniforms[scalar_offset / 4][scalar_offset % 4]);
+  bool tint_tmp_4 = (dstTexCoord.x < uniforms[1].x);
   if (!tint_tmp_4) {
-    const uint scalar_offset_1 = (20u) / 4;
-    tint_tmp_4 = (dstTexCoord.y < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]);
+    tint_tmp_4 = (dstTexCoord.y < uniforms[1].y);
   }
   bool tint_tmp_3 = (tint_tmp_4);
   if (!tint_tmp_3) {
-    const uint scalar_offset_2 = (16u) / 4;
-    const uint scalar_offset_3 = (24u) / 4;
-    tint_tmp_3 = (dstTexCoord.x >= (uniforms[scalar_offset_2 / 4][scalar_offset_2 % 4] + uniforms[scalar_offset_3 / 4][scalar_offset_3 % 4]));
+    tint_tmp_3 = (dstTexCoord.x >= (uniforms[1].x + uniforms[1].z));
   }
   bool tint_tmp_2 = (tint_tmp_3);
   if (!tint_tmp_2) {
-    const uint scalar_offset_4 = (20u) / 4;
-    const uint scalar_offset_5 = (28u) / 4;
-    tint_tmp_2 = (dstTexCoord.y >= (uniforms[scalar_offset_4 / 4][scalar_offset_4 % 4] + uniforms[scalar_offset_5 / 4][scalar_offset_5 % 4]));
+    tint_tmp_2 = (dstTexCoord.y >= (uniforms[1].y + uniforms[1].w));
   }
   if ((tint_tmp_2)) {
     bool tint_tmp_5 = success;
@@ -50,19 +44,13 @@
     }
     success = (tint_tmp_5);
   } else {
-    const uint scalar_offset_6 = (16u) / 4;
-    uint4 ubo_load = uniforms[scalar_offset_6 / 4];
-    const uint scalar_offset_7 = (8u) / 4;
-    uint4 ubo_load_1 = uniforms[scalar_offset_7 / 4];
-    uint2 srcTexCoord = ((dstTexCoord - ((scalar_offset_6 & 2) ? ubo_load.zw : ubo_load.xy)) + ((scalar_offset_7 & 2) ? ubo_load_1.zw : ubo_load_1.xy));
-    const uint scalar_offset_8 = (0u) / 4;
-    if ((uniforms[scalar_offset_8 / 4][scalar_offset_8 % 4] == 1u)) {
+    uint2 srcTexCoord = ((dstTexCoord - uniforms[1].xy) + uniforms[0].zw);
+    if ((uniforms[0].x == 1u)) {
       srcTexCoord.y = ((uint(srcSize.y) - srcTexCoord.y) - 1u);
     }
     const float4 srcColor = src.Load(int3(srcTexCoord, 0));
     const float4 dstColor = tint_symbol.Load(int3(dstTexCoord, 0));
-    const uint scalar_offset_9 = (4u) / 4;
-    if ((uniforms[scalar_offset_9 / 4][scalar_offset_9 % 4] == 2u)) {
+    if ((uniforms[0].y == 2u)) {
       bool tint_tmp_7 = success;
       if (tint_tmp_7) {
         tint_tmp_7 = aboutEqual(dstColor.r, srcColor.r);
diff --git a/test/bug/tint/914.wgsl.expected.hlsl b/test/bug/tint/914.wgsl.expected.hlsl
index 5e0ebb4..00be3a9 100644
--- a/test/bug/tint/914.wgsl.expected.hlsl
+++ b/test/bug/tint/914.wgsl.expected.hlsl
@@ -6,45 +6,36 @@
 };
 
 float mm_readA(uint row, uint col) {
-  const uint scalar_offset = (0u) / 4;
-  bool tint_tmp = (row < uniforms[scalar_offset / 4][scalar_offset % 4]);
+  bool tint_tmp = (row < uniforms[0].x);
   if (tint_tmp) {
-    const uint scalar_offset_1 = (4u) / 4;
-    tint_tmp = (col < uniforms[scalar_offset_1 / 4][scalar_offset_1 % 4]);
+    tint_tmp = (col < uniforms[0].y);
   }
   if ((tint_tmp)) {
-    const uint scalar_offset_2 = (4u) / 4;
-    const float result = asfloat(firstMatrix.Load((4u * ((row * uniforms[scalar_offset_2 / 4][scalar_offset_2 % 4]) + col))));
+    const float result = asfloat(firstMatrix.Load((4u * ((row * uniforms[0].y) + col))));
     return result;
   }
   return 0.0f;
 }
 
 float mm_readB(uint row, uint col) {
-  const uint scalar_offset_3 = (4u) / 4;
-  bool tint_tmp_1 = (row < uniforms[scalar_offset_3 / 4][scalar_offset_3 % 4]);
+  bool tint_tmp_1 = (row < uniforms[0].y);
   if (tint_tmp_1) {
-    const uint scalar_offset_4 = (8u) / 4;
-    tint_tmp_1 = (col < uniforms[scalar_offset_4 / 4][scalar_offset_4 % 4]);
+    tint_tmp_1 = (col < uniforms[0].z);
   }
   if ((tint_tmp_1)) {
-    const uint scalar_offset_5 = (8u) / 4;
-    const float result = asfloat(secondMatrix.Load((4u * ((row * uniforms[scalar_offset_5 / 4][scalar_offset_5 % 4]) + col))));
+    const float result = asfloat(secondMatrix.Load((4u * ((row * uniforms[0].z) + col))));
     return result;
   }
   return 0.0f;
 }
 
 void mm_write(uint row, uint col, float value) {
-  const uint scalar_offset_6 = (0u) / 4;
-  bool tint_tmp_2 = (row < uniforms[scalar_offset_6 / 4][scalar_offset_6 % 4]);
+  bool tint_tmp_2 = (row < uniforms[0].x);
   if (tint_tmp_2) {
-    const uint scalar_offset_7 = (8u) / 4;
-    tint_tmp_2 = (col < uniforms[scalar_offset_7 / 4][scalar_offset_7 % 4]);
+    tint_tmp_2 = (col < uniforms[0].z);
   }
   if ((tint_tmp_2)) {
-    const uint scalar_offset_8 = (8u) / 4;
-    const uint index = (col + (row * uniforms[scalar_offset_8 / 4][scalar_offset_8 % 4]));
+    const uint index = (col + (row * uniforms[0].z));
     resultMatrix.Store((4u * index), asuint(value));
   }
 }
@@ -93,8 +84,7 @@
   const uint tileCol = (local_id.x * ColPerThread);
   const uint globalRow = (global_id.y * RowPerThread);
   const uint globalCol = (global_id.x * ColPerThread);
-  const uint scalar_offset_9 = (4u) / 4;
-  const uint numTiles = (((uniforms[scalar_offset_9 / 4][scalar_offset_9 % 4] - 1u) / TileInner) + 1u);
+  const uint numTiles = (((uniforms[0].y - 1u) / TileInner) + 1u);
   float acc[16] = (float[16])0;
   float ACached = 0.0f;
   float BCached[4] = (float[4])0;
diff --git a/test/bug/tint/922.wgsl.expected.hlsl b/test/bug/tint/922.wgsl.expected.hlsl
index 964eb97..60eb716 100644
--- a/test/bug/tint/922.wgsl.expected.hlsl
+++ b/test/bug/tint/922.wgsl.expected.hlsl
@@ -196,8 +196,7 @@
   const float4 _e49 = Mul(_e35, _e48);
   gl_Position = _e49;
   v_Color = a_Color1;
-  const uint scalar_offset_9 = (32u) / 4;
-  const float4 _e52 = asfloat(global1[scalar_offset_9 / 4]);
+  const float4 _e52 = asfloat(global1[2]);
   if ((_e52.x == 2.0f)) {
     {
       const float3 _e59 = a_Normal1;
diff --git a/test/bug/tint/942.wgsl.expected.hlsl b/test/bug/tint/942.wgsl.expected.hlsl
index 695e8ae..ba42782 100644
--- a/test/bug/tint/942.wgsl.expected.hlsl
+++ b/test/bug/tint/942.wgsl.expected.hlsl
@@ -33,20 +33,17 @@
     }
   }
   GroupMemoryBarrierWithGroupSync();
-  const uint scalar_offset = (0u) / 4;
-  const uint filterOffset = ((params[scalar_offset / 4][scalar_offset % 4] - 1u) / 2u);
+  const uint filterOffset = ((params[0].x - 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));
+  const int2 baseIndex = (int2(((WorkGroupID.xy * uint2(params[0].y, 4u)) + (LocalInvocationID.xy * uint2(4u, 1u)))) - int2(int(filterOffset), 0));
   {
     for(uint r = 0u; (r < 4u); r = (r + 1u)) {
       {
         for(uint c = 0u; (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)) {
+          if ((flip[0].x != 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;
@@ -60,8 +57,7 @@
       {
         for(uint c = 0u; (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)) {
+          if ((flip[0].x != 0u)) {
             writeIndex = writeIndex.yx;
           }
           const uint center = ((4u * LocalInvocationID.x) + c);
@@ -76,14 +72,9 @@
           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; }
+              for(uint f = 0u; (f < params[0].x); f = (f + 1u)) {
                 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);
+                acc = (acc + ((1.0f / float(params[0].x)) * tile[r][i]));
               }
             }
             outputTex[writeIndex] = float4(acc, 1.0f);
diff --git a/test/bug/tint/943.spvasm.expected.hlsl b/test/bug/tint/943.spvasm.expected.hlsl
index a317e46..0087ca2 100644
--- a/test/bug/tint/943.spvasm.expected.hlsl
+++ b/test/bug/tint/943.spvasm.expected.hlsl
@@ -33,10 +33,8 @@
   int2 param_10 = int2(0, 0);
   int2 param_11 = int2(0, 0);
   float x_430 = 0.0f;
-  const uint scalar_offset = (20u) / 4;
-  const int x_417 = asint(x_48[scalar_offset / 4][scalar_offset % 4]);
-  const uint scalar_offset_1 = (24u) / 4;
-  const int x_419 = asint(x_48[scalar_offset_1 / 4][scalar_offset_1 % 4]);
+  const int x_417 = asint(x_48[1].y);
+  const int x_419 = asint(x_48[1].z);
   batchASize = (x_417 * x_419);
   const int x_421 = row;
   const int x_422 = col;
@@ -64,10 +62,8 @@
   int2 param_12 = int2(0, 0);
   int2 param_13 = int2(0, 0);
   float x_468 = 0.0f;
-  const uint scalar_offset_2 = (36u) / 4;
-  const int x_455 = asint(x_48[scalar_offset_2 / 4][scalar_offset_2 % 4]);
-  const uint scalar_offset_3 = (40u) / 4;
-  const int x_457 = asint(x_48[scalar_offset_3 / 4][scalar_offset_3 % 4]);
+  const int x_455 = asint(x_48[2].y);
+  const int x_457 = asint(x_48[2].z);
   batchBSize = (x_455 * x_457);
   const int x_459 = row_1;
   const int x_460 = col_1;
@@ -92,10 +88,8 @@
 
 int getOutputFlatIndex_vi3_(inout int3 coords) {
   const int3 x_99 = coords;
-  const uint scalar_offset_4 = (64u) / 4;
-  const int x_105 = asint(x_48[scalar_offset_4 / 4][scalar_offset_4 % 4]);
-  const uint scalar_offset_5 = (68u) / 4;
-  const int x_107 = asint(x_48[scalar_offset_5 / 4][scalar_offset_5 % 4]);
+  const int x_105 = asint(x_48[4].x);
+  const int x_107 = asint(x_48[4].y);
   return int(dot(float3(x_99), float3(int3(x_105, x_107, 1))));
 }
 
@@ -330,14 +324,11 @@
   int param_18 = 0;
   int param_19 = 0;
   int param_20 = 0;
-  const uint scalar_offset_6 = (20u) / 4;
-  const int x_67 = asint(x_48[scalar_offset_6 / 4][scalar_offset_6 % 4]);
+  const int x_67 = asint(x_48[1].y);
   dimAOuter_1 = x_67;
-  const uint scalar_offset_7 = (24u) / 4;
-  const int x_71 = asint(x_48[scalar_offset_7 / 4][scalar_offset_7 % 4]);
+  const int x_71 = asint(x_48[1].z);
   dimInner_1 = x_71;
-  const uint scalar_offset_8 = (40u) / 4;
-  const int x_75 = asint(x_48[scalar_offset_8 / 4][scalar_offset_8 % 4]);
+  const int x_75 = asint(x_48[2].z);
   dimBOuter_1 = x_75;
   const uint x_505 = gl_GlobalInvocationID.z;
   batch = asint(x_505);
diff --git a/test/bug/tint/948.wgsl.expected.hlsl b/test/bug/tint/948.wgsl.expected.hlsl
index d6f3a02..694cd53 100644
--- a/test/bug/tint/948.wgsl.expected.hlsl
+++ b/test/bug/tint/948.wgsl.expected.hlsl
@@ -22,8 +22,7 @@
 float4x4 getFrameData_f1_(inout float frameID) {
   float fX = 0.0f;
   const float x_15 = frameID;
-  const uint scalar_offset = (108u) / 4;
-  const float x_25 = asfloat(x_20[scalar_offset / 4][scalar_offset % 4]);
+  const float x_25 = asfloat(x_20[6].w);
   fX = (x_15 / x_25);
   const float4 x_40 = frameMapTexture.SampleBias(frameMapSampler, float2(fX, 0.0f), 0.0f);
   const float4 x_47 = frameMapTexture.SampleBias(frameMapSampler, float2(fX, 0.25f), 0.0f);
@@ -55,16 +54,11 @@
   const float x_91 = tileUV.y;
   tileUV.y = (1.0f - x_91);
   tileID = floor(tUV);
-  const uint scalar_offset_1 = (96u) / 4;
-  uint4 ubo_load = x_20[scalar_offset_1 / 4];
-  const float2 x_101 = asfloat(((scalar_offset_1 & 2) ? ubo_load.zw : ubo_load.xy));
+  const float2 x_101 = asfloat(x_20[6].xy);
   sheetUnits = (float2(1.0f, 1.0f) / x_101);
-  const uint scalar_offset_2 = (108u) / 4;
-  const float x_106 = asfloat(x_20[scalar_offset_2 / 4][scalar_offset_2 % 4]);
+  const float x_106 = asfloat(x_20[6].w);
   spriteUnits = (1.0f / x_106);
-  const uint scalar_offset_3 = (88u) / 4;
-  uint4 ubo_load_1 = x_20[scalar_offset_3 / 4];
-  const float2 x_111 = asfloat(((scalar_offset_3 & 2) ? ubo_load_1.zw : ubo_load_1.xy));
+  const float2 x_111 = asfloat(x_20[5].zw);
   stageUnits = (float2(1.0f, 1.0f) / x_111);
   i = 0;
   {
@@ -72,18 +66,14 @@
       switch(i) {
         case 1: {
           const float2 x_150 = tileID;
-          const uint scalar_offset_4 = (88u) / 4;
-          uint4 ubo_load_2 = x_20[scalar_offset_4 / 4];
-          const float2 x_154 = asfloat(((scalar_offset_4 & 2) ? ubo_load_2.zw : ubo_load_2.xy));
+          const float2 x_154 = asfloat(x_20[5].zw);
           const float4 x_156 = tileMapsTexture1.SampleBias(tileMapsSampler, ((x_150 + float2(0.5f, 0.5f)) / x_154), 0.0f);
           frameID_1 = x_156.x;
           break;
         }
         case 0: {
           const float2 x_136 = tileID;
-          const uint scalar_offset_5 = (88u) / 4;
-          uint4 ubo_load_3 = x_20[scalar_offset_5 / 4];
-          const float2 x_140 = asfloat(((scalar_offset_5 & 2) ? ubo_load_3.zw : ubo_load_3.xy));
+          const float2 x_140 = asfloat(x_20[5].zw);
           const float4 x_142 = tileMapsTexture0.SampleBias(tileMapsSampler, ((x_136 + float2(0.5f, 0.5f)) / x_140), 0.0f);
           frameID_1 = x_142.x;
           break;
@@ -93,14 +83,12 @@
         }
       }
       const float x_166 = frameID_1;
-      const uint scalar_offset_6 = (108u) / 4;
-      const float x_169 = asfloat(x_20[scalar_offset_6 / 4][scalar_offset_6 % 4]);
+      const float x_169 = asfloat(x_20[6].w);
       const float4 x_172 = animationMapTexture.SampleBias(animationMapSampler, float2(((x_166 + 0.5f) / x_169), 0.0f), 0.0f);
       animationData = x_172;
       const float x_174 = animationData.y;
       if ((x_174 > 0.0f)) {
-        const uint scalar_offset_7 = (0u) / 4;
-        const float x_181 = asfloat(x_20[scalar_offset_7 / 4][scalar_offset_7 % 4]);
+        const float x_181 = asfloat(x_20[0].x);
         const float x_184 = animationData.z;
         mt = ((x_181 * x_184) % 1.0f);
         f = 0.0f;
@@ -113,8 +101,7 @@
               break;
             }
             const float x_208 = frameID_1;
-            const uint scalar_offset_8 = (108u) / 4;
-            const float x_211 = asfloat(x_20[scalar_offset_8 / 4][scalar_offset_8 % 4]);
+            const float x_211 = asfloat(x_20[6].w);
             const float4 x_217 = animationMapTexture.SampleBias(animationMapSampler, float2(((x_208 + 0.5f) / x_211), (0.125f * f)), 0.0f);
             animationData = x_217;
           }
@@ -124,9 +111,7 @@
       const float4x4 x_225 = getFrameData_f1_(param);
       frameData = x_225;
       const float4 x_228 = frameData[0];
-      const uint scalar_offset_9 = (96u) / 4;
-      uint4 ubo_load_4 = x_20[scalar_offset_9 / 4];
-      const float2 x_231 = asfloat(((scalar_offset_9 & 2) ? ubo_load_4.zw : ubo_load_4.xy));
+      const float2 x_231 = asfloat(x_20[6].xy);
       frameSize = (float2(x_228.w, x_228.z) / x_231);
       const float4 x_235 = frameData[0];
       offset_1 = (float2(x_235.x, x_235.y) * sheetUnits);
@@ -156,8 +141,7 @@
       }
     }
   }
-  const uint scalar_offset_10 = (112u) / 4;
-  const float3 x_310 = asfloat(x_20[scalar_offset_10 / 4].xyz);
+  const float3 x_310 = asfloat(x_20[7].xyz);
   const float4 x_311 = color;
   const float3 x_313 = (float3(x_311.x, x_311.y, x_311.z) * x_310);
   color = float4(x_313.x, x_313.y, x_313.z, color.w);
diff --git a/test/bug/tint/949.wgsl.expected.hlsl b/test/bug/tint/949.wgsl.expected.hlsl
index 2bb720a..5aca5bf 100644
--- a/test/bug/tint/949.wgsl.expected.hlsl
+++ b/test/bug/tint/949.wgsl.expected.hlsl
@@ -200,17 +200,14 @@
   const float4 x_262 = TextureSamplerTexture.Sample(TextureSamplerSampler, vMainuv);
   tempTextureRead = x_262;
   const float4 x_264 = tempTextureRead;
-  const uint scalar_offset = (160u) / 4;
-  const float x_273 = asfloat(x_269[scalar_offset / 4][scalar_offset % 4]);
+  const float x_273 = asfloat(x_269[10].x);
   rgb = (float3(x_264.x, x_264.y, x_264.z) * x_273);
-  const uint scalar_offset_1 = (144u) / 4;
-  const float3 x_279 = asfloat(x_269[scalar_offset_1 / 4].xyz);
+  const float3 x_279 = asfloat(x_269[9].xyz);
   const float4 x_282 = v_output1;
   output5 = normalize((x_279 - float3(x_282.x, x_282.y, x_282.z)));
   output4 = float4(0.0f, 0.0f, 0.0f, 0.0f);
   uvOffset = float2(0.0f, 0.0f);
-  const uint scalar_offset_2 = (128u) / 4;
-  const float x_292 = asfloat(x_269[scalar_offset_2 / 4][scalar_offset_2 % 4]);
+  const float x_292 = asfloat(x_269[8].x);
   normalScale = (1.0f / x_292);
   if (gl_FrontFacing) {
     x_299 = v_uv;
@@ -223,9 +220,7 @@
   const float4 x_317 = v_output1;
   param_4 = float3(x_317.x, x_317.y, x_317.z);
   param_5 = TBNUV;
-  const uint scalar_offset_3 = (168u) / 4;
-  uint4 ubo_load = x_269[scalar_offset_3 / 4];
-  const float2 x_324 = asfloat(((scalar_offset_3 & 2) ? ubo_load.zw : ubo_load.xy));
+  const float2 x_324 = asfloat(x_269[10].zw);
   param_6 = x_324;
   const float3x3 x_325 = cotangent_frame_vf3_vf3_vf2_vf2_(param_3, param_4, param_5, param_6);
   TBN = x_325;
@@ -234,8 +229,7 @@
   invTBN = x_329;
   const float3 x_334 = mul(-(output5), invTBN);
   parallaxLimit = (length(float2(x_334.x, x_334.y)) / mul(-(output5), invTBN).z);
-  const uint scalar_offset_4 = (156u) / 4;
-  const float x_345 = asfloat(x_269[scalar_offset_4 / 4][scalar_offset_4 % 4]);
+  const float x_345 = asfloat(x_269[9].w);
   parallaxLimit = (parallaxLimit * x_345);
   const float3 x_352 = mul(-(output5), invTBN);
   vOffsetDir = normalize(float2(x_352.x, x_352.y));
@@ -270,8 +264,7 @@
   parallaxOcclusion_0 = vCurrOffset;
   uvOffset = parallaxOcclusion_0;
   const float4 x_452 = TextureSamplerTexture.Sample(TextureSamplerSampler, (v_uv + uvOffset));
-  const uint scalar_offset_5 = (128u) / 4;
-  const float x_454 = asfloat(x_269[scalar_offset_5 / 4][scalar_offset_5 % 4]);
+  const float x_454 = asfloat(x_269[8].x);
   param_8 = TBN;
   param_9 = float3(x_452.x, x_452.y, x_452.z);
   param_10 = (1.0f / x_454);
@@ -282,8 +275,7 @@
   tempTextureRead1 = x_475;
   const float4 x_477 = tempTextureRead1;
   rgb1 = float3(x_477.x, x_477.y, x_477.z);
-  const uint scalar_offset_6 = (144u) / 4;
-  const float3 x_481 = asfloat(x_269[scalar_offset_6 / 4].xyz);
+  const float3 x_481 = asfloat(x_269[9].xyz);
   const float4 x_482 = v_output1;
   viewDirectionW_1 = normalize((x_481 - float3(x_482.x, x_482.y, x_482.z)));
   shadow = 1.0f;
@@ -294,17 +286,13 @@
   normalW = float3(x_494.x, x_494.y, x_494.z);
   param_11 = viewDirectionW_1;
   param_12 = normalW;
-  const uint scalar_offset_7 = (0u) / 4;
-  const float4 x_507 = asfloat(light0[scalar_offset_7 / 4]);
+  const float4 x_507 = asfloat(light0[0]);
   param_13 = x_507;
-  const uint scalar_offset_8 = (16u) / 4;
-  const float4 x_510 = asfloat(light0[scalar_offset_8 / 4]);
+  const float4 x_510 = asfloat(light0[1]);
   param_14 = float3(x_510.x, x_510.y, x_510.z);
-  const uint scalar_offset_9 = (32u) / 4;
-  const float4 x_514 = asfloat(light0[scalar_offset_9 / 4]);
+  const float4 x_514 = asfloat(light0[2]);
   param_15 = float3(x_514.x, x_514.y, x_514.z);
-  const uint scalar_offset_10 = (48u) / 4;
-  const float3 x_518 = asfloat(light0[scalar_offset_10 / 4].xyz);
+  const float3 x_518 = asfloat(light0[3].xyz);
   param_16 = x_518;
   param_17 = glossiness_1;
   const lightingInfo x_521 = computeHemisphericLighting_vf3_vf3_vf4_vf3_vf3_vf3_f1_(param_11, param_12, param_13, param_14, param_15, param_16, param_17);
diff --git a/test/bug/tint/951.spvasm.expected.hlsl b/test/bug/tint/951.spvasm.expected.hlsl
index 24ddc28..86a56b5 100644
--- a/test/bug/tint/951.spvasm.expected.hlsl
+++ b/test/bug/tint/951.spvasm.expected.hlsl
@@ -36,8 +36,7 @@
   const uint x_61 = gl_GlobalInvocationID.x;
   index = asint(x_61);
   const int x_63 = index;
-  const uint scalar_offset = (16u) / 4;
-  const int x_70 = asint(x_24[scalar_offset / 4][scalar_offset % 4]);
+  const int x_70 = asint(x_24[1].x);
   if ((x_63 < x_70)) {
     const float x_75 = getAAtOutCoords_();
     a_1 = x_75;
diff --git a/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl
index 45ec648..634b49e 100644
--- a/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl
+++ b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl
@@ -4,7 +4,6 @@
 
 [numthreads(1, 1, 1)]
 void main() {
-  const uint scalar_offset = (0u) / 4;
-  const int use = (asint(v[scalar_offset / 4][scalar_offset % 4]) + 1);
+  const int use = (asint(v[0].x) + 1);
   return;
 }
diff --git a/test/samples/compute_boids.wgsl.expected.hlsl b/test/samples/compute_boids.wgsl.expected.hlsl
index 4643faa..3a88307 100644
--- a/test/samples/compute_boids.wgsl.expected.hlsl
+++ b/test/samples/compute_boids.wgsl.expected.hlsl
@@ -59,17 +59,14 @@
       }
       pos = asfloat(particlesA.Load2((16u * i))).xy;
       vel = asfloat(particlesA.Load2(((16u * i) + 8u))).xy;
-      const uint scalar_offset = (4u) / 4;
-      if ((distance(pos, vPos) < asfloat(params[scalar_offset / 4][scalar_offset % 4]))) {
+      if ((distance(pos, vPos) < asfloat(params[0].y))) {
         cMass = (cMass + pos);
         cMassCount = (cMassCount + 1);
       }
-      const uint scalar_offset_1 = (8u) / 4;
-      if ((distance(pos, vPos) < asfloat(params[scalar_offset_1 / 4][scalar_offset_1 % 4]))) {
+      if ((distance(pos, vPos) < asfloat(params[0].z))) {
         colVel = (colVel - (pos - vPos));
       }
-      const uint scalar_offset_2 = (12u) / 4;
-      if ((distance(pos, vPos) < asfloat(params[scalar_offset_2 / 4][scalar_offset_2 % 4]))) {
+      if ((distance(pos, vPos) < asfloat(params[0].w))) {
         cVel = (cVel + vel);
         cVelCount = (cVelCount + 1);
       }
@@ -81,13 +78,9 @@
   if ((cVelCount > 0)) {
     cVel = (cVel / float2(float(cVelCount), float(cVelCount)));
   }
-  const uint scalar_offset_3 = (16u) / 4;
-  const uint scalar_offset_4 = (20u) / 4;
-  const uint scalar_offset_5 = (24u) / 4;
-  vVel = (((vVel + (cMass * asfloat(params[scalar_offset_3 / 4][scalar_offset_3 % 4]))) + (colVel * asfloat(params[scalar_offset_4 / 4][scalar_offset_4 % 4]))) + (cVel * asfloat(params[scalar_offset_5 / 4][scalar_offset_5 % 4])));
+  vVel = (((vVel + (cMass * asfloat(params[1].x))) + (colVel * asfloat(params[1].y))) + (cVel * asfloat(params[1].z)));
   vVel = (normalize(vVel) * clamp(length(vVel), 0.0f, 0.100000001f));
-  const uint scalar_offset_6 = (0u) / 4;
-  vPos = (vPos + (vVel * asfloat(params[scalar_offset_6 / 4][scalar_offset_6 % 4])));
+  vPos = (vPos + (vVel * asfloat(params[0].x)));
   if ((vPos.x < -1.0f)) {
     vPos.x = 1.0f;
   }