diff --git a/docs/origin-trial-changes.md b/docs/origin-trial-changes.md
index bc03ac6..435934e 100644
--- a/docs/origin-trial-changes.md
+++ b/docs/origin-trial-changes.md
@@ -13,6 +13,7 @@
 * The `[[block]]` attribute has been deprecated. [tint:1324](https://crbug.com/tint/1324)
 * Attributes now use the `@decoration` syntax instead of the `[[decoration]]` syntax. [tint:1382](https://crbug.com/tint/1382)
 * `elseif` has been replaced with `else if`. [tint:1289](https://crbug.com/tint/1289)
+* The `[[stride]]` attribute has been deprecated. [tint:1381](https://crbug.com/tint/1381)
 
 ### New Features
 
diff --git a/src/reader/wgsl/parser_impl.cc b/src/reader/wgsl/parser_impl.cc
index 9f54309..6570e43 100644
--- a/src/reader/wgsl/parser_impl.cc
+++ b/src/reader/wgsl/parser_impl.cc
@@ -3076,7 +3076,9 @@
       auto val = expect_nonzero_positive_sint(use);
       if (val.errored)
         return Failure::kErrored;
-
+      deprecated(t.source(),
+                 "the @stride attribute is deprecated; use a larger type if "
+                 "necessary");
       return create<ast::StrideDecoration>(t.source(), val.value);
     });
   }
diff --git a/src/reader/wgsl/parser_impl_error_msg_test.cc b/src/reader/wgsl/parser_impl_error_msg_test.cc
index dfb5c57..3c7b91d 100644
--- a/src/reader/wgsl/parser_impl_error_msg_test.cc
+++ b/src/reader/wgsl/parser_impl_error_msg_test.cc
@@ -985,24 +985,24 @@
 }
 
 TEST_F(ParserImplErrorTest, GlobalDeclVarArrayDecoNotArray) {
-  EXPECT("var i : @stride(1) i32;",
+  EXPECT("var i : @location(1) i32;",
          R"(test.wgsl:1:10 error: unexpected decorations
-var i : @stride(1) i32;
-         ^^^^^^
+var i : @location(1) i32;
+         ^^^^^^^^
 )");
 }
 
 // TODO(crbug.com/tint/1382): Remove
 TEST_F(ParserImplErrorTest, DEPRECATED_GlobalDeclVarArrayDecoMissingEnd) {
   EXPECT(
-      "var i : [[stride(1) array<i32>;",
+      "var i : [[location(1) array<i32>;",
       R"(test.wgsl:1:9 warning: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style
-var i : [[stride(1) array<i32>;
+var i : [[location(1) array<i32>;
         ^^
 
-test.wgsl:1:21 error: expected ']]' for decoration list
-var i : [[stride(1) array<i32>;
-                    ^^^^^
+test.wgsl:1:23 error: expected ']]' for decoration list
+var i : [[location(1) array<i32>;
+                      ^^^^^
 )");
 }
 
@@ -1025,14 +1025,14 @@
 TEST_F(ParserImplErrorTest,
        DEPRECATED_GlobalDeclVarArrayDecoStrideMissingRParen) {
   EXPECT(
-      "var i : [[stride(1]] array<i32>;",
+      "var i : [[location(1]] array<i32>;",
       R"(test.wgsl:1:9 warning: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style
-var i : [[stride(1]] array<i32>;
+var i : [[location(1]] array<i32>;
         ^^
 
-test.wgsl:1:19 error: expected ')' for stride decoration
-var i : [[stride(1]] array<i32>;
-                  ^^
+test.wgsl:1:21 error: expected ')' for location decoration
+var i : [[location(1]] array<i32>;
+                    ^^
 )");
 }
 
diff --git a/src/reader/wgsl/parser_impl_type_decl_test.cc b/src/reader/wgsl/parser_impl_type_decl_test.cc
index 13e368a..476b19e 100644
--- a/src/reader/wgsl/parser_impl_type_decl_test.cc
+++ b/src/reader/wgsl/parser_impl_type_decl_test.cc
@@ -534,7 +534,10 @@
   EXPECT_FALSE(t.matched);
   ASSERT_EQ(t.value, nullptr);
   ASSERT_TRUE(p->has_error());
-  EXPECT_EQ(p->error(), "1:2: unexpected decorations");
+  EXPECT_EQ(
+      p->error(),
+      R"(1:2: use of deprecated language feature: the @stride attribute is deprecated; use a larger type if necessary
+1:2: unexpected decorations)");
 }
 
 TEST_F(ParserImplTest, TypeDecl_Array_Decoration_UnknownDecoration) {
@@ -546,6 +549,7 @@
   ASSERT_TRUE(p->has_error());
   EXPECT_EQ(p->error(), R"(1:2: expected decoration)");
 }
+
 TEST_F(ParserImplTest, TypeDecl_Array_Stride_MissingLeftParen) {
   auto p = parser("@stride 4) array<f32, 5>");
   auto t = p->type_decl();
@@ -563,7 +567,10 @@
   EXPECT_FALSE(t.matched);
   ASSERT_EQ(t.value, nullptr);
   ASSERT_TRUE(p->has_error());
-  EXPECT_EQ(p->error(), R"(1:11: expected ')' for stride decoration)");
+  EXPECT_EQ(
+      p->error(),
+      R"(1:2: use of deprecated language feature: the @stride attribute is deprecated; use a larger type if necessary
+1:11: expected ')' for stride decoration)");
 }
 
 TEST_F(ParserImplTest, TypeDecl_Array_Stride_MissingValue) {
@@ -610,6 +617,7 @@
   EXPECT_EQ(
       p->error(),
       R"(1:1: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style
+1:3: use of deprecated language feature: the @stride attribute is deprecated; use a larger type if necessary
 1:14: expected ']]' for decoration list)");
 }
 
@@ -638,6 +646,7 @@
   EXPECT_EQ(
       p->error(),
       R"(1:1: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style
+1:3: use of deprecated language feature: the @stride attribute is deprecated; use a larger type if necessary
 1:11: expected ')' for stride decoration)");
 }
 
diff --git a/src/reader/wgsl/parser_impl_variable_ident_decl_test.cc b/src/reader/wgsl/parser_impl_variable_ident_decl_test.cc
index 73efe8c..3c43ed8 100644
--- a/src/reader/wgsl/parser_impl_variable_ident_decl_test.cc
+++ b/src/reader/wgsl/parser_impl_variable_ident_decl_test.cc
@@ -77,7 +77,7 @@
 }
 
 TEST_F(ParserImplTest, VariableIdentDecl_NonAccessDecoFail) {
-  auto p = parser("my_var : @stride(1) S");
+  auto p = parser("my_var : @location(1) S");
 
   auto* mem = Member("a", ty.i32(), ast::DecorationList{});
   ast::StructMemberList members;
@@ -94,11 +94,11 @@
 }
 
 TEST_F(ParserImplTest, VariableIdentDecl_DecorationMissingRightParen) {
-  auto p = parser("my_var : @stride(4 S");
+  auto p = parser("my_var : @location(4 S");
   auto decl = p->expect_variable_ident_decl("test");
   ASSERT_TRUE(p->has_error());
   ASSERT_TRUE(decl.errored);
-  ASSERT_EQ(p->error(), "1:20: expected ')' for stride decoration");
+  ASSERT_EQ(p->error(), "1:22: expected ')' for location decoration");
 }
 
 TEST_F(ParserImplTest, VariableIdentDecl_DecorationMissingLeftParen) {
@@ -112,27 +112,27 @@
 // TODO(crbug.com/tint/1382): Remove
 TEST_F(ParserImplTest,
        DEPRECATED_VariableIdentDecl_DecorationMissingRightBlock) {
-  auto p = parser("my_var : [[stride(4) S");
+  auto p = parser("my_var : [[location(4) S");
   auto decl = p->expect_variable_ident_decl("test");
   ASSERT_TRUE(p->has_error());
   ASSERT_TRUE(decl.errored);
   ASSERT_EQ(
       p->error(),
       R"(1:10: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style
-1:22: expected ']]' for decoration list)");
+1:24: expected ']]' for decoration list)");
 }
 
 // TODO(crbug.com/tint/1382): Remove
 TEST_F(ParserImplTest,
        DEPRECATED_VariableIdentDecl_DecorationMissingRightParen) {
-  auto p = parser("my_var : [[stride(4]] S");
+  auto p = parser("my_var : [[location(4]] S");
   auto decl = p->expect_variable_ident_decl("test");
   ASSERT_TRUE(p->has_error());
   ASSERT_TRUE(decl.errored);
   ASSERT_EQ(
       p->error(),
       R"(1:10: use of deprecated language feature: [[decoration]] style decorations have been replaced with @decoration style
-1:20: expected ')' for stride decoration)");
+1:22: expected ')' for location decoration)");
 }
 
 // TODO(crbug.com/tint/1382): Remove
diff --git a/src/resolver/resolver_validation.cc b/src/resolver/resolver_validation.cc
index 5edc4ee..8864a26 100644
--- a/src/resolver/resolver_validation.cc
+++ b/src/resolver/resolver_validation.cc
@@ -321,15 +321,29 @@
         // bytes above, so we only need to validate that stride is a multiple of
         // 16 bytes.
         if (arr->Stride() % 16 != 0) {
+          // Since WGSL has no stride attribute, try to provide a useful hint
+          // for how the shader author can resolve the issue.
+          std::string hint;
+          if (arr->ElemType()->is_scalar()) {
+            hint =
+                "Consider using a vector or struct as the element type "
+                "instead.";
+          } else if (auto* vec = arr->ElemType()->As<sem::Vector>();
+                     vec && vec->type()->Size() == 4) {
+            hint = "Consider using a vec4 instead.";
+          } else if (arr->ElemType()->Is<sem::Struct>()) {
+            hint =
+                "Consider using the @size attribute on the last struct member.";
+          } else {
+            hint =
+                "Consider wrapping the element type in a struct and using the "
+                "@size attribute.";
+          }
           AddError(
               "uniform storage requires that array elements be aligned to 16 "
-              "bytes, but array stride of '" +
+              "bytes, but array element alignment of '" +
                   member_name_of(m) + "' is currently " +
-                  std::to_string(arr->Stride()) +
-                  ". Consider setting @stride(" +
-                  std::to_string(
-                      utils::RoundUp(required_align, arr->Stride())) +
-                  ") on the array type",
+                  std::to_string(arr->Stride()) + ". " + hint,
               m->Declaration()->type->source);
           AddNote("see layout of struct:\n" + str->Layout(builder_->Symbols()),
                   str->Declaration()->source);
diff --git a/src/resolver/storage_class_layout_validation_test.cc b/src/resolver/storage_class_layout_validation_test.cc
index f9fcc12..a8e7056 100644
--- a/src/resolver/storage_class_layout_validation_test.cc
+++ b/src/resolver/storage_class_layout_validation_test.cc
@@ -378,8 +378,8 @@
 
 // Detect array stride must be a multiple of 16 bytes for uniform buffers
 TEST_F(ResolverStorageClassLayoutValidationTest,
-       UniformBuffer_InvalidArrayStride) {
-  // type Inner = @stride(8) array<f32, 10>;
+       UniformBuffer_InvalidArrayStride_Scalar) {
+  // type Inner = array<f32, 10>;
   //
   // [[block]]
   // struct Outer {
@@ -390,7 +390,7 @@
   // @group(0) @binding(0)
   // var<uniform> a : Outer;
 
-  Alias("Inner", ty.array(ty.f32(), 10, 8));
+  Alias("Inner", ty.array(ty.f32(), 10));
 
   Structure(Source{{12, 34}}, "Outer",
             {
@@ -405,10 +405,93 @@
   ASSERT_FALSE(r()->Resolve());
   EXPECT_EQ(
       r()->error(),
-      R"(34:56 error: uniform storage requires that array elements be aligned to 16 bytes, but array stride of 'inner' is currently 8. Consider setting @stride(16) on the array type
+      R"(34:56 error: uniform storage requires that array elements be aligned to 16 bytes, but array element alignment of 'inner' is currently 4. Consider using a vector or struct as the element type instead.
+12:34 note: see layout of struct:
+/*            align(4) size(44) */ struct Outer {
+/* offset( 0) align(4) size(40) */   inner : array<f32, 10>;
+/* offset(40) align(4) size( 4) */   scalar : i32;
+/*                              */ };
+78:90 note: see declaration of variable)");
+}
+
+TEST_F(ResolverStorageClassLayoutValidationTest,
+       UniformBuffer_InvalidArrayStride_Vector) {
+  // type Inner = array<vec2<f32>, 10>;
+  //
+  // [[block]]
+  // struct Outer {
+  //   inner : Inner;
+  //   scalar : i32;
+  // };
+  //
+  // @group(0) @binding(0)
+  // var<uniform> a : Outer;
+
+  Alias("Inner", ty.array(ty.vec2<f32>(), 10));
+
+  Structure(Source{{12, 34}}, "Outer",
+            {
+                Member("inner", ty.type_name(Source{{34, 56}}, "Inner")),
+                Member("scalar", ty.i32()),
+            },
+            {StructBlock()});
+
+  Global(Source{{78, 90}}, "a", ty.type_name("Outer"),
+         ast::StorageClass::kUniform, GroupAndBinding(0, 0));
+
+  ASSERT_FALSE(r()->Resolve());
+  EXPECT_EQ(
+      r()->error(),
+      R"(34:56 error: uniform storage requires that array elements be aligned to 16 bytes, but array element alignment of 'inner' is currently 8. Consider using a vec4 instead.
+12:34 note: see layout of struct:
+/*            align(8) size(88) */ struct Outer {
+/* offset( 0) align(8) size(80) */   inner : array<vec2<f32>, 10>;
+/* offset(80) align(4) size( 4) */   scalar : i32;
+/* offset(84) align(1) size( 4) */   // -- implicit struct size padding --;
+/*                              */ };
+78:90 note: see declaration of variable)");
+}
+
+TEST_F(ResolverStorageClassLayoutValidationTest,
+       UniformBuffer_InvalidArrayStride_Struct) {
+  // struct ArrayElem {
+  //   a : f32;
+  //   b : i32;
+  // }
+  // type Inner = array<ArrayElem, 10>;
+  //
+  // [[block]]
+  // struct Outer {
+  //   inner : Inner;
+  //   scalar : i32;
+  // };
+  //
+  // @group(0) @binding(0)
+  // var<uniform> a : Outer;
+
+  auto* array_elem = Structure("ArrayElem", {
+                                                Member("a", ty.f32()),
+                                                Member("b", ty.i32()),
+                                            });
+  Alias("Inner", ty.array(ty.Of(array_elem), 10));
+
+  Structure(Source{{12, 34}}, "Outer",
+            {
+                Member("inner", ty.type_name(Source{{34, 56}}, "Inner")),
+                Member("scalar", ty.i32()),
+            },
+            {StructBlock()});
+
+  Global(Source{{78, 90}}, "a", ty.type_name("Outer"),
+         ast::StorageClass::kUniform, GroupAndBinding(0, 0));
+
+  ASSERT_FALSE(r()->Resolve());
+  EXPECT_EQ(
+      r()->error(),
+      R"(34:56 error: uniform storage requires that array elements be aligned to 16 bytes, but array element alignment of 'inner' is currently 8. Consider using the @size attribute on the last struct member.
 12:34 note: see layout of struct:
 /*            align(4) size(84) */ struct Outer {
-/* offset( 0) align(4) size(80) */   inner : @stride(8) array<f32, 10>;
+/* offset( 0) align(4) size(80) */   inner : array<ArrayElem, 10>;
 /* offset(80) align(4) size( 4) */   scalar : i32;
 /*                              */ };
 78:90 note: see declaration of variable)");
diff --git a/test/array/assign_to_function_var.wgsl b/test/array/assign_to_function_var.wgsl
index d5ccaa2..b2dc68d 100644
--- a/test/array/assign_to_function_var.wgsl
+++ b/test/array/assign_to_function_var.wgsl
@@ -1,4 +1,4 @@
-type ArrayType = @stride(16) array<i32, 4>;
+type ArrayType = array<vec4<i32>, 4>;
 
 struct S {
   arr : ArrayType;
@@ -23,7 +23,7 @@
   var dst : ArrayType;
 
   // Assign from type constructor.
-  dst = ArrayType(1, 2, 3, 3);
+  dst = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3));
 
   // Assign from parameter.
   dst = src_param;
diff --git a/test/array/assign_to_function_var.wgsl.expected.hlsl b/test/array/assign_to_function_var.wgsl.expected.hlsl
index f5c4ff0..eda01ad 100644
--- a/test/array/assign_to_function_var.wgsl.expected.hlsl
+++ b/test/array/assign_to_function_var.wgsl.expected.hlsl
@@ -3,23 +3,20 @@
   return;
 }
 
-struct tint_padded_array_element {
-  int el;
-};
 struct S {
-  tint_padded_array_element arr[4];
+  int4 arr[4];
 };
 
-static tint_padded_array_element src_private[4] = (tint_padded_array_element[4])0;
-groupshared tint_padded_array_element src_workgroup[4];
+static int4 src_private[4] = (int4[4])0;
+groupshared int4 src_workgroup[4];
 cbuffer cbuffer_src_uniform : register(b0, space0) {
   uint4 src_uniform[4];
 };
 RWByteAddressBuffer src_storage : register(u1, space0);
 
-typedef tint_padded_array_element ret_arr_ret[4];
+typedef int4 ret_arr_ret[4];
 ret_arr_ret ret_arr() {
-  const tint_padded_array_element tint_symbol_5[4] = (tint_padded_array_element[4])0;
+  const int4 tint_symbol_5[4] = (int4[4])0;
   return tint_symbol_5;
 }
 
@@ -28,37 +25,37 @@
   return tint_symbol_6;
 }
 
-typedef tint_padded_array_element tint_symbol_1_ret[4];
+typedef int4 tint_symbol_1_ret[4];
 tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
-  tint_padded_array_element arr_1[4] = (tint_padded_array_element[4])0;
+  int4 arr_1[4] = (int4[4])0;
   {
     [loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
       const uint scalar_offset = ((offset + (i * 16u))) / 4;
-      arr_1[i].el = asint(buffer[scalar_offset / 4][scalar_offset % 4]);
+      arr_1[i] = asint(buffer[scalar_offset / 4]);
     }
   }
   return arr_1;
 }
 
-typedef tint_padded_array_element tint_symbol_3_ret[4];
+typedef int4 tint_symbol_3_ret[4];
 tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
-  tint_padded_array_element arr_2[4] = (tint_padded_array_element[4])0;
+  int4 arr_2[4] = (int4[4])0;
   {
     [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
-      arr_2[i_1].el = asint(buffer.Load((offset + (i_1 * 16u))));
+      arr_2[i_1] = asint(buffer.Load4((offset + (i_1 * 16u))));
     }
   }
   return arr_2;
 }
 
-void foo(tint_padded_array_element src_param[4]) {
-  tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0;
-  tint_padded_array_element tint_symbol[4] = (tint_padded_array_element[4])0;
-  const tint_padded_array_element tint_symbol_7[4] = {{1}, {2}, {3}, {3}};
+void foo(int4 src_param[4]) {
+  int4 src_function[4] = (int4[4])0;
+  int4 tint_symbol[4] = (int4[4])0;
+  const int4 tint_symbol_7[4] = {int4((1).xxxx), int4((2).xxxx), int4((3).xxxx), int4((3).xxxx)};
   tint_symbol = tint_symbol_7;
   tint_symbol = src_param;
   tint_symbol = ret_arr();
-  const tint_padded_array_element src_let[4] = (tint_padded_array_element[4])0;
+  const int4 src_let[4] = (int4[4])0;
   tint_symbol = src_let;
   tint_symbol = src_function;
   tint_symbol = src_private;
diff --git a/test/array/assign_to_function_var.wgsl.expected.msl b/test/array/assign_to_function_var.wgsl.expected.msl
index e198433..7abc32d 100644
--- a/test/array/assign_to_function_var.wgsl.expected.msl
+++ b/test/array/assign_to_function_var.wgsl.expected.msl
@@ -1,12 +1,8 @@
 #include <metal_stdlib>
 
 using namespace metal;
-struct tint_padded_array_element {
-  /* 0x0000 */ int el;
-  /* 0x0004 */ int8_t tint_pad[12];
-};
 struct tint_array_wrapper {
-  /* 0x0000 */ tint_padded_array_element arr[4];
+  /* 0x0000 */ int4 arr[4];
 };
 struct S {
   /* 0x0000 */ tint_array_wrapper arr;
@@ -34,7 +30,7 @@
 void foo(tint_array_wrapper src_param, thread tint_array_wrapper* const tint_symbol_3, threadgroup tint_array_wrapper* const tint_symbol_4, const constant S* const tint_symbol_5, device S* const tint_symbol_6) {
   tint_array_wrapper src_function = {};
   tint_array_wrapper dst = {};
-  tint_array_wrapper const tint_symbol_2 = {.arr={{.el=1}, {.el=2}, {.el=3}, {.el=3}}};
+  tint_array_wrapper const tint_symbol_2 = {.arr={int4(1), int4(2), int4(3), int4(3)}};
   dst = tint_symbol_2;
   dst = src_param;
   dst = ret_arr();
diff --git a/test/array/assign_to_function_var.wgsl.expected.spvasm b/test/array/assign_to_function_var.wgsl.expected.spvasm
index 668fc9c..636346d 100644
--- a/test/array/assign_to_function_var.wgsl.expected.spvasm
+++ b/test/array/assign_to_function_var.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 60
+; Bound: 64
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -22,7 +22,7 @@
                OpName %dst "dst"
                OpName %dst_nested "dst_nested"
                OpName %src_nested "src_nested"
-               OpDecorate %_arr_int_uint_4 ArrayStride 16
+               OpDecorate %_arr_v4int_uint_4 ArrayStride 16
                OpDecorate %S Block
                OpMemberDecorate %S 0 Offset 0
                OpDecorate %src_uniform NonWritable
@@ -34,80 +34,84 @@
                OpDecorate %_arr__arr_int_uint_2_uint_3 ArrayStride 8
                OpDecorate %_arr__arr__arr_int_uint_2_uint_3_uint_4 ArrayStride 24
         %int = OpTypeInt 32 1
+      %v4int = OpTypeVector %int 4
        %uint = OpTypeInt 32 0
      %uint_4 = OpConstant %uint 4
-%_arr_int_uint_4 = OpTypeArray %int %uint_4
-%_ptr_Private__arr_int_uint_4 = OpTypePointer Private %_arr_int_uint_4
-          %7 = OpConstantNull %_arr_int_uint_4
-%src_private = OpVariable %_ptr_Private__arr_int_uint_4 Private %7
-%_ptr_Workgroup__arr_int_uint_4 = OpTypePointer Workgroup %_arr_int_uint_4
-%src_workgroup = OpVariable %_ptr_Workgroup__arr_int_uint_4 Workgroup
-          %S = OpTypeStruct %_arr_int_uint_4
+%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
+%_ptr_Private__arr_v4int_uint_4 = OpTypePointer Private %_arr_v4int_uint_4
+          %8 = OpConstantNull %_arr_v4int_uint_4
+%src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8
+%_ptr_Workgroup__arr_v4int_uint_4 = OpTypePointer Workgroup %_arr_v4int_uint_4
+%src_workgroup = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup
+          %S = OpTypeStruct %_arr_v4int_uint_4
 %_ptr_Uniform_S = OpTypePointer Uniform %S
 %src_uniform = OpVariable %_ptr_Uniform_S Uniform
 %_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
 %src_storage = OpVariable %_ptr_StorageBuffer_S StorageBuffer
        %void = OpTypeVoid
-         %15 = OpTypeFunction %void
-         %19 = OpTypeFunction %_arr_int_uint_4
-         %22 = OpTypeFunction %S
-         %25 = OpConstantNull %S
-         %26 = OpTypeFunction %void %_arr_int_uint_4
-%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4
+         %16 = OpTypeFunction %void
+         %20 = OpTypeFunction %_arr_v4int_uint_4
+         %23 = OpTypeFunction %S
+         %26 = OpConstantNull %S
+         %27 = OpTypeFunction %void %_arr_v4int_uint_4
+%_ptr_Function__arr_v4int_uint_4 = OpTypePointer Function %_arr_v4int_uint_4
       %int_1 = OpConstant %int 1
+         %35 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
       %int_2 = OpConstant %int 2
+         %37 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2
       %int_3 = OpConstant %int 3
-         %36 = OpConstantComposite %_arr_int_uint_4 %int_1 %int_2 %int_3 %int_3
+         %39 = OpConstantComposite %v4int %int_3 %int_3 %int_3 %int_3
+         %40 = OpConstantComposite %_arr_v4int_uint_4 %35 %37 %39 %39
      %uint_0 = OpConstant %uint 0
-%_ptr_Uniform__arr_int_uint_4 = OpTypePointer Uniform %_arr_int_uint_4
-%_ptr_StorageBuffer__arr_int_uint_4 = OpTypePointer StorageBuffer %_arr_int_uint_4
+%_ptr_Uniform__arr_v4int_uint_4 = OpTypePointer Uniform %_arr_v4int_uint_4
+%_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
      %uint_2 = OpConstant %uint 2
 %_arr_int_uint_2 = OpTypeArray %int %uint_2
      %uint_3 = OpConstant %uint 3
 %_arr__arr_int_uint_2_uint_3 = OpTypeArray %_arr_int_uint_2 %uint_3
 %_arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypeArray %_arr__arr_int_uint_2_uint_3 %uint_4
 %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Function %_arr__arr__arr_int_uint_2_uint_3_uint_4
-         %57 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
-%unused_entry_point = OpFunction %void None %15
-         %18 = OpLabel
+         %61 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
+%unused_entry_point = OpFunction %void None %16
+         %19 = OpLabel
                OpReturn
                OpFunctionEnd
-    %ret_arr = OpFunction %_arr_int_uint_4 None %19
-         %21 = OpLabel
-               OpReturnValue %7
+    %ret_arr = OpFunction %_arr_v4int_uint_4 None %20
+         %22 = OpLabel
+               OpReturnValue %8
                OpFunctionEnd
-%ret_struct_arr = OpFunction %S None %22
-         %24 = OpLabel
-               OpReturnValue %25
+%ret_struct_arr = OpFunction %S None %23
+         %25 = OpLabel
+               OpReturnValue %26
                OpFunctionEnd
-        %foo = OpFunction %void None %26
-  %src_param = OpFunctionParameter %_arr_int_uint_4
-         %29 = OpLabel
-%src_function = OpVariable %_ptr_Function__arr_int_uint_4 Function %7
-        %dst = OpVariable %_ptr_Function__arr_int_uint_4 Function %7
- %dst_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %57
- %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %57
-               OpStore %dst %36
-               OpStore %dst %src_param
-         %37 = OpFunctionCall %_arr_int_uint_4 %ret_arr
-               OpStore %dst %37
-               OpStore %dst %7
-         %38 = OpLoad %_arr_int_uint_4 %src_function
-               OpStore %dst %38
-         %39 = OpLoad %_arr_int_uint_4 %src_private
-               OpStore %dst %39
-         %40 = OpLoad %_arr_int_uint_4 %src_workgroup
+        %foo = OpFunction %void None %27
+  %src_param = OpFunctionParameter %_arr_v4int_uint_4
+         %30 = OpLabel
+%src_function = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %8
+        %dst = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %8
+ %dst_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %61
+ %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %61
                OpStore %dst %40
-         %41 = OpFunctionCall %S %ret_struct_arr
-         %42 = OpCompositeExtract %_arr_int_uint_4 %41 0
+               OpStore %dst %src_param
+         %41 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
+               OpStore %dst %41
+               OpStore %dst %8
+         %42 = OpLoad %_arr_v4int_uint_4 %src_function
                OpStore %dst %42
-         %45 = OpAccessChain %_ptr_Uniform__arr_int_uint_4 %src_uniform %uint_0
-         %46 = OpLoad %_arr_int_uint_4 %45
+         %43 = OpLoad %_arr_v4int_uint_4 %src_private
+               OpStore %dst %43
+         %44 = OpLoad %_arr_v4int_uint_4 %src_workgroup
+               OpStore %dst %44
+         %45 = OpFunctionCall %S %ret_struct_arr
+         %46 = OpCompositeExtract %_arr_v4int_uint_4 %45 0
                OpStore %dst %46
-         %48 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %src_storage %uint_0
-         %49 = OpLoad %_arr_int_uint_4 %48
-               OpStore %dst %49
-         %59 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
-               OpStore %dst_nested %59
+         %49 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0
+         %50 = OpLoad %_arr_v4int_uint_4 %49
+               OpStore %dst %50
+         %52 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0
+         %53 = OpLoad %_arr_v4int_uint_4 %52
+               OpStore %dst %53
+         %63 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
+               OpStore %dst_nested %63
                OpReturn
                OpFunctionEnd
diff --git a/test/array/assign_to_function_var.wgsl.expected.wgsl b/test/array/assign_to_function_var.wgsl.expected.wgsl
index 0033bea..b001ab6 100644
--- a/test/array/assign_to_function_var.wgsl.expected.wgsl
+++ b/test/array/assign_to_function_var.wgsl.expected.wgsl
@@ -1,4 +1,4 @@
-type ArrayType = @stride(16) array<i32, 4>;
+type ArrayType = array<vec4<i32>, 4>;
 
 struct S {
   arr : ArrayType;
@@ -23,7 +23,7 @@
 fn foo(src_param : ArrayType) {
   var src_function : ArrayType;
   var dst : ArrayType;
-  dst = ArrayType(1, 2, 3, 3);
+  dst = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3));
   dst = src_param;
   dst = ret_arr();
   let src_let : ArrayType = ArrayType();
diff --git a/test/array/assign_to_private_var.wgsl b/test/array/assign_to_private_var.wgsl
index 510da19..c19b3ee 100644
--- a/test/array/assign_to_private_var.wgsl
+++ b/test/array/assign_to_private_var.wgsl
@@ -1,4 +1,4 @@
-type ArrayType = @stride(16) array<i32, 4>;
+type ArrayType = array<vec4<i32>, 4>;
 
 struct S {
   arr : ArrayType;
@@ -24,7 +24,7 @@
   var src_function : ArrayType;
 
   // Assign from type constructor.
-  dst = ArrayType(1, 2, 3, 3);
+  dst = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3));
 
   // Assign from parameter.
   dst = src_param;
diff --git a/test/array/assign_to_private_var.wgsl.expected.hlsl b/test/array/assign_to_private_var.wgsl.expected.hlsl
index 4fd9d45..ddb084e 100644
--- a/test/array/assign_to_private_var.wgsl.expected.hlsl
+++ b/test/array/assign_to_private_var.wgsl.expected.hlsl
@@ -3,25 +3,22 @@
   return;
 }
 
-struct tint_padded_array_element {
-  int el;
-};
 struct S {
-  tint_padded_array_element arr[4];
+  int4 arr[4];
 };
 
-static tint_padded_array_element src_private[4] = (tint_padded_array_element[4])0;
-groupshared tint_padded_array_element src_workgroup[4];
+static int4 src_private[4] = (int4[4])0;
+groupshared int4 src_workgroup[4];
 cbuffer cbuffer_src_uniform : register(b0, space0) {
   uint4 src_uniform[4];
 };
 RWByteAddressBuffer src_storage : register(u1, space0);
-static tint_padded_array_element tint_symbol[4] = (tint_padded_array_element[4])0;
+static int4 tint_symbol[4] = (int4[4])0;
 static int dst_nested[4][3][2] = (int[4][3][2])0;
 
-typedef tint_padded_array_element ret_arr_ret[4];
+typedef int4 ret_arr_ret[4];
 ret_arr_ret ret_arr() {
-  const tint_padded_array_element tint_symbol_5[4] = (tint_padded_array_element[4])0;
+  const int4 tint_symbol_5[4] = (int4[4])0;
   return tint_symbol_5;
 }
 
@@ -30,36 +27,36 @@
   return tint_symbol_6;
 }
 
-typedef tint_padded_array_element tint_symbol_1_ret[4];
+typedef int4 tint_symbol_1_ret[4];
 tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
-  tint_padded_array_element arr_1[4] = (tint_padded_array_element[4])0;
+  int4 arr_1[4] = (int4[4])0;
   {
     [loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
       const uint scalar_offset = ((offset + (i * 16u))) / 4;
-      arr_1[i].el = asint(buffer[scalar_offset / 4][scalar_offset % 4]);
+      arr_1[i] = asint(buffer[scalar_offset / 4]);
     }
   }
   return arr_1;
 }
 
-typedef tint_padded_array_element tint_symbol_3_ret[4];
+typedef int4 tint_symbol_3_ret[4];
 tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
-  tint_padded_array_element arr_2[4] = (tint_padded_array_element[4])0;
+  int4 arr_2[4] = (int4[4])0;
   {
     [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
-      arr_2[i_1].el = asint(buffer.Load((offset + (i_1 * 16u))));
+      arr_2[i_1] = asint(buffer.Load4((offset + (i_1 * 16u))));
     }
   }
   return arr_2;
 }
 
-void foo(tint_padded_array_element src_param[4]) {
-  tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0;
-  const tint_padded_array_element tint_symbol_7[4] = {{1}, {2}, {3}, {3}};
+void foo(int4 src_param[4]) {
+  int4 src_function[4] = (int4[4])0;
+  const int4 tint_symbol_7[4] = {int4((1).xxxx), int4((2).xxxx), int4((3).xxxx), int4((3).xxxx)};
   tint_symbol = tint_symbol_7;
   tint_symbol = src_param;
   tint_symbol = ret_arr();
-  const tint_padded_array_element src_let[4] = (tint_padded_array_element[4])0;
+  const int4 src_let[4] = (int4[4])0;
   tint_symbol = src_let;
   tint_symbol = src_function;
   tint_symbol = src_private;
diff --git a/test/array/assign_to_private_var.wgsl.expected.msl b/test/array/assign_to_private_var.wgsl.expected.msl
index dd91dca..7d8a3ab 100644
--- a/test/array/assign_to_private_var.wgsl.expected.msl
+++ b/test/array/assign_to_private_var.wgsl.expected.msl
@@ -1,12 +1,8 @@
 #include <metal_stdlib>
 
 using namespace metal;
-struct tint_padded_array_element {
-  /* 0x0000 */ int el;
-  /* 0x0004 */ int8_t tint_pad[12];
-};
 struct tint_array_wrapper {
-  /* 0x0000 */ tint_padded_array_element arr[4];
+  /* 0x0000 */ int4 arr[4];
 };
 struct S {
   /* 0x0000 */ tint_array_wrapper arr;
@@ -33,7 +29,7 @@
 
 void foo(tint_array_wrapper src_param, thread tint_array_wrapper* const tint_symbol_3, thread tint_array_wrapper* const tint_symbol_4, threadgroup tint_array_wrapper* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7, thread tint_array_wrapper_1* const tint_symbol_8) {
   tint_array_wrapper src_function = {};
-  tint_array_wrapper const tint_symbol_2 = {.arr={{.el=1}, {.el=2}, {.el=3}, {.el=3}}};
+  tint_array_wrapper const tint_symbol_2 = {.arr={int4(1), int4(2), int4(3), int4(3)}};
   *(tint_symbol_3) = tint_symbol_2;
   *(tint_symbol_3) = src_param;
   *(tint_symbol_3) = ret_arr();
diff --git a/test/array/assign_to_private_var.wgsl.expected.spvasm b/test/array/assign_to_private_var.wgsl.expected.spvasm
index 258a540..3427ea7 100644
--- a/test/array/assign_to_private_var.wgsl.expected.spvasm
+++ b/test/array/assign_to_private_var.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 61
+; Bound: 65
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -22,7 +22,7 @@
                OpName %src_param "src_param"
                OpName %src_function "src_function"
                OpName %src_nested "src_nested"
-               OpDecorate %_arr_int_uint_4 ArrayStride 16
+               OpDecorate %_arr_v4int_uint_4 ArrayStride 16
                OpDecorate %S Block
                OpMemberDecorate %S 0 Offset 0
                OpDecorate %src_uniform NonWritable
@@ -34,81 +34,85 @@
                OpDecorate %_arr__arr_int_uint_2_uint_3 ArrayStride 8
                OpDecorate %_arr__arr__arr_int_uint_2_uint_3_uint_4 ArrayStride 24
         %int = OpTypeInt 32 1
+      %v4int = OpTypeVector %int 4
        %uint = OpTypeInt 32 0
      %uint_4 = OpConstant %uint 4
-%_arr_int_uint_4 = OpTypeArray %int %uint_4
-%_ptr_Private__arr_int_uint_4 = OpTypePointer Private %_arr_int_uint_4
-          %7 = OpConstantNull %_arr_int_uint_4
-%src_private = OpVariable %_ptr_Private__arr_int_uint_4 Private %7
-%_ptr_Workgroup__arr_int_uint_4 = OpTypePointer Workgroup %_arr_int_uint_4
-%src_workgroup = OpVariable %_ptr_Workgroup__arr_int_uint_4 Workgroup
-          %S = OpTypeStruct %_arr_int_uint_4
+%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
+%_ptr_Private__arr_v4int_uint_4 = OpTypePointer Private %_arr_v4int_uint_4
+          %8 = OpConstantNull %_arr_v4int_uint_4
+%src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8
+%_ptr_Workgroup__arr_v4int_uint_4 = OpTypePointer Workgroup %_arr_v4int_uint_4
+%src_workgroup = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup
+          %S = OpTypeStruct %_arr_v4int_uint_4
 %_ptr_Uniform_S = OpTypePointer Uniform %S
 %src_uniform = OpVariable %_ptr_Uniform_S Uniform
 %_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
 %src_storage = OpVariable %_ptr_StorageBuffer_S StorageBuffer
-        %dst = OpVariable %_ptr_Private__arr_int_uint_4 Private %7
+        %dst = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8
      %uint_2 = OpConstant %uint 2
 %_arr_int_uint_2 = OpTypeArray %int %uint_2
      %uint_3 = OpConstant %uint 3
 %_arr__arr_int_uint_2_uint_3 = OpTypeArray %_arr_int_uint_2 %uint_3
 %_arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypeArray %_arr__arr_int_uint_2_uint_3 %uint_4
 %_ptr_Private__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Private %_arr__arr__arr_int_uint_2_uint_3_uint_4
-         %23 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
- %dst_nested = OpVariable %_ptr_Private__arr__arr__arr_int_uint_2_uint_3_uint_4 Private %23
+         %24 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
+ %dst_nested = OpVariable %_ptr_Private__arr__arr__arr_int_uint_2_uint_3_uint_4 Private %24
        %void = OpTypeVoid
-         %24 = OpTypeFunction %void
-         %28 = OpTypeFunction %_arr_int_uint_4
-         %31 = OpTypeFunction %S
-         %34 = OpConstantNull %S
-         %35 = OpTypeFunction %void %_arr_int_uint_4
-%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4
+         %25 = OpTypeFunction %void
+         %29 = OpTypeFunction %_arr_v4int_uint_4
+         %32 = OpTypeFunction %S
+         %35 = OpConstantNull %S
+         %36 = OpTypeFunction %void %_arr_v4int_uint_4
+%_ptr_Function__arr_v4int_uint_4 = OpTypePointer Function %_arr_v4int_uint_4
       %int_1 = OpConstant %int 1
+         %43 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
       %int_2 = OpConstant %int 2
+         %45 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2
       %int_3 = OpConstant %int 3
-         %44 = OpConstantComposite %_arr_int_uint_4 %int_1 %int_2 %int_3 %int_3
+         %47 = OpConstantComposite %v4int %int_3 %int_3 %int_3 %int_3
+         %48 = OpConstantComposite %_arr_v4int_uint_4 %43 %45 %47 %47
      %uint_0 = OpConstant %uint 0
-%_ptr_Uniform__arr_int_uint_4 = OpTypePointer Uniform %_arr_int_uint_4
-%_ptr_StorageBuffer__arr_int_uint_4 = OpTypePointer StorageBuffer %_arr_int_uint_4
+%_ptr_Uniform__arr_v4int_uint_4 = OpTypePointer Uniform %_arr_v4int_uint_4
+%_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
 %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Function %_arr__arr__arr_int_uint_2_uint_3_uint_4
-%unused_entry_point = OpFunction %void None %24
-         %27 = OpLabel
+%unused_entry_point = OpFunction %void None %25
+         %28 = OpLabel
                OpReturn
                OpFunctionEnd
-    %ret_arr = OpFunction %_arr_int_uint_4 None %28
-         %30 = OpLabel
-               OpReturnValue %7
+    %ret_arr = OpFunction %_arr_v4int_uint_4 None %29
+         %31 = OpLabel
+               OpReturnValue %8
                OpFunctionEnd
-%ret_struct_arr = OpFunction %S None %31
-         %33 = OpLabel
-               OpReturnValue %34
+%ret_struct_arr = OpFunction %S None %32
+         %34 = OpLabel
+               OpReturnValue %35
                OpFunctionEnd
-        %foo = OpFunction %void None %35
-  %src_param = OpFunctionParameter %_arr_int_uint_4
-         %38 = OpLabel
-%src_function = OpVariable %_ptr_Function__arr_int_uint_4 Function %7
- %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %23
-               OpStore %dst %44
-               OpStore %dst %src_param
-         %45 = OpFunctionCall %_arr_int_uint_4 %ret_arr
-               OpStore %dst %45
-               OpStore %dst %7
-         %46 = OpLoad %_arr_int_uint_4 %src_function
-               OpStore %dst %46
-         %47 = OpLoad %_arr_int_uint_4 %src_private
-               OpStore %dst %47
-         %48 = OpLoad %_arr_int_uint_4 %src_workgroup
+        %foo = OpFunction %void None %36
+  %src_param = OpFunctionParameter %_arr_v4int_uint_4
+         %39 = OpLabel
+%src_function = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %8
+ %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %24
                OpStore %dst %48
-         %49 = OpFunctionCall %S %ret_struct_arr
-         %50 = OpCompositeExtract %_arr_int_uint_4 %49 0
+               OpStore %dst %src_param
+         %49 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
+               OpStore %dst %49
+               OpStore %dst %8
+         %50 = OpLoad %_arr_v4int_uint_4 %src_function
                OpStore %dst %50
-         %53 = OpAccessChain %_ptr_Uniform__arr_int_uint_4 %src_uniform %uint_0
-         %54 = OpLoad %_arr_int_uint_4 %53
+         %51 = OpLoad %_arr_v4int_uint_4 %src_private
+               OpStore %dst %51
+         %52 = OpLoad %_arr_v4int_uint_4 %src_workgroup
+               OpStore %dst %52
+         %53 = OpFunctionCall %S %ret_struct_arr
+         %54 = OpCompositeExtract %_arr_v4int_uint_4 %53 0
                OpStore %dst %54
-         %56 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %src_storage %uint_0
-         %57 = OpLoad %_arr_int_uint_4 %56
-               OpStore %dst %57
-         %60 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
-               OpStore %dst_nested %60
+         %57 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0
+         %58 = OpLoad %_arr_v4int_uint_4 %57
+               OpStore %dst %58
+         %60 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0
+         %61 = OpLoad %_arr_v4int_uint_4 %60
+               OpStore %dst %61
+         %64 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
+               OpStore %dst_nested %64
                OpReturn
                OpFunctionEnd
diff --git a/test/array/assign_to_private_var.wgsl.expected.wgsl b/test/array/assign_to_private_var.wgsl.expected.wgsl
index fd2057b..7f88483 100644
--- a/test/array/assign_to_private_var.wgsl.expected.wgsl
+++ b/test/array/assign_to_private_var.wgsl.expected.wgsl
@@ -1,4 +1,4 @@
-type ArrayType = @stride(16) array<i32, 4>;
+type ArrayType = array<vec4<i32>, 4>;
 
 struct S {
   arr : ArrayType;
@@ -26,7 +26,7 @@
 
 fn foo(src_param : ArrayType) {
   var src_function : ArrayType;
-  dst = ArrayType(1, 2, 3, 3);
+  dst = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3));
   dst = src_param;
   dst = ret_arr();
   let src_let : ArrayType = ArrayType();
diff --git a/test/array/assign_to_storage_var.wgsl b/test/array/assign_to_storage_var.wgsl
index 6d001d8..9ed30b2 100644
--- a/test/array/assign_to_storage_var.wgsl
+++ b/test/array/assign_to_storage_var.wgsl
@@ -1,4 +1,4 @@
-type ArrayType = @stride(16) array<i32, 4>;
+type ArrayType = array<vec4<i32>, 4>;
 
 struct S {
   arr : ArrayType;
@@ -28,7 +28,7 @@
   var src_function : ArrayType;
 
   // Assign from type constructor.
-  dst.arr = ArrayType(1, 2, 3, 3);
+  dst.arr = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3));
 
   // Assign from parameter.
   dst.arr = src_param;
diff --git a/test/array/assign_to_storage_var.wgsl.expected.hlsl b/test/array/assign_to_storage_var.wgsl.expected.hlsl
index 8e9d0b6..c940b98 100644
--- a/test/array/assign_to_storage_var.wgsl.expected.hlsl
+++ b/test/array/assign_to_storage_var.wgsl.expected.hlsl
@@ -3,15 +3,12 @@
   return;
 }
 
-struct tint_padded_array_element {
-  int el;
-};
 struct S {
-  tint_padded_array_element arr[4];
+  int4 arr[4];
 };
 
-static tint_padded_array_element src_private[4] = (tint_padded_array_element[4])0;
-groupshared tint_padded_array_element src_workgroup[4];
+static int4 src_private[4] = (int4[4])0;
+groupshared int4 src_workgroup[4];
 cbuffer cbuffer_src_uniform : register(b0, space0) {
   uint4 src_uniform[4];
 };
@@ -19,9 +16,9 @@
 RWByteAddressBuffer tint_symbol : register(u2, space0);
 RWByteAddressBuffer dst_nested : register(u3, space0);
 
-typedef tint_padded_array_element ret_arr_ret[4];
+typedef int4 ret_arr_ret[4];
 ret_arr_ret ret_arr() {
-  const tint_padded_array_element tint_symbol_11[4] = (tint_padded_array_element[4])0;
+  const int4 tint_symbol_11[4] = (int4[4])0;
   return tint_symbol_11;
 }
 
@@ -30,33 +27,33 @@
   return tint_symbol_12;
 }
 
-void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[4]) {
-  tint_padded_array_element array[4] = value;
+void tint_symbol_1(RWByteAddressBuffer buffer, uint offset, int4 value[4]) {
+  int4 array[4] = value;
   {
     [loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
-      buffer.Store((offset + (i * 16u)), asuint(array[i].el));
+      buffer.Store4((offset + (i * 16u)), asuint(array[i]));
     }
   }
 }
 
-typedef tint_padded_array_element tint_symbol_3_ret[4];
+typedef int4 tint_symbol_3_ret[4];
 tint_symbol_3_ret tint_symbol_3(uint4 buffer[4], uint offset) {
-  tint_padded_array_element arr_1[4] = (tint_padded_array_element[4])0;
+  int4 arr_1[4] = (int4[4])0;
   {
     [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
       const uint scalar_offset = ((offset + (i_1 * 16u))) / 4;
-      arr_1[i_1].el = asint(buffer[scalar_offset / 4][scalar_offset % 4]);
+      arr_1[i_1] = asint(buffer[scalar_offset / 4]);
     }
   }
   return arr_1;
 }
 
-typedef tint_padded_array_element tint_symbol_5_ret[4];
+typedef int4 tint_symbol_5_ret[4];
 tint_symbol_5_ret tint_symbol_5(RWByteAddressBuffer buffer, uint offset) {
-  tint_padded_array_element arr_2[4] = (tint_padded_array_element[4])0;
+  int4 arr_2[4] = (int4[4])0;
   {
     [loop] for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) {
-      arr_2[i_2].el = asint(buffer.Load((offset + (i_2 * 16u))));
+      arr_2[i_2] = asint(buffer.Load4((offset + (i_2 * 16u))));
     }
   }
   return arr_2;
@@ -89,13 +86,13 @@
   }
 }
 
-void foo(tint_padded_array_element src_param[4]) {
-  tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0;
-  const tint_padded_array_element tint_symbol_13[4] = {{1}, {2}, {3}, {3}};
+void foo(int4 src_param[4]) {
+  int4 src_function[4] = (int4[4])0;
+  const int4 tint_symbol_13[4] = {int4((1).xxxx), int4((2).xxxx), int4((3).xxxx), int4((3).xxxx)};
   tint_symbol_1(tint_symbol, 0u, tint_symbol_13);
   tint_symbol_1(tint_symbol, 0u, src_param);
   tint_symbol_1(tint_symbol, 0u, ret_arr());
-  const tint_padded_array_element src_let[4] = (tint_padded_array_element[4])0;
+  const int4 src_let[4] = (int4[4])0;
   tint_symbol_1(tint_symbol, 0u, src_let);
   tint_symbol_1(tint_symbol, 0u, src_function);
   tint_symbol_1(tint_symbol, 0u, src_private);
diff --git a/test/array/assign_to_storage_var.wgsl.expected.msl b/test/array/assign_to_storage_var.wgsl.expected.msl
index c0d7dba..c1a31f9 100644
--- a/test/array/assign_to_storage_var.wgsl.expected.msl
+++ b/test/array/assign_to_storage_var.wgsl.expected.msl
@@ -1,12 +1,8 @@
 #include <metal_stdlib>
 
 using namespace metal;
-struct tint_padded_array_element {
-  /* 0x0000 */ int el;
-  /* 0x0004 */ int8_t tint_pad[12];
-};
 struct tint_array_wrapper {
-  /* 0x0000 */ tint_padded_array_element arr[4];
+  /* 0x0000 */ int4 arr[4];
 };
 struct S {
   /* 0x0000 */ tint_array_wrapper arr;
@@ -36,7 +32,7 @@
 
 void foo(tint_array_wrapper src_param, device S* const tint_symbol_3, thread tint_array_wrapper* const tint_symbol_4, threadgroup tint_array_wrapper* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7, device S_nested* const tint_symbol_8) {
   tint_array_wrapper src_function = {};
-  tint_array_wrapper const tint_symbol_2 = {.arr={{.el=1}, {.el=2}, {.el=3}, {.el=3}}};
+  tint_array_wrapper const tint_symbol_2 = {.arr={int4(1), int4(2), int4(3), int4(3)}};
   (*(tint_symbol_3)).arr = tint_symbol_2;
   (*(tint_symbol_3)).arr = src_param;
   (*(tint_symbol_3)).arr = ret_arr();
diff --git a/test/array/assign_to_storage_var.wgsl.expected.spvasm b/test/array/assign_to_storage_var.wgsl.expected.spvasm
index 1764efa..e4da62a 100644
--- a/test/array/assign_to_storage_var.wgsl.expected.spvasm
+++ b/test/array/assign_to_storage_var.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 74
+; Bound: 78
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -24,7 +24,7 @@
                OpName %src_param "src_param"
                OpName %src_function "src_function"
                OpName %src_nested "src_nested"
-               OpDecorate %_arr_int_uint_4 ArrayStride 16
+               OpDecorate %_arr_v4int_uint_4 ArrayStride 16
                OpDecorate %S Block
                OpMemberDecorate %S 0 Offset 0
                OpDecorate %src_uniform NonWritable
@@ -42,15 +42,16 @@
                OpDecorate %dst_nested DescriptorSet 0
                OpDecorate %dst_nested Binding 3
         %int = OpTypeInt 32 1
+      %v4int = OpTypeVector %int 4
        %uint = OpTypeInt 32 0
      %uint_4 = OpConstant %uint 4
-%_arr_int_uint_4 = OpTypeArray %int %uint_4
-%_ptr_Private__arr_int_uint_4 = OpTypePointer Private %_arr_int_uint_4
-          %7 = OpConstantNull %_arr_int_uint_4
-%src_private = OpVariable %_ptr_Private__arr_int_uint_4 Private %7
-%_ptr_Workgroup__arr_int_uint_4 = OpTypePointer Workgroup %_arr_int_uint_4
-%src_workgroup = OpVariable %_ptr_Workgroup__arr_int_uint_4 Workgroup
-          %S = OpTypeStruct %_arr_int_uint_4
+%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
+%_ptr_Private__arr_v4int_uint_4 = OpTypePointer Private %_arr_v4int_uint_4
+          %8 = OpConstantNull %_arr_v4int_uint_4
+%src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8
+%_ptr_Workgroup__arr_v4int_uint_4 = OpTypePointer Workgroup %_arr_v4int_uint_4
+%src_workgroup = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup
+          %S = OpTypeStruct %_arr_v4int_uint_4
 %_ptr_Uniform_S = OpTypePointer Uniform %S
 %src_uniform = OpVariable %_ptr_Uniform_S Uniform
 %_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
@@ -65,71 +66,74 @@
 %_ptr_StorageBuffer_S_nested = OpTypePointer StorageBuffer %S_nested
  %dst_nested = OpVariable %_ptr_StorageBuffer_S_nested StorageBuffer
        %void = OpTypeVoid
-         %24 = OpTypeFunction %void
-         %28 = OpTypeFunction %_arr_int_uint_4
-         %31 = OpTypeFunction %S
-         %34 = OpConstantNull %S
-         %35 = OpTypeFunction %void %_arr_int_uint_4
-%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4
+         %25 = OpTypeFunction %void
+         %29 = OpTypeFunction %_arr_v4int_uint_4
+         %32 = OpTypeFunction %S
+         %35 = OpConstantNull %S
+         %36 = OpTypeFunction %void %_arr_v4int_uint_4
+%_ptr_Function__arr_v4int_uint_4 = OpTypePointer Function %_arr_v4int_uint_4
      %uint_0 = OpConstant %uint 0
-%_ptr_StorageBuffer__arr_int_uint_4 = OpTypePointer StorageBuffer %_arr_int_uint_4
+%_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
       %int_1 = OpConstant %int 1
+         %46 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
       %int_2 = OpConstant %int 2
+         %48 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2
       %int_3 = OpConstant %int 3
-         %47 = OpConstantComposite %_arr_int_uint_4 %int_1 %int_2 %int_3 %int_3
-%_ptr_Uniform__arr_int_uint_4 = OpTypePointer Uniform %_arr_int_uint_4
+         %50 = OpConstantComposite %v4int %int_3 %int_3 %int_3 %int_3
+         %51 = OpConstantComposite %_arr_v4int_uint_4 %46 %48 %50 %50
+%_ptr_Uniform__arr_v4int_uint_4 = OpTypePointer Uniform %_arr_v4int_uint_4
 %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Function %_arr__arr__arr_int_uint_2_uint_3_uint_4
-         %70 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
+         %74 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
 %_ptr_StorageBuffer__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer StorageBuffer %_arr__arr__arr_int_uint_2_uint_3_uint_4
-%unused_entry_point = OpFunction %void None %24
-         %27 = OpLabel
+%unused_entry_point = OpFunction %void None %25
+         %28 = OpLabel
                OpReturn
                OpFunctionEnd
-    %ret_arr = OpFunction %_arr_int_uint_4 None %28
-         %30 = OpLabel
-               OpReturnValue %7
+    %ret_arr = OpFunction %_arr_v4int_uint_4 None %29
+         %31 = OpLabel
+               OpReturnValue %8
                OpFunctionEnd
-%ret_struct_arr = OpFunction %S None %31
-         %33 = OpLabel
-               OpReturnValue %34
+%ret_struct_arr = OpFunction %S None %32
+         %34 = OpLabel
+               OpReturnValue %35
                OpFunctionEnd
-        %foo = OpFunction %void None %35
-  %src_param = OpFunctionParameter %_arr_int_uint_4
-         %38 = OpLabel
-%src_function = OpVariable %_ptr_Function__arr_int_uint_4 Function %7
- %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %70
-         %43 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0
-               OpStore %43 %47
-         %48 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0
-               OpStore %48 %src_param
-         %49 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0
-         %50 = OpFunctionCall %_arr_int_uint_4 %ret_arr
-               OpStore %49 %50
-         %51 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0
-               OpStore %51 %7
-         %52 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0
-         %53 = OpLoad %_arr_int_uint_4 %src_function
-               OpStore %52 %53
-         %54 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0
-         %55 = OpLoad %_arr_int_uint_4 %src_private
-               OpStore %54 %55
-         %56 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0
-         %57 = OpLoad %_arr_int_uint_4 %src_workgroup
+        %foo = OpFunction %void None %36
+  %src_param = OpFunctionParameter %_arr_v4int_uint_4
+         %39 = OpLabel
+%src_function = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %8
+ %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %74
+         %44 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
+               OpStore %44 %51
+         %52 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
+               OpStore %52 %src_param
+         %53 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
+         %54 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
+               OpStore %53 %54
+         %55 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
+               OpStore %55 %8
+         %56 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
+         %57 = OpLoad %_arr_v4int_uint_4 %src_function
                OpStore %56 %57
-         %58 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0
-         %59 = OpFunctionCall %S %ret_struct_arr
-         %60 = OpCompositeExtract %_arr_int_uint_4 %59 0
-               OpStore %58 %60
-         %61 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0
-         %63 = OpAccessChain %_ptr_Uniform__arr_int_uint_4 %src_uniform %uint_0
-         %64 = OpLoad %_arr_int_uint_4 %63
-               OpStore %61 %64
-         %65 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %dst %uint_0
-         %66 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %src_storage %uint_0
-         %67 = OpLoad %_arr_int_uint_4 %66
-               OpStore %65 %67
-         %72 = OpAccessChain %_ptr_StorageBuffer__arr__arr__arr_int_uint_2_uint_3_uint_4 %dst_nested %uint_0
-         %73 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
-               OpStore %72 %73
+         %58 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
+         %59 = OpLoad %_arr_v4int_uint_4 %src_private
+               OpStore %58 %59
+         %60 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
+         %61 = OpLoad %_arr_v4int_uint_4 %src_workgroup
+               OpStore %60 %61
+         %62 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
+         %63 = OpFunctionCall %S %ret_struct_arr
+         %64 = OpCompositeExtract %_arr_v4int_uint_4 %63 0
+               OpStore %62 %64
+         %65 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
+         %67 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0
+         %68 = OpLoad %_arr_v4int_uint_4 %67
+               OpStore %65 %68
+         %69 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %dst %uint_0
+         %70 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0
+         %71 = OpLoad %_arr_v4int_uint_4 %70
+               OpStore %69 %71
+         %76 = OpAccessChain %_ptr_StorageBuffer__arr__arr__arr_int_uint_2_uint_3_uint_4 %dst_nested %uint_0
+         %77 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
+               OpStore %76 %77
                OpReturn
                OpFunctionEnd
diff --git a/test/array/assign_to_storage_var.wgsl.expected.wgsl b/test/array/assign_to_storage_var.wgsl.expected.wgsl
index 88902fa..fe47458 100644
--- a/test/array/assign_to_storage_var.wgsl.expected.wgsl
+++ b/test/array/assign_to_storage_var.wgsl.expected.wgsl
@@ -1,4 +1,4 @@
-type ArrayType = @stride(16) array<i32, 4>;
+type ArrayType = array<vec4<i32>, 4>;
 
 struct S {
   arr : ArrayType;
@@ -30,7 +30,7 @@
 
 fn foo(src_param : ArrayType) {
   var src_function : ArrayType;
-  dst.arr = ArrayType(1, 2, 3, 3);
+  dst.arr = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3));
   dst.arr = src_param;
   dst.arr = ret_arr();
   let src_let : ArrayType = ArrayType();
diff --git a/test/array/assign_to_workgroup_var.wgsl b/test/array/assign_to_workgroup_var.wgsl
index 9fb94eb..1615348 100644
--- a/test/array/assign_to_workgroup_var.wgsl
+++ b/test/array/assign_to_workgroup_var.wgsl
@@ -1,4 +1,4 @@
-type ArrayType = @stride(16) array<i32, 4>;
+type ArrayType = array<vec4<i32>, 4>;
 
 struct S {
   arr : ArrayType;
@@ -24,7 +24,7 @@
   var src_function : ArrayType;
 
   // Assign from type constructor.
-  dst = ArrayType(1, 2, 3, 3);
+  dst = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3));
 
   // Assign from parameter.
   dst = src_param;
diff --git a/test/array/assign_to_workgroup_var.wgsl.expected.hlsl b/test/array/assign_to_workgroup_var.wgsl.expected.hlsl
index b60b974..4243ec1 100644
--- a/test/array/assign_to_workgroup_var.wgsl.expected.hlsl
+++ b/test/array/assign_to_workgroup_var.wgsl.expected.hlsl
@@ -3,25 +3,22 @@
   return;
 }
 
-struct tint_padded_array_element {
-  int el;
-};
 struct S {
-  tint_padded_array_element arr[4];
+  int4 arr[4];
 };
 
-static tint_padded_array_element src_private[4] = (tint_padded_array_element[4])0;
-groupshared tint_padded_array_element src_workgroup[4];
+static int4 src_private[4] = (int4[4])0;
+groupshared int4 src_workgroup[4];
 cbuffer cbuffer_src_uniform : register(b0, space0) {
   uint4 src_uniform[4];
 };
 RWByteAddressBuffer src_storage : register(u1, space0);
-groupshared tint_padded_array_element tint_symbol[4];
+groupshared int4 tint_symbol[4];
 groupshared int dst_nested[4][3][2];
 
-typedef tint_padded_array_element ret_arr_ret[4];
+typedef int4 ret_arr_ret[4];
 ret_arr_ret ret_arr() {
-  const tint_padded_array_element tint_symbol_5[4] = (tint_padded_array_element[4])0;
+  const int4 tint_symbol_5[4] = (int4[4])0;
   return tint_symbol_5;
 }
 
@@ -30,36 +27,36 @@
   return tint_symbol_6;
 }
 
-typedef tint_padded_array_element tint_symbol_1_ret[4];
+typedef int4 tint_symbol_1_ret[4];
 tint_symbol_1_ret tint_symbol_1(uint4 buffer[4], uint offset) {
-  tint_padded_array_element arr_1[4] = (tint_padded_array_element[4])0;
+  int4 arr_1[4] = (int4[4])0;
   {
     [loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
       const uint scalar_offset = ((offset + (i * 16u))) / 4;
-      arr_1[i].el = asint(buffer[scalar_offset / 4][scalar_offset % 4]);
+      arr_1[i] = asint(buffer[scalar_offset / 4]);
     }
   }
   return arr_1;
 }
 
-typedef tint_padded_array_element tint_symbol_3_ret[4];
+typedef int4 tint_symbol_3_ret[4];
 tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) {
-  tint_padded_array_element arr_2[4] = (tint_padded_array_element[4])0;
+  int4 arr_2[4] = (int4[4])0;
   {
     [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
-      arr_2[i_1].el = asint(buffer.Load((offset + (i_1 * 16u))));
+      arr_2[i_1] = asint(buffer.Load4((offset + (i_1 * 16u))));
     }
   }
   return arr_2;
 }
 
-void foo(tint_padded_array_element src_param[4]) {
-  tint_padded_array_element src_function[4] = (tint_padded_array_element[4])0;
-  const tint_padded_array_element tint_symbol_7[4] = {{1}, {2}, {3}, {3}};
+void foo(int4 src_param[4]) {
+  int4 src_function[4] = (int4[4])0;
+  const int4 tint_symbol_7[4] = {int4((1).xxxx), int4((2).xxxx), int4((3).xxxx), int4((3).xxxx)};
   tint_symbol = tint_symbol_7;
   tint_symbol = src_param;
   tint_symbol = ret_arr();
-  const tint_padded_array_element src_let[4] = (tint_padded_array_element[4])0;
+  const int4 src_let[4] = (int4[4])0;
   tint_symbol = src_let;
   tint_symbol = src_function;
   tint_symbol = src_private;
diff --git a/test/array/assign_to_workgroup_var.wgsl.expected.msl b/test/array/assign_to_workgroup_var.wgsl.expected.msl
index 0511506..2c45f26 100644
--- a/test/array/assign_to_workgroup_var.wgsl.expected.msl
+++ b/test/array/assign_to_workgroup_var.wgsl.expected.msl
@@ -1,12 +1,8 @@
 #include <metal_stdlib>
 
 using namespace metal;
-struct tint_padded_array_element {
-  /* 0x0000 */ int el;
-  /* 0x0004 */ int8_t tint_pad[12];
-};
 struct tint_array_wrapper {
-  /* 0x0000 */ tint_padded_array_element arr[4];
+  /* 0x0000 */ int4 arr[4];
 };
 struct S {
   /* 0x0000 */ tint_array_wrapper arr;
@@ -33,7 +29,7 @@
 
 void foo(tint_array_wrapper src_param, threadgroup tint_array_wrapper* const tint_symbol_3, thread tint_array_wrapper* const tint_symbol_4, threadgroup tint_array_wrapper* const tint_symbol_5, const constant S* const tint_symbol_6, device S* const tint_symbol_7, threadgroup tint_array_wrapper_1* const tint_symbol_8) {
   tint_array_wrapper src_function = {};
-  tint_array_wrapper const tint_symbol_2 = {.arr={{.el=1}, {.el=2}, {.el=3}, {.el=3}}};
+  tint_array_wrapper const tint_symbol_2 = {.arr={int4(1), int4(2), int4(3), int4(3)}};
   *(tint_symbol_3) = tint_symbol_2;
   *(tint_symbol_3) = src_param;
   *(tint_symbol_3) = ret_arr();
diff --git a/test/array/assign_to_workgroup_var.wgsl.expected.spvasm b/test/array/assign_to_workgroup_var.wgsl.expected.spvasm
index 5f46701..a834fb2 100644
--- a/test/array/assign_to_workgroup_var.wgsl.expected.spvasm
+++ b/test/array/assign_to_workgroup_var.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 61
+; Bound: 65
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -22,7 +22,7 @@
                OpName %src_param "src_param"
                OpName %src_function "src_function"
                OpName %src_nested "src_nested"
-               OpDecorate %_arr_int_uint_4 ArrayStride 16
+               OpDecorate %_arr_v4int_uint_4 ArrayStride 16
                OpDecorate %S Block
                OpMemberDecorate %S 0 Offset 0
                OpDecorate %src_uniform NonWritable
@@ -34,20 +34,21 @@
                OpDecorate %_arr__arr_int_uint_2_uint_3 ArrayStride 8
                OpDecorate %_arr__arr__arr_int_uint_2_uint_3_uint_4 ArrayStride 24
         %int = OpTypeInt 32 1
+      %v4int = OpTypeVector %int 4
        %uint = OpTypeInt 32 0
      %uint_4 = OpConstant %uint 4
-%_arr_int_uint_4 = OpTypeArray %int %uint_4
-%_ptr_Private__arr_int_uint_4 = OpTypePointer Private %_arr_int_uint_4
-          %7 = OpConstantNull %_arr_int_uint_4
-%src_private = OpVariable %_ptr_Private__arr_int_uint_4 Private %7
-%_ptr_Workgroup__arr_int_uint_4 = OpTypePointer Workgroup %_arr_int_uint_4
-%src_workgroup = OpVariable %_ptr_Workgroup__arr_int_uint_4 Workgroup
-          %S = OpTypeStruct %_arr_int_uint_4
+%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
+%_ptr_Private__arr_v4int_uint_4 = OpTypePointer Private %_arr_v4int_uint_4
+          %8 = OpConstantNull %_arr_v4int_uint_4
+%src_private = OpVariable %_ptr_Private__arr_v4int_uint_4 Private %8
+%_ptr_Workgroup__arr_v4int_uint_4 = OpTypePointer Workgroup %_arr_v4int_uint_4
+%src_workgroup = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup
+          %S = OpTypeStruct %_arr_v4int_uint_4
 %_ptr_Uniform_S = OpTypePointer Uniform %S
 %src_uniform = OpVariable %_ptr_Uniform_S Uniform
 %_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
 %src_storage = OpVariable %_ptr_StorageBuffer_S StorageBuffer
-        %dst = OpVariable %_ptr_Workgroup__arr_int_uint_4 Workgroup
+        %dst = OpVariable %_ptr_Workgroup__arr_v4int_uint_4 Workgroup
      %uint_2 = OpConstant %uint 2
 %_arr_int_uint_2 = OpTypeArray %int %uint_2
      %uint_3 = OpConstant %uint 3
@@ -56,59 +57,62 @@
 %_ptr_Workgroup__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Workgroup %_arr__arr__arr_int_uint_2_uint_3_uint_4
  %dst_nested = OpVariable %_ptr_Workgroup__arr__arr__arr_int_uint_2_uint_3_uint_4 Workgroup
        %void = OpTypeVoid
-         %23 = OpTypeFunction %void
-         %27 = OpTypeFunction %_arr_int_uint_4
-         %30 = OpTypeFunction %S
-         %33 = OpConstantNull %S
-         %34 = OpTypeFunction %void %_arr_int_uint_4
-%_ptr_Function__arr_int_uint_4 = OpTypePointer Function %_arr_int_uint_4
+         %24 = OpTypeFunction %void
+         %28 = OpTypeFunction %_arr_v4int_uint_4
+         %31 = OpTypeFunction %S
+         %34 = OpConstantNull %S
+         %35 = OpTypeFunction %void %_arr_v4int_uint_4
+%_ptr_Function__arr_v4int_uint_4 = OpTypePointer Function %_arr_v4int_uint_4
       %int_1 = OpConstant %int 1
+         %42 = OpConstantComposite %v4int %int_1 %int_1 %int_1 %int_1
       %int_2 = OpConstant %int 2
+         %44 = OpConstantComposite %v4int %int_2 %int_2 %int_2 %int_2
       %int_3 = OpConstant %int 3
-         %43 = OpConstantComposite %_arr_int_uint_4 %int_1 %int_2 %int_3 %int_3
+         %46 = OpConstantComposite %v4int %int_3 %int_3 %int_3 %int_3
+         %47 = OpConstantComposite %_arr_v4int_uint_4 %42 %44 %46 %46
      %uint_0 = OpConstant %uint 0
-%_ptr_Uniform__arr_int_uint_4 = OpTypePointer Uniform %_arr_int_uint_4
-%_ptr_StorageBuffer__arr_int_uint_4 = OpTypePointer StorageBuffer %_arr_int_uint_4
+%_ptr_Uniform__arr_v4int_uint_4 = OpTypePointer Uniform %_arr_v4int_uint_4
+%_ptr_StorageBuffer__arr_v4int_uint_4 = OpTypePointer StorageBuffer %_arr_v4int_uint_4
 %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 = OpTypePointer Function %_arr__arr__arr_int_uint_2_uint_3_uint_4
-         %59 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
-%unused_entry_point = OpFunction %void None %23
-         %26 = OpLabel
+         %63 = OpConstantNull %_arr__arr__arr_int_uint_2_uint_3_uint_4
+%unused_entry_point = OpFunction %void None %24
+         %27 = OpLabel
                OpReturn
                OpFunctionEnd
-    %ret_arr = OpFunction %_arr_int_uint_4 None %27
-         %29 = OpLabel
-               OpReturnValue %7
+    %ret_arr = OpFunction %_arr_v4int_uint_4 None %28
+         %30 = OpLabel
+               OpReturnValue %8
                OpFunctionEnd
-%ret_struct_arr = OpFunction %S None %30
-         %32 = OpLabel
-               OpReturnValue %33
+%ret_struct_arr = OpFunction %S None %31
+         %33 = OpLabel
+               OpReturnValue %34
                OpFunctionEnd
-        %foo = OpFunction %void None %34
-  %src_param = OpFunctionParameter %_arr_int_uint_4
-         %37 = OpLabel
-%src_function = OpVariable %_ptr_Function__arr_int_uint_4 Function %7
- %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %59
-               OpStore %dst %43
-               OpStore %dst %src_param
-         %44 = OpFunctionCall %_arr_int_uint_4 %ret_arr
-               OpStore %dst %44
-               OpStore %dst %7
-         %45 = OpLoad %_arr_int_uint_4 %src_function
-               OpStore %dst %45
-         %46 = OpLoad %_arr_int_uint_4 %src_private
-               OpStore %dst %46
-         %47 = OpLoad %_arr_int_uint_4 %src_workgroup
+        %foo = OpFunction %void None %35
+  %src_param = OpFunctionParameter %_arr_v4int_uint_4
+         %38 = OpLabel
+%src_function = OpVariable %_ptr_Function__arr_v4int_uint_4 Function %8
+ %src_nested = OpVariable %_ptr_Function__arr__arr__arr_int_uint_2_uint_3_uint_4 Function %63
                OpStore %dst %47
-         %48 = OpFunctionCall %S %ret_struct_arr
-         %49 = OpCompositeExtract %_arr_int_uint_4 %48 0
+               OpStore %dst %src_param
+         %48 = OpFunctionCall %_arr_v4int_uint_4 %ret_arr
+               OpStore %dst %48
+               OpStore %dst %8
+         %49 = OpLoad %_arr_v4int_uint_4 %src_function
                OpStore %dst %49
-         %52 = OpAccessChain %_ptr_Uniform__arr_int_uint_4 %src_uniform %uint_0
-         %53 = OpLoad %_arr_int_uint_4 %52
+         %50 = OpLoad %_arr_v4int_uint_4 %src_private
+               OpStore %dst %50
+         %51 = OpLoad %_arr_v4int_uint_4 %src_workgroup
+               OpStore %dst %51
+         %52 = OpFunctionCall %S %ret_struct_arr
+         %53 = OpCompositeExtract %_arr_v4int_uint_4 %52 0
                OpStore %dst %53
-         %55 = OpAccessChain %_ptr_StorageBuffer__arr_int_uint_4 %src_storage %uint_0
-         %56 = OpLoad %_arr_int_uint_4 %55
-               OpStore %dst %56
-         %60 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
-               OpStore %dst_nested %60
+         %56 = OpAccessChain %_ptr_Uniform__arr_v4int_uint_4 %src_uniform %uint_0
+         %57 = OpLoad %_arr_v4int_uint_4 %56
+               OpStore %dst %57
+         %59 = OpAccessChain %_ptr_StorageBuffer__arr_v4int_uint_4 %src_storage %uint_0
+         %60 = OpLoad %_arr_v4int_uint_4 %59
+               OpStore %dst %60
+         %64 = OpLoad %_arr__arr__arr_int_uint_2_uint_3_uint_4 %src_nested
+               OpStore %dst_nested %64
                OpReturn
                OpFunctionEnd
diff --git a/test/array/assign_to_workgroup_var.wgsl.expected.wgsl b/test/array/assign_to_workgroup_var.wgsl.expected.wgsl
index 4fecbc8..6888710 100644
--- a/test/array/assign_to_workgroup_var.wgsl.expected.wgsl
+++ b/test/array/assign_to_workgroup_var.wgsl.expected.wgsl
@@ -1,4 +1,4 @@
-type ArrayType = @stride(16) array<i32, 4>;
+type ArrayType = array<vec4<i32>, 4>;
 
 struct S {
   arr : ArrayType;
@@ -26,7 +26,7 @@
 
 fn foo(src_param : ArrayType) {
   var src_function : ArrayType;
-  dst = ArrayType(1, 2, 3, 3);
+  dst = ArrayType(vec4(1), vec4(2), vec4(3), vec4(3));
   dst = src_param;
   dst = ret_arr();
   let src_let : ArrayType = ArrayType();
diff --git a/test/buffer/storage/dynamic_index/read.wgsl b/test/buffer/storage/dynamic_index/read.wgsl
index 94c2c1f..5cd0316 100644
--- a/test/buffer/storage/dynamic_index/read.wgsl
+++ b/test/buffer/storage/dynamic_index/read.wgsl
@@ -7,7 +7,7 @@
     f : f32;
     g : mat2x3<f32>;
     h : mat3x2<f32>;
-    i : @stride(16) array<vec4<i32>, 4>;
+    i : array<vec4<i32>, 4>;
 };
 
 struct S {
diff --git a/test/buffer/storage/dynamic_index/read.wgsl.expected.wgsl b/test/buffer/storage/dynamic_index/read.wgsl.expected.wgsl
index fd0c451..f9c274c 100644
--- a/test/buffer/storage/dynamic_index/read.wgsl.expected.wgsl
+++ b/test/buffer/storage/dynamic_index/read.wgsl.expected.wgsl
@@ -7,7 +7,7 @@
   f : f32;
   g : mat2x3<f32>;
   h : mat3x2<f32>;
-  i : @stride(16) array<vec4<i32>, 4>;
+  i : array<vec4<i32>, 4>;
 }
 
 struct S {
diff --git a/test/buffer/storage/dynamic_index/write.wgsl b/test/buffer/storage/dynamic_index/write.wgsl
index 3770015..4167c1e 100644
--- a/test/buffer/storage/dynamic_index/write.wgsl
+++ b/test/buffer/storage/dynamic_index/write.wgsl
@@ -7,7 +7,7 @@
     f : f32;
     g : mat2x3<f32>;
     h : mat3x2<f32>;
-    i : @stride(16) array<vec4<i32>, 4>;
+    i : array<vec4<i32>, 4>;
 };
 
 struct S {
@@ -26,5 +26,5 @@
     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>();
+    s.arr[idx].i = array<vec4<i32>, 4>();
 }
diff --git a/test/buffer/storage/dynamic_index/write.wgsl.expected.wgsl b/test/buffer/storage/dynamic_index/write.wgsl.expected.wgsl
index d51db21..661916d 100644
--- a/test/buffer/storage/dynamic_index/write.wgsl.expected.wgsl
+++ b/test/buffer/storage/dynamic_index/write.wgsl.expected.wgsl
@@ -7,7 +7,7 @@
   f : f32;
   g : mat2x3<f32>;
   h : mat3x2<f32>;
-  i : @stride(16) array<vec4<i32>, 4>;
+  i : array<vec4<i32>, 4>;
 }
 
 struct S {
@@ -26,5 +26,5 @@
   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>();
+  s.arr[idx].i = array<vec4<i32>, 4>();
 }
diff --git a/test/buffer/storage/static_index/read.wgsl b/test/buffer/storage/static_index/read.wgsl
index 773fbbb..06ea57e 100644
--- a/test/buffer/storage/static_index/read.wgsl
+++ b/test/buffer/storage/static_index/read.wgsl
@@ -12,7 +12,7 @@
     g : mat2x3<f32>;
     h : mat3x2<f32>;
     i : Inner;
-    j : @stride(16) array<Inner, 4>;
+    j : array<Inner, 4>;
 };
 
 @binding(0) @group(0) var<storage, read> s : S;
diff --git a/test/buffer/storage/static_index/read.wgsl.expected.glsl b/test/buffer/storage/static_index/read.wgsl.expected.glsl
index f481a98..1b4c469 100644
--- a/test/buffer/storage/static_index/read.wgsl.expected.glsl
+++ b/test/buffer/storage/static_index/read.wgsl.expected.glsl
@@ -4,9 +4,6 @@
 struct Inner {
   int x;
 };
-struct tint_padded_array_element {
-  Inner el;
-};
 struct S {
   ivec3 a;
   int b;
@@ -17,7 +14,7 @@
   mat2x3 g;
   mat3x2 h;
   Inner i;
-  tint_padded_array_element j[4];
+  Inner j[4];
 };
 
 layout (binding = 0) buffer S_1 {
@@ -30,7 +27,7 @@
   mat2x3 g;
   mat3x2 h;
   Inner i;
-  tint_padded_array_element j[4];
+  Inner j[4];
 } s;
 
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@@ -44,7 +41,7 @@
   mat2x3 g = s.g;
   mat3x2 h = s.h;
   Inner i = s.i;
-  tint_padded_array_element j[4] = s.j;
+  Inner j[4] = s.j;
   return;
 }
 void main() {
diff --git a/test/buffer/storage/static_index/read.wgsl.expected.hlsl b/test/buffer/storage/static_index/read.wgsl.expected.hlsl
index 667739b..9062d00 100644
--- a/test/buffer/storage/static_index/read.wgsl.expected.hlsl
+++ b/test/buffer/storage/static_index/read.wgsl.expected.hlsl
@@ -1,9 +1,6 @@
 struct Inner {
   int x;
 };
-struct tint_padded_array_element {
-  Inner el;
-};
 
 ByteAddressBuffer s : register(t0, space0);
 
@@ -20,12 +17,12 @@
   return tint_symbol_11;
 }
 
-typedef tint_padded_array_element tint_symbol_10_ret[4];
+typedef Inner tint_symbol_10_ret[4];
 tint_symbol_10_ret tint_symbol_10(ByteAddressBuffer buffer, uint offset) {
-  tint_padded_array_element arr[4] = (tint_padded_array_element[4])0;
+  Inner arr[4] = (Inner[4])0;
   {
     [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
-      arr[i_1].el = tint_symbol_9(buffer, (offset + (i_1 * 16u)));
+      arr[i_1] = tint_symbol_9(buffer, (offset + (i_1 * 4u)));
     }
   }
   return arr;
@@ -42,6 +39,6 @@
   const float2x3 g = tint_symbol_6(s, 48u);
   const float3x2 h = tint_symbol_7(s, 80u);
   const Inner i = tint_symbol_9(s, 104u);
-  const tint_padded_array_element j[4] = tint_symbol_10(s, 108u);
+  const Inner 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
index d60818e..efc8a38 100644
--- a/test/buffer/storage/static_index/read.wgsl.expected.msl
+++ b/test/buffer/storage/static_index/read.wgsl.expected.msl
@@ -15,12 +15,8 @@
 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];
+  /* 0x0000 */ Inner arr[4];
 };
 struct S {
   /* 0x0000 */ packed_int3 a;
@@ -33,7 +29,7 @@
   /* 0x0050 */ float3x2 h;
   /* 0x0068 */ Inner i;
   /* 0x006c */ tint_array_wrapper j;
-  /* 0x00ac */ int8_t tint_pad_1[4];
+  /* 0x007c */ int8_t tint_pad[4];
 };
 
 kernel void tint_symbol(const device S* tint_symbol_1 [[buffer(0)]]) {
diff --git a/test/buffer/storage/static_index/read.wgsl.expected.spvasm b/test/buffer/storage/static_index/read.wgsl.expected.spvasm
index 98bf962..42262ef 100644
--- a/test/buffer/storage/static_index/read.wgsl.expected.spvasm
+++ b/test/buffer/storage/static_index/read.wgsl.expected.spvasm
@@ -38,7 +38,7 @@
                OpMemberDecorate %S 8 Offset 104
                OpMemberDecorate %Inner 0 Offset 0
                OpMemberDecorate %S 9 Offset 108
-               OpDecorate %_arr_Inner_uint_4 ArrayStride 16
+               OpDecorate %_arr_Inner_uint_4 ArrayStride 4
                OpDecorate %s NonWritable
                OpDecorate %s Binding 0
                OpDecorate %s DescriptorSet 0
diff --git a/test/buffer/storage/static_index/read.wgsl.expected.wgsl b/test/buffer/storage/static_index/read.wgsl.expected.wgsl
index 3413a96..6c48125 100644
--- a/test/buffer/storage/static_index/read.wgsl.expected.wgsl
+++ b/test/buffer/storage/static_index/read.wgsl.expected.wgsl
@@ -12,7 +12,7 @@
   g : mat2x3<f32>;
   h : mat3x2<f32>;
   i : Inner;
-  j : @stride(16) array<Inner, 4>;
+  j : array<Inner, 4>;
 }
 
 @binding(0) @group(0) var<storage, read> s : S;
diff --git a/test/buffer/storage/static_index/write.wgsl b/test/buffer/storage/static_index/write.wgsl
index ff18cbc..47827bf 100644
--- a/test/buffer/storage/static_index/write.wgsl
+++ b/test/buffer/storage/static_index/write.wgsl
@@ -12,7 +12,7 @@
     g : mat2x3<f32>;
     h : mat3x2<f32>;
     i : Inner;
-    j : @stride(16) array<Inner, 4>;
+    j : array<Inner, 4>;
 };
 
 @binding(0) @group(0) var<storage, write> s : S;
@@ -28,5 +28,5 @@
     s.g = mat2x3<f32>();
     s.h = mat3x2<f32>();
     s.i = Inner();
-    s.j = @stride(16) array<Inner, 4>();
+    s.j = array<Inner, 4>();
 }
diff --git a/test/buffer/storage/static_index/write.wgsl.expected.glsl b/test/buffer/storage/static_index/write.wgsl.expected.glsl
index 4434c96..b4db83b 100644
--- a/test/buffer/storage/static_index/write.wgsl.expected.glsl
+++ b/test/buffer/storage/static_index/write.wgsl.expected.glsl
@@ -4,9 +4,6 @@
 struct Inner {
   int x;
 };
-struct tint_padded_array_element {
-  Inner el;
-};
 struct S {
   ivec3 a;
   int b;
@@ -17,7 +14,7 @@
   mat2x3 g;
   mat3x2 h;
   Inner i;
-  tint_padded_array_element j[4];
+  Inner j[4];
 };
 
 layout (binding = 0) buffer S_1 {
@@ -30,7 +27,7 @@
   mat2x3 g;
   mat3x2 h;
   Inner i;
-  tint_padded_array_element j[4];
+  Inner j[4];
 } s;
 
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@@ -45,7 +42,7 @@
   s.h = mat3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
   Inner tint_symbol_1 = Inner(0);
   s.i = tint_symbol_1;
-  tint_padded_array_element tint_symbol_2[4] = tint_padded_array_element[4](tint_padded_array_element(Inner(0)), tint_padded_array_element(Inner(0)), tint_padded_array_element(Inner(0)), tint_padded_array_element(Inner(0)));
+  Inner tint_symbol_2[4] = Inner[4](Inner(0), Inner(0), Inner(0), Inner(0));
   s.j = tint_symbol_2;
   return;
 }
diff --git a/test/buffer/storage/static_index/write.wgsl.expected.hlsl b/test/buffer/storage/static_index/write.wgsl.expected.hlsl
index 8a2e683..f4d1745 100644
--- a/test/buffer/storage/static_index/write.wgsl.expected.hlsl
+++ b/test/buffer/storage/static_index/write.wgsl.expected.hlsl
@@ -1,9 +1,6 @@
 struct Inner {
   int x;
 };
-struct tint_padded_array_element {
-  Inner el;
-};
 
 RWByteAddressBuffer s : register(u0, space0);
 
@@ -22,11 +19,11 @@
   buffer.Store((offset + 0u), asuint(value.x));
 }
 
-void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[4]) {
-  tint_padded_array_element array[4] = value;
+void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, Inner value[4]) {
+  Inner array[4] = value;
   {
     [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
-      tint_symbol_9(buffer, (offset + (i_1 * 16u)), array[i_1].el);
+      tint_symbol_9(buffer, (offset + (i_1 * 4u)), array[i_1]);
     }
   }
 }
@@ -43,7 +40,7 @@
   tint_symbol_7(s, 80u, float3x2(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f));
   const Inner tint_symbol_11 = (Inner)0;
   tint_symbol_9(s, 104u, tint_symbol_11);
-  const tint_padded_array_element tint_symbol_12[4] = (tint_padded_array_element[4])0;
+  const Inner tint_symbol_12[4] = (Inner[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
index bae3fcd..a339682 100644
--- a/test/buffer/storage/static_index/write.wgsl.expected.msl
+++ b/test/buffer/storage/static_index/write.wgsl.expected.msl
@@ -15,12 +15,8 @@
 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];
+  /* 0x0000 */ Inner arr[4];
 };
 struct S {
   /* 0x0000 */ packed_int3 a;
@@ -33,7 +29,7 @@
   /* 0x0050 */ float3x2 h;
   /* 0x0068 */ Inner i;
   /* 0x006c */ tint_array_wrapper j;
-  /* 0x00ac */ int8_t tint_pad_1[4];
+  /* 0x007c */ int8_t tint_pad[4];
 };
 
 kernel void tint_symbol(device S* tint_symbol_3 [[buffer(0)]]) {
diff --git a/test/buffer/storage/static_index/write.wgsl.expected.spvasm b/test/buffer/storage/static_index/write.wgsl.expected.spvasm
index faa44c5..c099dbd 100644
--- a/test/buffer/storage/static_index/write.wgsl.expected.spvasm
+++ b/test/buffer/storage/static_index/write.wgsl.expected.spvasm
@@ -38,7 +38,7 @@
                OpMemberDecorate %S 8 Offset 104
                OpMemberDecorate %Inner 0 Offset 0
                OpMemberDecorate %S 9 Offset 108
-               OpDecorate %_arr_Inner_uint_4 ArrayStride 16
+               OpDecorate %_arr_Inner_uint_4 ArrayStride 4
                OpDecorate %s NonReadable
                OpDecorate %s Binding 0
                OpDecorate %s DescriptorSet 0
diff --git a/test/buffer/storage/static_index/write.wgsl.expected.wgsl b/test/buffer/storage/static_index/write.wgsl.expected.wgsl
index 595efce..68a73e0 100644
--- a/test/buffer/storage/static_index/write.wgsl.expected.wgsl
+++ b/test/buffer/storage/static_index/write.wgsl.expected.wgsl
@@ -12,7 +12,7 @@
   g : mat2x3<f32>;
   h : mat3x2<f32>;
   i : Inner;
-  j : @stride(16) array<Inner, 4>;
+  j : array<Inner, 4>;
 }
 
 @binding(0) @group(0) var<storage, write> s : S;
@@ -28,5 +28,5 @@
   s.g = mat2x3<f32>();
   s.h = mat3x2<f32>();
   s.i = Inner();
-  s.j = @stride(16) array<Inner, 4>();
+  s.j = array<Inner, 4>();
 }
diff --git a/test/buffer/uniform/dynamic_index/read.wgsl b/test/buffer/uniform/dynamic_index/read.wgsl
index 3e9142b..d5b7564 100644
--- a/test/buffer/uniform/dynamic_index/read.wgsl
+++ b/test/buffer/uniform/dynamic_index/read.wgsl
@@ -9,7 +9,7 @@
     h : vec2<i32>;
     i : mat2x3<f32>;
     @align(16) j : mat3x2<f32>;
-    @align(16) k : @stride(16) array<vec4<i32>, 4>;
+    @align(16) k : array<vec4<i32>, 4>;
 };
 
 struct S {
diff --git a/test/buffer/uniform/dynamic_index/read.wgsl.expected.wgsl b/test/buffer/uniform/dynamic_index/read.wgsl.expected.wgsl
index 1bee626..dbf4c9e 100644
--- a/test/buffer/uniform/dynamic_index/read.wgsl.expected.wgsl
+++ b/test/buffer/uniform/dynamic_index/read.wgsl.expected.wgsl
@@ -11,7 +11,7 @@
   @align(16)
   j : mat3x2<f32>;
   @align(16)
-  k : @stride(16) array<vec4<i32>, 4>;
+  k : array<vec4<i32>, 4>;
 }
 
 struct S {
diff --git a/test/buffer/uniform/static_index/read.wgsl b/test/buffer/uniform/static_index/read.wgsl
index e1d10e8..6a27ee9 100644
--- a/test/buffer/uniform/static_index/read.wgsl
+++ b/test/buffer/uniform/static_index/read.wgsl
@@ -1,5 +1,5 @@
 struct Inner {
-    x : i32;
+    @size(16) x : i32;
 };
 
 struct S {
@@ -14,7 +14,7 @@
     i : mat2x3<f32>;
     j : mat3x2<f32>;
     @align(16) k : Inner;
-    @align(16) l : @stride(16) array<Inner, 4>;
+    @align(16) l : array<Inner, 4>;
 };
 
 @binding(0) @group(0) var<uniform> s : S;
diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.glsl b/test/buffer/uniform/static_index/read.wgsl.expected.glsl
index c7b4979..ee1c6d8 100644
--- a/test/buffer/uniform/static_index/read.wgsl.expected.glsl
+++ b/test/buffer/uniform/static_index/read.wgsl.expected.glsl
@@ -4,9 +4,6 @@
 struct Inner {
   int x;
 };
-struct tint_padded_array_element {
-  Inner el;
-};
 struct S {
   ivec3 a;
   int b;
@@ -19,7 +16,7 @@
   mat2x3 i;
   mat3x2 j;
   Inner k;
-  tint_padded_array_element l[4];
+  Inner l[4];
 };
 
 layout (binding = 0) uniform S_1 {
@@ -34,7 +31,7 @@
   mat2x3 i;
   mat3x2 j;
   Inner k;
-  tint_padded_array_element l[4];
+  Inner l[4];
 } s;
 
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
@@ -50,7 +47,7 @@
   mat2x3 i = s.i;
   mat3x2 j = s.j;
   Inner k = s.k;
-  tint_padded_array_element l[4] = s.l;
+  Inner l[4] = s.l;
   return;
 }
 void main() {
diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.hlsl b/test/buffer/uniform/static_index/read.wgsl.expected.hlsl
index 9fa1d03..1446321 100644
--- a/test/buffer/uniform/static_index/read.wgsl.expected.hlsl
+++ b/test/buffer/uniform/static_index/read.wgsl.expected.hlsl
@@ -1,9 +1,6 @@
 struct Inner {
   int x;
 };
-struct tint_padded_array_element {
-  Inner el;
-};
 
 cbuffer cbuffer_s : register(b0, space0) {
   uint4 s[13];
@@ -31,12 +28,12 @@
   return tint_symbol_12;
 }
 
-typedef tint_padded_array_element tint_symbol_11_ret[4];
+typedef Inner tint_symbol_11_ret[4];
 tint_symbol_11_ret tint_symbol_11(uint4 buffer[13], uint offset) {
-  tint_padded_array_element arr[4] = (tint_padded_array_element[4])0;
+  Inner arr[4] = (Inner[4])0;
   {
     [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
-      arr[i_1].el = tint_symbol_10(buffer, (offset + (i_1 * 16u)));
+      arr[i_1] = tint_symbol_10(buffer, (offset + (i_1 * 16u)));
     }
   }
   return arr;
@@ -55,6 +52,6 @@
   const float2x3 i = tint_symbol_7(s, 64u);
   const float3x2 j = tint_symbol_8(s, 96u);
   const Inner k = tint_symbol_10(s, 128u);
-  const tint_padded_array_element l[4] = tint_symbol_11(s, 144u);
+  const Inner 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
index 670ba5d..46934da 100644
--- a/test/buffer/uniform/static_index/read.wgsl.expected.msl
+++ b/test/buffer/uniform/static_index/read.wgsl.expected.msl
@@ -14,13 +14,10 @@
 
 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];
+  /* 0x0000 */ Inner arr[4];
 };
 struct S {
   /* 0x0000 */ packed_int3 a;
@@ -35,7 +32,6 @@
   /* 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;
 };
 
diff --git a/test/buffer/uniform/static_index/read.wgsl.expected.wgsl b/test/buffer/uniform/static_index/read.wgsl.expected.wgsl
index b790c19..e23dc3b 100644
--- a/test/buffer/uniform/static_index/read.wgsl.expected.wgsl
+++ b/test/buffer/uniform/static_index/read.wgsl.expected.wgsl
@@ -1,4 +1,5 @@
 struct Inner {
+  @size(16)
   x : i32;
 }
 
@@ -16,7 +17,7 @@
   @align(16)
   k : Inner;
   @align(16)
-  l : @stride(16) array<Inner, 4>;
+  l : array<Inner, 4>;
 }
 
 @binding(0) @group(0) var<uniform> s : S;
diff --git a/test/bug/chromium/1273230.wgsl b/test/bug/chromium/1273230.wgsl
index 90ed840..dad0cba 100644
--- a/test/bug/chromium/1273230.wgsl
+++ b/test/bug/chromium/1273230.wgsl
@@ -37,23 +37,23 @@
 };
 
 struct F32s {
-  values : @stride(4) array<f32>;
+  values : array<f32>;
 };
 
 struct U32s {
-  values : @stride(4) array<u32>;
+  values : array<u32>;
 };
 
 struct I32s {
-  values : @stride(4) array<i32>;
+  values : array<i32>;
 };
 
 struct AU32s {
-  values : @stride(4) array<atomic<u32>>;
+  values : array<atomic<u32>>;
 };
 
 struct AI32s {
-  values : @stride(4) array<atomic<i32>>;
+  values : array<atomic<i32>>;
 };
 
 @binding(0) @group(0) var<uniform> uniforms : Uniforms;
diff --git a/test/bug/chromium/1273230.wgsl.expected.wgsl b/test/bug/chromium/1273230.wgsl.expected.wgsl
index 3cdfc2b..2d9a271 100644
--- a/test/bug/chromium/1273230.wgsl.expected.wgsl
+++ b/test/bug/chromium/1273230.wgsl.expected.wgsl
@@ -54,23 +54,23 @@
 }
 
 struct F32s {
-  values : @stride(4) array<f32>;
+  values : array<f32>;
 }
 
 struct U32s {
-  values : @stride(4) array<u32>;
+  values : array<u32>;
 }
 
 struct I32s {
-  values : @stride(4) array<i32>;
+  values : array<i32>;
 }
 
 struct AU32s {
-  values : @stride(4) array<atomic<u32>>;
+  values : array<atomic<u32>>;
 }
 
 struct AI32s {
-  values : @stride(4) array<atomic<i32>>;
+  values : array<atomic<i32>>;
 }
 
 @binding(0) @group(0) var<uniform> uniforms : Uniforms;
diff --git a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl
index be26146..d05dd5b 100644
--- a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl
+++ b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl
@@ -1,5 +1,5 @@
 struct UBO {
-  data: @stride(16) array<i32, 4>;
+  data: array<vec4<i32>, 4>;
   dynamic_idx: i32;
 };
 @group(0) @binding(0) var<uniform> ubo: UBO;
@@ -10,5 +10,5 @@
 
 @stage(compute) @workgroup_size(1)
 fn f() {
-  result.out = ubo.data[ubo.dynamic_idx];
+  result.out = ubo.data[ubo.dynamic_idx].x;
 }
diff --git a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.glsl b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.glsl
index cb10e51..df1ea99 100644
--- a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.glsl
+++ b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.glsl
@@ -1,16 +1,13 @@
 #version 310 es
 precision mediump float;
 
-struct tint_padded_array_element {
-  int el;
-};
 struct UBO {
-  tint_padded_array_element data[4];
+  ivec4 data[4];
   int dynamic_idx;
 };
 
 layout (binding = 0) uniform UBO_1 {
-  tint_padded_array_element data[4];
+  ivec4 data[4];
   int dynamic_idx;
 } ubo;
 
@@ -24,7 +21,7 @@
 
 layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 void f() {
-  result.tint_symbol = ubo.data[ubo.dynamic_idx].el;
+  result.tint_symbol = ubo.data[ubo.dynamic_idx].x;
   return;
 }
 void main() {
diff --git a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.msl b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.msl
index 3c75f01..3721713 100644
--- a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.msl
+++ b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.msl
@@ -1,23 +1,20 @@
 #include <metal_stdlib>
 
 using namespace metal;
-struct tint_padded_array_element {
-  /* 0x0000 */ int el;
-  /* 0x0004 */ int8_t tint_pad[12];
-};
 struct tint_array_wrapper {
-  /* 0x0000 */ tint_padded_array_element arr[4];
+  /* 0x0000 */ int4 arr[4];
 };
 struct UBO {
   /* 0x0000 */ tint_array_wrapper data;
   /* 0x0040 */ int dynamic_idx;
+  /* 0x0044 */ int8_t tint_pad[12];
 };
 struct Result {
   /* 0x0000 */ int out;
 };
 
 kernel void f(device Result* tint_symbol [[buffer(1)]], const constant UBO* tint_symbol_1 [[buffer(0)]]) {
-  (*(tint_symbol)).out = (*(tint_symbol_1)).data.arr[(*(tint_symbol_1)).dynamic_idx].el;
+  (*(tint_symbol)).out = (*(tint_symbol_1)).data.arr[(*(tint_symbol_1)).dynamic_idx][0];
   return;
 }
 
diff --git a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.spvasm b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.spvasm
index 1026040..8e837ac 100644
--- a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.spvasm
+++ b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.spvasm
@@ -1,7 +1,7 @@
 ; SPIR-V
 ; Version: 1.3
 ; Generator: Google Tint Compiler; 0
-; Bound: 24
+; Bound: 25
 ; Schema: 0
                OpCapability Shader
                OpMemoryModel Logical GLSL450
@@ -17,7 +17,7 @@
                OpName %f "f"
                OpDecorate %UBO Block
                OpMemberDecorate %UBO 0 Offset 0
-               OpDecorate %_arr_int_uint_4 ArrayStride 16
+               OpDecorate %_arr_v4int_uint_4 ArrayStride 16
                OpMemberDecorate %UBO 1 Offset 64
                OpDecorate %ubo NonWritable
                OpDecorate %ubo DescriptorSet 0
@@ -27,28 +27,29 @@
                OpDecorate %result DescriptorSet 0
                OpDecorate %result Binding 2
         %int = OpTypeInt 32 1
+      %v4int = OpTypeVector %int 4
        %uint = OpTypeInt 32 0
      %uint_4 = OpConstant %uint 4
-%_arr_int_uint_4 = OpTypeArray %int %uint_4
-        %UBO = OpTypeStruct %_arr_int_uint_4 %int
+%_arr_v4int_uint_4 = OpTypeArray %v4int %uint_4
+        %UBO = OpTypeStruct %_arr_v4int_uint_4 %int
 %_ptr_Uniform_UBO = OpTypePointer Uniform %UBO
         %ubo = OpVariable %_ptr_Uniform_UBO Uniform
      %Result = OpTypeStruct %int
 %_ptr_StorageBuffer_Result = OpTypePointer StorageBuffer %Result
      %result = OpVariable %_ptr_StorageBuffer_Result StorageBuffer
        %void = OpTypeVoid
-         %11 = OpTypeFunction %void
+         %12 = OpTypeFunction %void
      %uint_0 = OpConstant %uint 0
 %_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
      %uint_1 = OpConstant %uint 1
 %_ptr_Uniform_int = OpTypePointer Uniform %int
-          %f = OpFunction %void None %11
-         %14 = OpLabel
-         %17 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0
-         %20 = OpAccessChain %_ptr_Uniform_int %ubo %uint_1
-         %21 = OpLoad %int %20
-         %22 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0 %21
-         %23 = OpLoad %int %22
-               OpStore %17 %23
+          %f = OpFunction %void None %12
+         %15 = OpLabel
+         %18 = OpAccessChain %_ptr_StorageBuffer_int %result %uint_0
+         %21 = OpAccessChain %_ptr_Uniform_int %ubo %uint_1
+         %22 = OpLoad %int %21
+         %23 = OpAccessChain %_ptr_Uniform_int %ubo %uint_0 %22 %uint_0
+         %24 = OpLoad %int %23
+               OpStore %18 %24
                OpReturn
                OpFunctionEnd
diff --git a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.wgsl b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.wgsl
index 7f3ca0b..0dbf26b 100644
--- a/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.wgsl
+++ b/test/bug/fxc/dyn_array_idx/read/uniform.wgsl.expected.wgsl
@@ -1,5 +1,5 @@
 struct UBO {
-  data : @stride(16) array<i32, 4>;
+  data : array<vec4<i32>, 4>;
   dynamic_idx : i32;
 }
 
@@ -13,5 +13,5 @@
 
 @stage(compute) @workgroup_size(1)
 fn f() {
-  result.out = ubo.data[ubo.dynamic_idx];
+  result.out = ubo.data[ubo.dynamic_idx].x;
 }
diff --git a/test/bug/tint/1046.wgsl b/test/bug/tint/1046.wgsl
index a294edf..40eeee6 100644
--- a/test/bug/tint/1046.wgsl
+++ b/test/bug/tint/1046.wgsl
@@ -4,7 +4,7 @@
 };
 
  struct PointLights {
-    values : @stride(16) array<PointLight>;
+    values : array<PointLight>;
 };
 
  struct Uniforms {
diff --git a/test/bug/tint/1046.wgsl.expected.wgsl b/test/bug/tint/1046.wgsl.expected.wgsl
index b7000e6..f7ef8d8 100644
--- a/test/bug/tint/1046.wgsl.expected.wgsl
+++ b/test/bug/tint/1046.wgsl.expected.wgsl
@@ -3,7 +3,7 @@
 }
 
 struct PointLights {
-  values : @stride(16) array<PointLight>;
+  values : array<PointLight>;
 }
 
 struct Uniforms {
diff --git a/test/bug/tint/1113.wgsl b/test/bug/tint/1113.wgsl
index f3a080a..87912ff 100644
--- a/test/bug/tint/1113.wgsl
+++ b/test/bug/tint/1113.wgsl
@@ -23,11 +23,11 @@
     value_f32_3 : f32;
 };
 
- struct F32s { values : @stride(4) array<f32>; };
- struct U32s { values : @stride(4) array<u32>; };
- struct I32s { values : @stride(4) array<i32>; };
- struct AU32s { values : @stride(4) array<atomic<u32>>; };
- struct AI32s { values : @stride(4) array<atomic<i32>>; };
+ struct F32s { values : array<f32>; };
+ struct U32s { values : array<u32>; };
+ struct I32s { values : array<i32>; };
+ struct AU32s { values : array<atomic<u32>>; };
+ struct AI32s { values : array<atomic<i32>>; };
 
 // IN
 @binding(0) @group(0) var<uniform> uniforms : Uniforms;
diff --git a/test/bug/tint/1113.wgsl.expected.wgsl b/test/bug/tint/1113.wgsl.expected.wgsl
index 9eeb16c..c709577 100644
--- a/test/bug/tint/1113.wgsl.expected.wgsl
+++ b/test/bug/tint/1113.wgsl.expected.wgsl
@@ -23,23 +23,23 @@
 }
 
 struct F32s {
-  values : @stride(4) array<f32>;
+  values : array<f32>;
 }
 
 struct U32s {
-  values : @stride(4) array<u32>;
+  values : array<u32>;
 }
 
 struct I32s {
-  values : @stride(4) array<i32>;
+  values : array<i32>;
 }
 
 struct AU32s {
-  values : @stride(4) array<atomic<u32>>;
+  values : array<atomic<u32>>;
 }
 
 struct AI32s {
-  values : @stride(4) array<atomic<i32>>;
+  values : array<atomic<i32>>;
 }
 
 @binding(0) @group(0) var<uniform> uniforms : Uniforms;
diff --git a/test/bug/tint/294.wgsl b/test/bug/tint/294.wgsl
index 4917c47..beff423 100644
--- a/test/bug/tint/294.wgsl
+++ b/test/bug/tint/294.wgsl
@@ -3,6 +3,6 @@
   colour : vec3<f32>;
 };
  struct Lights {
-  light : @stride(32) array<Light>;
+  light : array<Light>;
 };
 @group(0) @binding(1) var<storage, read> lights : Lights;
diff --git a/test/bug/tint/294.wgsl.expected.wgsl b/test/bug/tint/294.wgsl.expected.wgsl
index 87f08de..f73ce3c 100644
--- a/test/bug/tint/294.wgsl.expected.wgsl
+++ b/test/bug/tint/294.wgsl.expected.wgsl
@@ -4,7 +4,7 @@
 }
 
 struct Lights {
-  light : @stride(32) array<Light>;
+  light : array<Light>;
 }
 
 @group(0) @binding(1) var<storage, read> lights : Lights;
diff --git a/test/bug/tint/534.wgsl b/test/bug/tint/534.wgsl
index 782f193..77e9f6b 100644
--- a/test/bug/tint/534.wgsl
+++ b/test/bug/tint/534.wgsl
@@ -5,7 +5,7 @@
     channelCount : u32;
 };
  struct OutputBuf {
-    result : @stride(4) array<u32>;
+    result : array<u32>;
 };
 @group(0) @binding(0) var src : texture_2d<f32>;
 @group(0) @binding(1) var dst : texture_2d<f32>;
diff --git a/test/bug/tint/534.wgsl.expected.wgsl b/test/bug/tint/534.wgsl.expected.wgsl
index dfd0f17..7e86208 100644
--- a/test/bug/tint/534.wgsl.expected.wgsl
+++ b/test/bug/tint/534.wgsl.expected.wgsl
@@ -6,7 +6,7 @@
 }
 
 struct OutputBuf {
-  result : @stride(4) array<u32>;
+  result : array<u32>;
 }
 
 @group(0) @binding(0) var src : texture_2d<f32>;
diff --git a/test/bug/tint/757.wgsl b/test/bug/tint/757.wgsl
index b2a05de..24b0418 100644
--- a/test/bug/tint/757.wgsl
+++ b/test/bug/tint/757.wgsl
@@ -7,7 +7,7 @@
 @group(0) @binding(1) var myTexture : texture_2d_array<f32>;
 
  struct Result {
-  values : @stride(4) array<f32>;
+  values : array<f32>;
 };
 @group(0) @binding(3) var<storage, read_write> result : Result;
 
diff --git a/test/bug/tint/757.wgsl.expected.wgsl b/test/bug/tint/757.wgsl.expected.wgsl
index b237cd5..6beca1d 100644
--- a/test/bug/tint/757.wgsl.expected.wgsl
+++ b/test/bug/tint/757.wgsl.expected.wgsl
@@ -7,7 +7,7 @@
 @group(0) @binding(1) var myTexture : texture_2d_array<f32>;
 
 struct Result {
-  values : @stride(4) array<f32>;
+  values : array<f32>;
 }
 
 @group(0) @binding(3) var<storage, read_write> result : Result;
diff --git a/test/bug/tint/782.wgsl b/test/bug/tint/782.wgsl
index 3fe1f23..0a4f31b 100644
--- a/test/bug/tint/782.wgsl
+++ b/test/bug/tint/782.wgsl
@@ -1,4 +1,4 @@
-type ArrayExplicitStride = @stride(4) array<i32, 2>;
+type ArrayExplicitStride = array<i32, 2>;
 type ArrayImplicitStride =               array<i32, 2>;
 
 fn foo() {
diff --git a/test/bug/tint/782.wgsl.expected.wgsl b/test/bug/tint/782.wgsl.expected.wgsl
index a4f37bc..9e89537 100644
--- a/test/bug/tint/782.wgsl.expected.wgsl
+++ b/test/bug/tint/782.wgsl.expected.wgsl
@@ -1,4 +1,4 @@
-type ArrayExplicitStride = @stride(4) array<i32, 2>;
+type ArrayExplicitStride = array<i32, 2>;
 
 type ArrayImplicitStride = array<i32, 2>;
 
diff --git a/test/bug/tint/922.wgsl b/test/bug/tint/922.wgsl
index 450aba1..1ef80df 100644
--- a/test/bug/tint/922.wgsl
+++ b/test/bug/tint/922.wgsl
@@ -21,12 +21,12 @@
 };
 
 struct ub_MaterialParams {
-    u_TexMtx: @stride(32) array<Mat4x2_,1>;
+    u_TexMtx: array<Mat4x2_,1>;
     u_Misc0_: vec4<f32>;
 };
 
 struct ub_PacketParams {
-    u_PosMtx: @stride(48) array<Mat4x3_,32>;
+    u_PosMtx: array<Mat4x3_,32>;
 };
 
 struct VertexOutput {
diff --git a/test/bug/tint/922.wgsl.expected.wgsl b/test/bug/tint/922.wgsl.expected.wgsl
index ac8b48c..8151b6e 100644
--- a/test/bug/tint/922.wgsl.expected.wgsl
+++ b/test/bug/tint/922.wgsl.expected.wgsl
@@ -21,12 +21,12 @@
 }
 
 struct ub_MaterialParams {
-  u_TexMtx : @stride(32) array<Mat4x2_, 1>;
+  u_TexMtx : array<Mat4x2_, 1>;
   u_Misc0_ : vec4<f32>;
 }
 
 struct ub_PacketParams {
-  u_PosMtx : @stride(48) array<Mat4x3_, 32>;
+  u_PosMtx : array<Mat4x3_, 32>;
 }
 
 struct VertexOutput {