diff --git a/docs/tint/origin-trial-changes.md b/docs/tint/origin-trial-changes.md
index 2371a83..ff3c4f6 100644
--- a/docs/tint/origin-trial-changes.md
+++ b/docs/tint/origin-trial-changes.md
@@ -5,6 +5,7 @@
 ### Breaking changes
 
 * The `@block` attribute has been removed. [tint:1324](crbug.com/tint/1324)
+* The `@stride` attribute has been removed. [tint:1381](crbug.com/tint/1381)
 * Attributes using `[[attribute]]` syntax are no longer supported. [tint:1382](crbug.com/tint/1382)
 
 ## Changes for M101
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 54578e7..91bee00 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -467,8 +467,6 @@
     "transform/multiplanar_external_texture.h",
     "transform/num_workgroups_from_uniform.cc",
     "transform/num_workgroups_from_uniform.h",
-    "transform/pad_array_elements.cc",
-    "transform/pad_array_elements.h",
     "transform/promote_initializers_to_const_var.cc",
     "transform/promote_initializers_to_const_var.h",
     "transform/promote_side_effects_to_decl.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 011fc39..5e077e6 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -345,8 +345,6 @@
   transform/multiplanar_external_texture.h
   transform/num_workgroups_from_uniform.cc
   transform/num_workgroups_from_uniform.h
-  transform/pad_array_elements.cc
-  transform/pad_array_elements.h
   transform/promote_initializers_to_const_var.cc
   transform/promote_initializers_to_const_var.h
   transform/promote_side_effects_to_decl.cc
@@ -1023,7 +1021,6 @@
       transform/module_scope_var_to_entry_point_param_test.cc
       transform/multiplanar_external_texture_test.cc
       transform/num_workgroups_from_uniform_test.cc
-      transform/pad_array_elements_test.cc
       transform/promote_initializers_to_const_var_test.cc
       transform/promote_side_effects_to_decl_test.cc
       transform/remove_phonies_test.cc
diff --git a/src/tint/ast/module_clone_test.cc b/src/tint/ast/module_clone_test.cc
index 12853ba..983d1b0 100644
--- a/src/tint/ast/module_clone_test.cc
+++ b/src/tint/ast/module_clone_test.cc
@@ -41,7 +41,7 @@
 let c0 : i32 = 10;
 let c1 : bool = true;
 
-type t0 = @stride(16) array<vec4<f32>>;
+type t0 = array<vec4<f32>>;
 type t1 = array<vec4<f32>>;
 
 var<private> g0 : u32 = 20u;
diff --git a/src/tint/ast/stride_attribute.h b/src/tint/ast/stride_attribute.h
index 4edc54c..0aa3baf 100644
--- a/src/tint/ast/stride_attribute.h
+++ b/src/tint/ast/stride_attribute.h
@@ -18,11 +18,13 @@
 #include <string>
 
 #include "src/tint/ast/attribute.h"
+#include "src/tint/ast/internal_attribute.h"
 
 namespace tint {
 namespace ast {
 
-/// A stride attribute
+/// A stride attribute used by the SPIR-V reader for strided arrays and
+/// matrices.
 class StrideAttribute final : public Castable<StrideAttribute, Attribute> {
  public:
   /// constructor
diff --git a/src/tint/reader/wgsl/parser_impl.cc b/src/tint/reader/wgsl/parser_impl.cc
index eaf5fb7..0f62a2f 100644
--- a/src/tint/reader/wgsl/parser_impl.cc
+++ b/src/tint/reader/wgsl/parser_impl.cc
@@ -119,7 +119,6 @@
 const char kSizeAttribute[] = "size";
 const char kAlignAttribute[] = "align";
 const char kStageAttribute[] = "stage";
-const char kStrideAttribute[] = "stride";
 const char kWorkgroupSizeAttribute[] = "workgroup_size";
 
 // https://gpuweb.github.io/gpuweb/wgsl.html#reserved-keywords
@@ -799,7 +798,7 @@
 }
 
 // variable_ident_decl
-//   : IDENT COLON variable_attribute_list* type_decl
+//   : IDENT COLON type_decl
 Expect<ParserImpl::TypedIdentifier> ParserImpl::expect_variable_ident_decl(
     std::string_view use,
     bool allow_inferred) {
@@ -814,20 +813,13 @@
   if (!expect(use, Token::Type::kColon))
     return Failure::kErrored;
 
-  auto attrs = attribute_list();
-  if (attrs.errored)
-    return Failure::kErrored;
-
   auto t = peek();
-  auto type = type_decl(attrs.value);
+  auto type = type_decl();
   if (type.errored)
     return Failure::kErrored;
   if (!type.matched)
     return add_error(t.source(), "invalid type", use);
 
-  if (!expect_attributes_consumed(attrs.value))
-    return Failure::kErrored;
-
   return TypedIdentifier{type.value, ident.value, ident.source};
 }
 
@@ -929,25 +921,6 @@
 //   | MAT4x4 LESS_THAN type_decl GREATER_THAN
 //   | texture_sampler_types
 Maybe<const ast::Type*> ParserImpl::type_decl() {
-  auto attrs = attribute_list();
-  if (attrs.errored)
-    return Failure::kErrored;
-
-  auto type = type_decl(attrs.value);
-  if (type.errored) {
-    return Failure::kErrored;
-  }
-  if (!expect_attributes_consumed(attrs.value)) {
-    return Failure::kErrored;
-  }
-  if (!type.matched) {
-    return Failure::kNoMatch;
-  }
-
-  return type;
-}
-
-Maybe<const ast::Type*> ParserImpl::type_decl(ast::AttributeList& attrs) {
   auto t = peek();
   Source source;
   if (match(Token::Type::kIdentifier, &source)) {
@@ -981,7 +954,7 @@
   }
 
   if (match(Token::Type::kArray, &source)) {
-    return expect_type_decl_array(t, std::move(attrs));
+    return expect_type_decl_array(t);
   }
 
   if (t.IsMatrix()) {
@@ -1080,9 +1053,7 @@
   return builder_.ty.vec(make_source_range_from(t.source()), subtype, count);
 }
 
-Expect<const ast::Type*> ParserImpl::expect_type_decl_array(
-    Token t,
-    ast::AttributeList attrs) {
+Expect<const ast::Type*> ParserImpl::expect_type_decl_array(Token t) {
   const char* use = "array declaration";
 
   const ast::Expression* size = nullptr;
@@ -1111,7 +1082,7 @@
   }
 
   return builder_.ty.array(make_source_range_from(t.source()), subtype.value,
-                           size, std::move(attrs));
+                           size);
 }
 
 Expect<const ast::Type*> ParserImpl::expect_type_decl_matrix(Token t) {
@@ -1321,19 +1292,7 @@
     }
     return_attributes = attrs.value;
 
-    // Apply stride attributes to the type node instead of the function.
-    ast::AttributeList type_attributes;
-    auto itr =
-        std::find_if(return_attributes.begin(), return_attributes.end(),
-                     [](auto* attr) { return Is<ast::StrideAttribute>(attr); });
-    if (itr != return_attributes.end()) {
-      type_attributes.emplace_back(*itr);
-      return_attributes.erase(itr);
-    }
-
-    auto tok = peek();
-
-    auto type = type_decl(type_attributes);
+    auto type = type_decl();
     if (type.errored) {
       errored = true;
     } else if (!type.matched) {
@@ -2974,19 +2933,6 @@
     });
   }
 
-  if (t == kStrideAttribute) {
-    const char* use = "stride attribute";
-    return expect_paren_block(use, [&]() -> Result {
-      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::StrideAttribute>(t.source(), val.value);
-    });
-  }
-
   if (t == kSizeAttribute) {
     const char* use = "size attribute";
     return expect_paren_block(use, [&]() -> Result {
diff --git a/src/tint/reader/wgsl/parser_impl.h b/src/tint/reader/wgsl/parser_impl.h
index 50344e7..7cce8fd 100644
--- a/src/tint/reader/wgsl/parser_impl.h
+++ b/src/tint/reader/wgsl/parser_impl.h
@@ -419,11 +419,6 @@
   /// Parses a `type_decl` grammar element
   /// @returns the parsed Type or nullptr if none matched.
   Maybe<const ast::Type*> type_decl();
-  /// Parses a `type_decl` grammar element with the given pre-parsed
-  /// attributes.
-  /// @param attrs the list of attributes for the type.
-  /// @returns the parsed Type or nullptr if none matched.
-  Maybe<const ast::Type*> type_decl(ast::AttributeList& attrs);
   /// Parses a `storage_class` grammar element, erroring on parse failure.
   /// @param use a description of what was being parsed if an error was raised.
   /// @returns the storage class or StorageClass::kNone if none matched
@@ -849,8 +844,7 @@
   Expect<const ast::Type*> expect_type_decl_pointer(Token t);
   Expect<const ast::Type*> expect_type_decl_atomic(Token t);
   Expect<const ast::Type*> expect_type_decl_vector(Token t);
-  Expect<const ast::Type*> expect_type_decl_array(Token t,
-                                                  ast::AttributeList attrs);
+  Expect<const ast::Type*> expect_type_decl_array(Token t);
   Expect<const ast::Type*> expect_type_decl_matrix(Token t);
 
   Expect<const ast::Type*> expect_type(std::string_view use);
diff --git a/src/tint/reader/wgsl/parser_impl_error_msg_test.cc b/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
index 7c21634..f68c973 100644
--- a/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_error_msg_test.cc
@@ -445,9 +445,9 @@
 
 TEST_F(ParserImplErrorTest, FunctionScopeUnusedDecl) {
   EXPECT("fn f(a:i32)->i32{return a;@size(1)}",
-         R"(test.wgsl:1:28 error: unexpected attributes
+         R"(test.wgsl:1:27 error: expected '}'
 fn f(a:i32)->i32{return a;@size(1)}
-                           ^^^^
+                          ^
 )");
 }
 
@@ -774,31 +774,6 @@
 )");
 }
 
-TEST_F(ParserImplErrorTest, GlobalDeclVarArrayAttrNotArray) {
-  EXPECT("var i : @location(1) i32;",
-         R"(test.wgsl:1:10 error: unexpected attributes
-var i : @location(1) i32;
-         ^^^^^^^^
-)");
-}
-
-TEST_F(ParserImplErrorTest, GlobalDeclVarArrayStrideInvalid) {
-  EXPECT(
-      "var i : @stride(x) array<i32>;",
-      R"(test.wgsl:1:17 error: expected signed integer literal for stride attribute
-var i : @stride(x) array<i32>;
-                ^
-)");
-}
-
-TEST_F(ParserImplErrorTest, GlobalDeclVarArrayStrideNegative) {
-  EXPECT("var i : @stride(-1) array<i32>;",
-         R"(test.wgsl:1:17 error: stride attribute must be greater than 0
-var i : @stride(-1) array<i32>;
-                ^^
-)");
-}
-
 TEST_F(ParserImplErrorTest, GlobalDeclVarArrayMissingType) {
   EXPECT("var i : array<1, 3>;",
          R"(test.wgsl:1:15 error: invalid type for array declaration
diff --git a/src/tint/reader/wgsl/parser_impl_function_header_test.cc b/src/tint/reader/wgsl/parser_impl_function_header_test.cc
index 45e4d5d..eff2184 100644
--- a/src/tint/reader/wgsl/parser_impl_function_header_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_function_header_test.cc
@@ -75,27 +75,6 @@
   EXPECT_TRUE(f->return_type_attributes[0]->Is<ast::InvariantAttribute>());
 }
 
-TEST_F(ParserImplTest, FunctionHeader_AttributeReturnType_WithArrayStride) {
-  auto p = parser("fn main() -> @location(1) @stride(16) array<f32, 4>");
-  auto f = p->function_header();
-  ASSERT_FALSE(p->has_error()) << p->error();
-  EXPECT_TRUE(f.matched);
-  EXPECT_FALSE(f.errored);
-
-  EXPECT_EQ(f->name, "main");
-  EXPECT_EQ(f->params.size(), 0u);
-  ASSERT_EQ(f->return_type_attributes.size(), 1u);
-  auto* loc = f->return_type_attributes[0]->As<ast::LocationAttribute>();
-  ASSERT_TRUE(loc != nullptr);
-  EXPECT_EQ(loc->value, 1u);
-
-  auto* array_type = f->return_type->As<ast::Array>();
-  ASSERT_EQ(array_type->attributes.size(), 1u);
-  auto* stride = array_type->attributes[0]->As<ast::StrideAttribute>();
-  ASSERT_TRUE(stride != nullptr);
-  EXPECT_EQ(stride->stride, 16u);
-}
-
 TEST_F(ParserImplTest, FunctionHeader_MissingIdent) {
   auto p = parser("fn ()");
   auto f = p->function_header();
diff --git a/src/tint/reader/wgsl/parser_impl_global_decl_test.cc b/src/tint/reader/wgsl/parser_impl_global_decl_test.cc
index 964f498..29d29a6 100644
--- a/src/tint/reader/wgsl/parser_impl_global_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_global_decl_test.cc
@@ -166,33 +166,6 @@
   EXPECT_EQ(str->members.size(), 2u);
 }
 
-TEST_F(ParserImplTest, GlobalDecl_Struct_WithStride) {
-  auto p = parser("struct A { data: @stride(4) array<f32>; }");
-
-  p->expect_global_decl();
-  ASSERT_FALSE(p->has_error()) << p->error();
-
-  auto program = p->program();
-  ASSERT_EQ(program.AST().TypeDecls().size(), 1u);
-
-  auto* t = program.AST().TypeDecls()[0];
-  ASSERT_NE(t, nullptr);
-  ASSERT_TRUE(t->Is<ast::Struct>());
-
-  auto* str = t->As<ast::Struct>();
-  EXPECT_EQ(str->name, program.Symbols().Get("A"));
-  EXPECT_EQ(str->members.size(), 1u);
-
-  const auto* ty = str->members[0]->type;
-  ASSERT_TRUE(ty->Is<ast::Array>());
-  const auto* arr = ty->As<ast::Array>();
-
-  ASSERT_EQ(arr->attributes.size(), 1u);
-  auto* stride = arr->attributes[0];
-  ASSERT_TRUE(stride->Is<ast::StrideAttribute>());
-  ASSERT_EQ(stride->As<ast::StrideAttribute>()->stride, 4u);
-}
-
 TEST_F(ParserImplTest, GlobalDecl_Struct_Invalid) {
   auto p = parser("A {}");
   p->expect_global_decl();
diff --git a/src/tint/reader/wgsl/parser_impl_type_decl_test.cc b/src/tint/reader/wgsl/parser_impl_type_decl_test.cc
index d74ddf6..5845c40 100644
--- a/src/tint/reader/wgsl/parser_impl_type_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_type_decl_test.cc
@@ -439,172 +439,6 @@
   EXPECT_EQ(p->builder().Symbols().NameFor(count_expr->symbol), "size");
 }
 
-TEST_F(ParserImplTest, TypeDecl_Array_Stride) {
-  auto p = parser("@stride(16) array<f32, 5>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.matched);
-  EXPECT_FALSE(t.errored);
-  ASSERT_NE(t.value, nullptr) << p->error();
-  ASSERT_FALSE(p->has_error());
-  ASSERT_TRUE(t.value->Is<ast::Array>());
-
-  auto* a = t.value->As<ast::Array>();
-  ASSERT_FALSE(a->IsRuntimeArray());
-  ASSERT_TRUE(a->type->Is<ast::F32>());
-
-  auto* size = a->count->As<ast::SintLiteralExpression>();
-  ASSERT_NE(size, nullptr);
-  EXPECT_EQ(size->ValueAsI32(), 5);
-
-  ASSERT_EQ(a->attributes.size(), 1u);
-  auto* stride = a->attributes[0];
-  ASSERT_TRUE(stride->Is<ast::StrideAttribute>());
-  ASSERT_EQ(stride->As<ast::StrideAttribute>()->stride, 16u);
-  EXPECT_EQ(t.value->source.range, (Source::Range{{1u, 13u}, {1u, 26u}}));
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_Runtime_Stride) {
-  auto p = parser("@stride(16) array<f32>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.matched);
-  EXPECT_FALSE(t.errored);
-  ASSERT_NE(t.value, nullptr) << p->error();
-  ASSERT_FALSE(p->has_error());
-  ASSERT_TRUE(t.value->Is<ast::Array>());
-
-  auto* a = t.value->As<ast::Array>();
-  ASSERT_TRUE(a->IsRuntimeArray());
-  ASSERT_TRUE(a->type->Is<ast::F32>());
-
-  ASSERT_EQ(a->attributes.size(), 1u);
-  auto* stride = a->attributes[0];
-  ASSERT_TRUE(stride->Is<ast::StrideAttribute>());
-  ASSERT_EQ(stride->As<ast::StrideAttribute>()->stride, 16u);
-  EXPECT_EQ(t.value->source.range, (Source::Range{{1u, 13u}, {1u, 23u}}));
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_MultipleAttributes_OneBlock) {
-  auto p = parser("@stride(16) @stride(32) array<f32>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.matched);
-  EXPECT_FALSE(t.errored);
-  ASSERT_NE(t.value, nullptr) << p->error();
-  ASSERT_FALSE(p->has_error());
-  ASSERT_TRUE(t.value->Is<ast::Array>());
-
-  auto* a = t.value->As<ast::Array>();
-  ASSERT_TRUE(a->IsRuntimeArray());
-  ASSERT_TRUE(a->type->Is<ast::F32>());
-
-  auto& attrs = a->attributes;
-  ASSERT_EQ(attrs.size(), 2u);
-  EXPECT_TRUE(attrs[0]->Is<ast::StrideAttribute>());
-  EXPECT_EQ(attrs[0]->As<ast::StrideAttribute>()->stride, 16u);
-  EXPECT_TRUE(attrs[1]->Is<ast::StrideAttribute>());
-  EXPECT_EQ(attrs[1]->As<ast::StrideAttribute>()->stride, 32u);
-  EXPECT_EQ(t.value->source.range, (Source::Range{{1u, 25u}, {1u, 35u}}));
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_MultipleAttributes_MultipleBlocks) {
-  auto p = parser("@stride(16) @stride(32) array<f32>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.matched);
-  EXPECT_FALSE(t.errored);
-  ASSERT_NE(t.value, nullptr) << p->error();
-  ASSERT_FALSE(p->has_error());
-  ASSERT_TRUE(t.value->Is<ast::Array>());
-
-  auto* a = t.value->As<ast::Array>();
-  ASSERT_TRUE(a->IsRuntimeArray());
-  ASSERT_TRUE(a->type->Is<ast::F32>());
-
-  auto& attrs = a->attributes;
-  ASSERT_EQ(attrs.size(), 2u);
-  EXPECT_TRUE(attrs[0]->Is<ast::StrideAttribute>());
-  EXPECT_EQ(attrs[0]->As<ast::StrideAttribute>()->stride, 16u);
-  EXPECT_TRUE(attrs[1]->Is<ast::StrideAttribute>());
-  EXPECT_EQ(attrs[1]->As<ast::StrideAttribute>()->stride, 32u);
-  EXPECT_EQ(t.value->source.range, (Source::Range{{1u, 25u}, {1u, 35u}}));
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_Attribute_MissingArray) {
-  auto p = parser("@stride(16) f32");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.errored);
-  EXPECT_FALSE(t.matched);
-  ASSERT_EQ(t.value, nullptr);
-  ASSERT_TRUE(p->has_error());
-  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 attributes)");
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_Attribute_UnknownAttribute) {
-  auto p = parser("@unknown(16) array<f32, 5>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.errored);
-  EXPECT_FALSE(t.matched);
-  ASSERT_EQ(t.value, nullptr);
-  ASSERT_TRUE(p->has_error());
-  EXPECT_EQ(p->error(), R"(1:2: expected attribute)");
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_Stride_MissingLeftParen) {
-  auto p = parser("@stride 4) array<f32, 5>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.errored);
-  EXPECT_FALSE(t.matched);
-  ASSERT_EQ(t.value, nullptr);
-  ASSERT_TRUE(p->has_error());
-  EXPECT_EQ(p->error(), R"(1:9: expected '(' for stride attribute)");
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_Stride_MissingRightParen) {
-  auto p = parser("@stride(4 array<f32, 5>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.errored);
-  EXPECT_FALSE(t.matched);
-  ASSERT_EQ(t.value, nullptr);
-  ASSERT_TRUE(p->has_error());
-  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 attribute)");
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_Stride_MissingValue) {
-  auto p = parser("@stride() array<f32, 5>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.errored);
-  EXPECT_FALSE(t.matched);
-  ASSERT_EQ(t.value, nullptr);
-  ASSERT_TRUE(p->has_error());
-  EXPECT_EQ(p->error(),
-            "1:9: expected signed integer literal for stride attribute");
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_Stride_InvalidValue) {
-  auto p = parser("@stride(invalid) array<f32, 5>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.errored);
-  EXPECT_FALSE(t.matched);
-  ASSERT_EQ(t.value, nullptr);
-  ASSERT_TRUE(p->has_error());
-  EXPECT_EQ(p->error(),
-            "1:9: expected signed integer literal for stride attribute");
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_Stride_InvalidValue_Negative) {
-  auto p = parser("@stride(-1) array<f32, 5>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.errored);
-  EXPECT_FALSE(t.matched);
-  ASSERT_EQ(t.value, nullptr);
-  ASSERT_TRUE(p->has_error());
-  EXPECT_EQ(p->error(), "1:9: stride attribute must be greater than 0");
-}
-
 TEST_F(ParserImplTest, TypeDecl_Array_Runtime) {
   auto p = parser("array<u32>");
   auto t = p->type_decl();
diff --git a/src/tint/reader/wgsl/parser_impl_variable_ident_decl_test.cc b/src/tint/reader/wgsl/parser_impl_variable_ident_decl_test.cc
index 976284e..6f07375 100644
--- a/src/tint/reader/wgsl/parser_impl_variable_ident_decl_test.cc
+++ b/src/tint/reader/wgsl/parser_impl_variable_ident_decl_test.cc
@@ -75,35 +75,6 @@
   ASSERT_EQ(p->error(), "1:1: expected identifier for test");
 }
 
-TEST_F(ParserImplTest, VariableIdentDecl_NonAccessAttrFail) {
-  auto p = parser("my_var : @location(1) S");
-
-  auto* mem = Member("a", ty.i32(), ast::AttributeList{});
-  ast::StructMemberList members;
-  members.push_back(mem);
-
-  auto decl = p->expect_variable_ident_decl("test");
-  ASSERT_TRUE(p->has_error());
-  ASSERT_TRUE(decl.errored);
-  ASSERT_EQ(p->error(), "1:11: unexpected attributes");
-}
-
-TEST_F(ParserImplTest, VariableIdentDecl_AttributeMissingRightParen) {
-  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:22: expected ')' for location attribute");
-}
-
-TEST_F(ParserImplTest, VariableIdentDecl_AttributeMissingLeftParen) {
-  auto p = parser("my_var : @stride 4) S");
-  auto decl = p->expect_variable_ident_decl("test");
-  ASSERT_TRUE(p->has_error());
-  ASSERT_TRUE(decl.errored);
-  ASSERT_EQ(p->error(), "1:18: expected '(' for stride attribute");
-}
-
 }  // namespace
 }  // namespace wgsl
 }  // namespace reader
diff --git a/src/tint/transform/array_length_from_uniform_test.cc b/src/tint/transform/array_length_from_uniform_test.cc
index ffb63af..8bd6659 100644
--- a/src/tint/transform/array_length_from_uniform_test.cc
+++ b/src/tint/transform/array_length_from_uniform_test.cc
@@ -177,94 +177,6 @@
             got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
 }
 
-TEST_F(ArrayLengthFromUniformTest, WithStride) {
-  auto* src = R"(
-@group(0) @binding(0) var<storage, read> sb : @stride(64) array<i32>;
-
-@stage(compute) @workgroup_size(1)
-fn main() {
-  var len : u32 = arrayLength(&sb);
-}
-)";
-
-  auto* expect = R"(
-struct tint_symbol {
-  buffer_size : array<vec4<u32>, 1u>;
-}
-
-@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
-
-@group(0) @binding(0) var<storage, read> sb : @stride(64) array<i32>;
-
-@stage(compute) @workgroup_size(1)
-fn main() {
-  var len : u32 = (tint_symbol_1.buffer_size[0u][0u] / 64u);
-}
-)";
-
-  ArrayLengthFromUniform::Config cfg({0, 30u});
-  cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 0}, 0);
-
-  DataMap data;
-  data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
-
-  auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
-
-  EXPECT_EQ(expect, str(got));
-  EXPECT_EQ(std::unordered_set<uint32_t>({0}),
-            got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
-}
-
-TEST_F(ArrayLengthFromUniformTest, WithStride_InStruct) {
-  auto* src = R"(
-struct SB {
-  x : i32;
-  y : f32;
-  arr : @stride(64) array<i32>;
-};
-
-@group(0) @binding(0) var<storage, read> sb : SB;
-
-@stage(compute) @workgroup_size(1)
-fn main() {
-  var len : u32 = arrayLength(&sb.arr);
-}
-)";
-
-  auto* expect = R"(
-struct tint_symbol {
-  buffer_size : array<vec4<u32>, 1u>;
-}
-
-@group(0) @binding(30) var<uniform> tint_symbol_1 : tint_symbol;
-
-struct SB {
-  x : i32;
-  y : f32;
-  arr : @stride(64) array<i32>;
-}
-
-@group(0) @binding(0) var<storage, read> sb : SB;
-
-@stage(compute) @workgroup_size(1)
-fn main() {
-  var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 8u) / 64u);
-}
-)";
-
-  ArrayLengthFromUniform::Config cfg({0, 30u});
-  cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 0}, 0);
-
-  DataMap data;
-  data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
-
-  auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
-
-  EXPECT_EQ(expect, str(got));
-  EXPECT_EQ(std::unordered_set<uint32_t>({0}),
-            got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
-}
-
 TEST_F(ArrayLengthFromUniformTest, MultipleStorageBuffers) {
   auto* src = R"(
 struct SB1 {
diff --git a/src/tint/transform/calculate_array_length_test.cc b/src/tint/transform/calculate_array_length_test.cc
index 64589a0..c2ecbc9 100644
--- a/src/tint/transform/calculate_array_length_test.cc
+++ b/src/tint/transform/calculate_array_length_test.cc
@@ -287,78 +287,6 @@
   EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(CalculateArrayLengthTest, WithStride) {
-  auto* src = R"(
-@group(0) @binding(0) var<storage, read> sb : @stride(64) array<i32>;
-
-@stage(compute) @workgroup_size(1)
-fn main() {
-  var len : u32 = arrayLength(&sb);
-}
-)";
-
-  auto* expect = R"(
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : @stride(64) array<i32>, result : ptr<function, u32>)
-
-@group(0) @binding(0) var<storage, read> sb : @stride(64) array<i32>;
-
-@stage(compute) @workgroup_size(1)
-fn main() {
-  var tint_symbol_1 : u32 = 0u;
-  tint_symbol(sb, &(tint_symbol_1));
-  let tint_symbol_2 : u32 = (tint_symbol_1 / 64u);
-  var len : u32 = tint_symbol_2;
-}
-)";
-
-  auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(CalculateArrayLengthTest, WithStride_InStruct) {
-  auto* src = R"(
-struct SB {
-  x : i32;
-  y : f32;
-  arr : @stride(64) array<i32>;
-};
-
-@group(0) @binding(0) var<storage, read> sb : SB;
-
-@stage(compute) @workgroup_size(1)
-fn main() {
-  var len : u32 = arrayLength(&sb.arr);
-}
-)";
-
-  auto* expect = R"(
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
-
-struct SB {
-  x : i32;
-  y : f32;
-  arr : @stride(64) array<i32>;
-}
-
-@group(0) @binding(0) var<storage, read> sb : SB;
-
-@stage(compute) @workgroup_size(1)
-fn main() {
-  var tint_symbol_1 : u32 = 0u;
-  tint_symbol(sb, &(tint_symbol_1));
-  let tint_symbol_2 : u32 = ((tint_symbol_1 - 8u) / 64u);
-  var len : u32 = tint_symbol_2;
-}
-)";
-
-  auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
 TEST_F(CalculateArrayLengthTest, Nested) {
   auto* src = R"(
 struct SB {
diff --git a/src/tint/transform/decompose_memory_access_test.cc b/src/tint/transform/decompose_memory_access_test.cc
index 8ca3c2b..9d1956b 100644
--- a/src/tint/transform/decompose_memory_access_test.cc
+++ b/src/tint/transform/decompose_memory_access_test.cc
@@ -1970,22 +1970,26 @@
 
 TEST_F(DecomposeMemoryAccessTest, ComplexStaticAccessChain) {
   auto* src = R"(
+// sizeof(S1) == 32
+// alignof(S1) == 16
 struct S1 {
   a : i32;
   b : vec3<f32>;
   c : i32;
 };
 
+// sizeof(S2) == 116
+// alignof(S2) == 16
 struct S2 {
   a : i32;
-  b : @stride(32) array<S1, 3>;
+  b : array<S1, 3>;
   c : i32;
 };
 
 struct SB {
   @size(128)
   a : i32;
-  b : @stride(256) array<S2>;
+  b : array<S2>;
 };
 
 @group(0) @binding(0) var<storage, read_write> sb : SB;
@@ -1999,9 +2003,9 @@
   // sb.b[4].b[1].b.z
   //    ^  ^ ^  ^ ^ ^
   //    |  | |  | | |
-  //  128  | |1200| 1224
+  //  128  | |688 | 712
   //       | |    |
-  //    1152 1168 1216
+  //     640 656  704
 
   auto* expect = R"(
 struct S1 {
@@ -2012,14 +2016,14 @@
 
 struct S2 {
   a : i32;
-  b : @stride(32) array<S1, 3>;
+  b : array<S1, 3>;
   c : i32;
 }
 
 struct SB {
   @size(128)
   a : i32;
-  b : @stride(256) array<S2>;
+  b : array<S2>;
 }
 
 @group(0) @binding(0) var<storage, read_write> sb : SB;
@@ -2029,7 +2033,7 @@
 
 @stage(compute) @workgroup_size(1)
 fn main() {
-  var x : f32 = tint_symbol(sb, 1224u);
+  var x : f32 = tint_symbol(sb, 712u);
 }
 )";
 
@@ -2050,12 +2054,12 @@
 struct SB {
   @size(128)
   a : i32;
-  b : @stride(256) array<S2>;
+  b : array<S2>;
 };
 
 struct S2 {
   a : i32;
-  b : @stride(32) array<S1, 3>;
+  b : array<S1, 3>;
   c : i32;
 };
 
@@ -2069,9 +2073,9 @@
   // sb.b[4].b[1].b.z
   //    ^  ^ ^  ^ ^ ^
   //    |  | |  | | |
-  //  128  | |1200| 1224
+  //  128  | |688 | 712
   //       | |    |
-  //    1152 1168 1216
+  //     640 656  704
 
   auto* expect = R"(
 @internal(intrinsic_load_storage_f32) @internal(disable_validation__function_has_no_body)
@@ -2079,7 +2083,7 @@
 
 @stage(compute) @workgroup_size(1)
 fn main() {
-  var x : f32 = tint_symbol(sb, 1224u);
+  var x : f32 = tint_symbol(sb, 712u);
 }
 
 @group(0) @binding(0) var<storage, read_write> sb : SB;
@@ -2087,12 +2091,12 @@
 struct SB {
   @size(128)
   a : i32;
-  b : @stride(256) array<S2>;
+  b : array<S2>;
 }
 
 struct S2 {
   a : i32;
-  b : @stride(32) array<S1, 3>;
+  b : array<S1, 3>;
   c : i32;
 }
 
@@ -2118,14 +2122,14 @@
 
 struct S2 {
   a : i32;
-  b : @stride(32) array<S1, 3>;
+  b : array<S1, 3>;
   c : i32;
 };
 
 struct SB {
   @size(128)
   a : i32;
-  b : @stride(256) array<S2>;
+  b : array<S2>;
 };
 
 @group(0) @binding(0) var<storage, read_write> sb : SB;
@@ -2148,14 +2152,14 @@
 
 struct S2 {
   a : i32;
-  b : @stride(32) array<S1, 3>;
+  b : array<S1, 3>;
   c : i32;
 }
 
 struct SB {
   @size(128)
   a : i32;
-  b : @stride(256) array<S2>;
+  b : array<S2>;
 }
 
 @group(0) @binding(0) var<storage, read_write> sb : SB;
@@ -2168,7 +2172,7 @@
   var i : i32 = 4;
   var j : u32 = 1u;
   var k : i32 = 2;
-  var x : f32 = tint_symbol(sb, (((((128u + (256u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
+  var x : f32 = tint_symbol(sb, (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
 }
 )";
 
@@ -2192,12 +2196,12 @@
 struct SB {
   @size(128)
   a : i32;
-  b : @stride(256) array<S2>;
+  b :array<S2>;
 };
 
 struct S2 {
   a : i32;
-  b : @stride(32) array<S1, 3>;
+  b : array<S1, 3>;
   c : i32;
 };
 
@@ -2217,7 +2221,7 @@
   var i : i32 = 4;
   var j : u32 = 1u;
   var k : i32 = 2;
-  var x : f32 = tint_symbol(sb, (((((128u + (256u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
+  var x : f32 = tint_symbol(sb, (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
 }
 
 @group(0) @binding(0) var<storage, read_write> sb : SB;
@@ -2225,12 +2229,12 @@
 struct SB {
   @size(128)
   a : i32;
-  b : @stride(256) array<S2>;
+  b : array<S2>;
 }
 
 struct S2 {
   a : i32;
-  b : @stride(32) array<S1, 3>;
+  b : array<S1, 3>;
   c : i32;
 }
 
@@ -2256,7 +2260,7 @@
 
 type A1 = S1;
 
-type A1_Array = @stride(32) array<S1, 3>;
+type A1_Array = array<S1, 3>;
 
 struct S2 {
   a : i32;
@@ -2266,7 +2270,7 @@
 
 type A2 = S2;
 
-type A2_Array = @stride(256) array<S2>;
+type A2_Array = array<S2>;
 
 struct SB {
   @size(128)
@@ -2294,7 +2298,7 @@
 
 type A1 = S1;
 
-type A1_Array = @stride(32) array<S1, 3>;
+type A1_Array = array<S1, 3>;
 
 struct S2 {
   a : i32;
@@ -2304,7 +2308,7 @@
 
 type A2 = S2;
 
-type A2_Array = @stride(256) array<S2>;
+type A2_Array = array<S2>;
 
 struct SB {
   @size(128)
@@ -2322,7 +2326,7 @@
   var i : i32 = 4;
   var j : u32 = 1u;
   var k : i32 = 2;
-  var x : f32 = tint_symbol(sb, (((((128u + (256u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
+  var x : f32 = tint_symbol(sb, (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
 }
 )";
 
@@ -2350,7 +2354,7 @@
   b : A2_Array;
 };
 
-type A2_Array = @stride(256) array<S2>;
+type A2_Array = array<S2>;
 
 type A2 = S2;
 
@@ -2362,7 +2366,7 @@
 
 type A1 = S1;
 
-type A1_Array = @stride(32) array<S1, 3>;
+type A1_Array = array<S1, 3>;
 
 struct S1 {
   a : i32;
@@ -2380,7 +2384,7 @@
   var i : i32 = 4;
   var j : u32 = 1u;
   var k : i32 = 2;
-  var x : f32 = tint_symbol(sb, (((((128u + (256u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
+  var x : f32 = tint_symbol(sb, (((((128u + (128u * u32(i))) + 16u) + (32u * j)) + 16u) + (4u * u32(k))));
 }
 
 @group(0) @binding(0) var<storage, read_write> sb : SB;
@@ -2391,7 +2395,7 @@
   b : A2_Array;
 }
 
-type A2_Array = @stride(256) array<S2>;
+type A2_Array = array<S2>;
 
 type A2 = S2;
 
@@ -2403,7 +2407,7 @@
 
 type A1 = S1;
 
-type A1_Array = @stride(32) array<S1, 3>;
+type A1_Array = array<S1, 3>;
 
 struct S1 {
   a : i32;
diff --git a/src/tint/transform/glsl.cc b/src/tint/transform/glsl.cc
index 39f81e6..c7352c4 100644
--- a/src/tint/transform/glsl.cc
+++ b/src/tint/transform/glsl.cc
@@ -27,7 +27,6 @@
 #include "src/tint/transform/fold_trivial_single_use_lets.h"
 #include "src/tint/transform/loop_to_for_loop.h"
 #include "src/tint/transform/manager.h"
-#include "src/tint/transform/pad_array_elements.h"
 #include "src/tint/transform/promote_initializers_to_const_var.h"
 #include "src/tint/transform/promote_side_effects_to_decl.h"
 #include "src/tint/transform/remove_phonies.h"
@@ -105,7 +104,6 @@
   }
   manager.Add<PromoteInitializersToConstVar>();
 
-  manager.Add<PadArrayElements>();
   manager.Add<AddEmptyEntryPoint>();
   manager.Add<AddSpirvBlockAttribute>();
 
diff --git a/src/tint/transform/pad_array_elements.cc b/src/tint/transform/pad_array_elements.cc
deleted file mode 100644
index bf406df..0000000
--- a/src/tint/transform/pad_array_elements.cc
+++ /dev/null
@@ -1,177 +0,0 @@
-// Copyright 2021 The Tint Authors.
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "src/tint/transform/pad_array_elements.h"
-
-#include <unordered_map>
-#include <utility>
-
-#include "src/tint/program_builder.h"
-#include "src/tint/sem/array.h"
-#include "src/tint/sem/call.h"
-#include "src/tint/sem/expression.h"
-#include "src/tint/sem/type_constructor.h"
-#include "src/tint/utils/map.h"
-
-TINT_INSTANTIATE_TYPEINFO(tint::transform::PadArrayElements);
-
-namespace tint {
-namespace transform {
-namespace {
-
-using ArrayBuilder = std::function<const ast::Array*()>;
-
-/// PadArray returns a function that constructs a new array in `ctx.dst` with
-/// the element type padded to account for the explicit stride. PadArray will
-/// recursively pad arrays-of-arrays. The new array element type will be added
-/// to module-scope type declarations of `ctx.dst`.
-/// @param ctx the CloneContext
-/// @param create_ast_type_for Transform::CreateASTTypeFor()
-/// @param padded_arrays a map of src array type to the new array name
-/// @param array the array type
-/// @return the new AST array
-template <typename CREATE_AST_TYPE_FOR>
-ArrayBuilder PadArray(
-    CloneContext& ctx,
-    CREATE_AST_TYPE_FOR&& create_ast_type_for,
-    std::unordered_map<const sem::Array*, ArrayBuilder>& padded_arrays,
-    const sem::Array* array) {
-  if (array->IsStrideImplicit()) {
-    // We don't want to wrap arrays that have an implicit stride
-    return nullptr;
-  }
-
-  return utils::GetOrCreate(padded_arrays, array, [&] {
-    // Generate a unique name for the array element type
-    auto name = ctx.dst->Symbols().New("tint_padded_array_element");
-
-    // Examine the element type. Is it also an array?
-    const ast::Type* el_ty = nullptr;
-    if (auto* el_array = array->ElemType()->As<sem::Array>()) {
-      // Array of array - call PadArray() on the element type
-      if (auto p =
-              PadArray(ctx, create_ast_type_for, padded_arrays, el_array)) {
-        el_ty = p();
-      }
-    }
-
-    // If the element wasn't a padded array, just create the typical AST type
-    // for it
-    if (el_ty == nullptr) {
-      el_ty = create_ast_type_for(ctx, array->ElemType());
-    }
-
-    // Structure() will create and append the ast::Struct to the
-    // global declarations of `ctx.dst`. As we haven't finished building the
-    // current module-scope statement or function, this will be placed
-    // immediately before the usage.
-    ctx.dst->Structure(
-        name,
-        {ctx.dst->Member("el", el_ty, {ctx.dst->MemberSize(array->Stride())})});
-
-    auto* dst = ctx.dst;
-    return [=] {
-      if (array->IsRuntimeSized()) {
-        return dst->ty.array(dst->create<ast::TypeName>(name));
-      } else {
-        return dst->ty.array(dst->create<ast::TypeName>(name), array->Count());
-      }
-    };
-  });
-}
-
-}  // namespace
-
-PadArrayElements::PadArrayElements() = default;
-
-PadArrayElements::~PadArrayElements() = default;
-
-bool PadArrayElements::ShouldRun(const Program* program, const DataMap&) const {
-  for (auto* node : program->ASTNodes().Objects()) {
-    if (auto* var = node->As<ast::Type>()) {
-      if (auto* arr = program->Sem().Get<sem::Array>(var)) {
-        if (!arr->IsStrideImplicit()) {
-          return true;
-        }
-      }
-    }
-  }
-  return false;
-}
-
-void PadArrayElements::Run(CloneContext& ctx, const DataMap&, DataMap&) const {
-  auto& sem = ctx.src->Sem();
-
-  std::unordered_map<const sem::Array*, ArrayBuilder> padded_arrays;
-  auto pad = [&](const sem::Array* array) {
-    return PadArray(ctx, CreateASTTypeFor, padded_arrays, array);
-  };
-
-  // Replace all array types with their corresponding padded array type
-  ctx.ReplaceAll([&](const ast::Type* ast_type) -> const ast::Type* {
-    auto* type = ctx.src->TypeOf(ast_type);
-    if (auto* array = type->UnwrapRef()->As<sem::Array>()) {
-      if (auto p = pad(array)) {
-        return p();
-      }
-    }
-    return nullptr;
-  });
-
-  // Fix up index accessors so `a[1]` becomes `a[1].el`
-  ctx.ReplaceAll([&](const ast::IndexAccessorExpression* accessor)
-                     -> const ast::Expression* {
-    if (auto* array = tint::As<sem::Array>(
-            sem.Get(accessor->object)->Type()->UnwrapRef())) {
-      if (pad(array)) {
-        // Array element is wrapped in a structure. Emit a member accessor
-        // to get to the actual array element.
-        auto* idx = ctx.CloneWithoutTransform(accessor);
-        return ctx.dst->MemberAccessor(idx, "el");
-      }
-    }
-    return nullptr;
-  });
-
-  // Fix up array constructors so `A(1,2)` becomes
-  // `A(padded(1), padded(2))`
-  ctx.ReplaceAll(
-      [&](const ast::CallExpression* expr) -> const ast::Expression* {
-        auto* call = sem.Get(expr);
-        if (auto* ctor = call->Target()->As<sem::TypeConstructor>()) {
-          if (auto* array = ctor->ReturnType()->As<sem::Array>()) {
-            if (auto p = pad(array)) {
-              auto* arr_ty = p();
-              auto el_typename = arr_ty->type->As<ast::TypeName>()->name;
-
-              ast::ExpressionList args;
-              args.reserve(call->Arguments().size());
-              for (auto* arg : call->Arguments()) {
-                auto* val = ctx.Clone(arg->Declaration());
-                args.emplace_back(ctx.dst->Construct(
-                    ctx.dst->create<ast::TypeName>(el_typename), val));
-              }
-
-              return ctx.dst->Construct(arr_ty, args);
-            }
-          }
-        }
-        return nullptr;
-      });
-
-  ctx.Clone();
-}
-
-}  // namespace transform
-}  // namespace tint
diff --git a/src/tint/transform/pad_array_elements.h b/src/tint/transform/pad_array_elements.h
deleted file mode 100644
index cf5bfee..0000000
--- a/src/tint/transform/pad_array_elements.h
+++ /dev/null
@@ -1,62 +0,0 @@
-// Copyright 2021 The Tint Authors.
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#ifndef SRC_TINT_TRANSFORM_PAD_ARRAY_ELEMENTS_H_
-#define SRC_TINT_TRANSFORM_PAD_ARRAY_ELEMENTS_H_
-
-#include "src/tint/transform/transform.h"
-
-namespace tint {
-namespace transform {
-
-/// PadArrayElements is a transform that replaces array types with an explicit
-/// stride that is larger than the implicit stride, with an array of a new
-/// structure type. This structure holds with a single field of the element
-/// type, decorated with a `@size` attribute to pad the structure to the
-/// required array stride. The new array types have no explicit stride,
-/// structure size is equal to the desired stride.
-/// Array index expressions and constructors are also adjusted to deal with this
-/// structure element type.
-/// This transform helps with backends that cannot directly return arrays or use
-/// them as parameters.
-class PadArrayElements : public Castable<PadArrayElements, Transform> {
- public:
-  /// Constructor
-  PadArrayElements();
-
-  /// Destructor
-  ~PadArrayElements() override;
-
-  /// @param program the program to inspect
-  /// @param data optional extra transform-specific input data
-  /// @returns true if this transform should be run for the given program
-  bool ShouldRun(const Program* program,
-                 const DataMap& data = {}) const override;
-
- protected:
-  /// Runs the transform using the CloneContext built for transforming a
-  /// program. Run() is responsible for calling Clone() on the CloneContext.
-  /// @param ctx the CloneContext primed with the input program and
-  /// ProgramBuilder
-  /// @param inputs optional extra transform-specific input data
-  /// @param outputs optional extra transform-specific output data
-  void Run(CloneContext& ctx,
-           const DataMap& inputs,
-           DataMap& outputs) const override;
-};
-
-}  // namespace transform
-}  // namespace tint
-
-#endif  // SRC_TINT_TRANSFORM_PAD_ARRAY_ELEMENTS_H_
diff --git a/src/tint/transform/pad_array_elements_test.cc b/src/tint/transform/pad_array_elements_test.cc
deleted file mode 100644
index f5e921b..0000000
--- a/src/tint/transform/pad_array_elements_test.cc
+++ /dev/null
@@ -1,518 +0,0 @@
-// Copyright 2021 The Tint Authors.
-//
-// Licensed under the Apache License, Version 2.0 (the "License");
-// you may not use this file except in compliance with the License.
-// You may obtain a copy of the License at
-//
-//     http://www.apache.org/licenses/LICENSE-2.0
-//
-// Unless required by applicable law or agreed to in writing, software
-// distributed under the License is distributed on an "AS IS" BASIS,
-// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
-// See the License for the specific language governing permissions and
-// limitations under the License.
-
-#include "src/tint/transform/pad_array_elements.h"
-
-#include <utility>
-
-#include "src/tint/transform/test_helper.h"
-
-namespace tint {
-namespace transform {
-namespace {
-
-using PadArrayElementsTest = TransformTest;
-
-TEST_F(PadArrayElementsTest, ShouldRunEmptyModule) {
-  auto* src = R"()";
-
-  EXPECT_FALSE(ShouldRun<PadArrayElements>(src));
-}
-
-TEST_F(PadArrayElementsTest, ShouldRunHasImplicitArrayStride) {
-  auto* src = R"(
-var<private> arr : array<i32, 4>;
-)";
-
-  EXPECT_FALSE(ShouldRun<PadArrayElements>(src));
-}
-
-TEST_F(PadArrayElementsTest, ShouldRunHasExplicitArrayStride) {
-  auto* src = R"(
-var<private> arr : @stride(8) array<i32, 4>;
-)";
-
-  EXPECT_TRUE(ShouldRun<PadArrayElements>(src));
-}
-
-TEST_F(PadArrayElementsTest, EmptyModule) {
-  auto* src = "";
-  auto* expect = "";
-
-  auto got = Run<PadArrayElements>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, ImplicitArrayStride) {
-  auto* src = R"(
-var<private> arr : array<i32, 4>;
-)";
-  auto* expect = src;
-
-  auto got = Run<PadArrayElements>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, ArrayAsGlobal) {
-  auto* src = R"(
-var<private> arr : @stride(8) array<i32, 4>;
-)";
-  auto* expect = R"(
-struct tint_padded_array_element {
-  @size(8)
-  el : i32;
-}
-
-var<private> arr : array<tint_padded_array_element, 4u>;
-)";
-
-  auto got = Run<PadArrayElements>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, RuntimeArray) {
-  auto* src = R"(
-struct S {
-  rta : @stride(8) array<i32>;
-};
-)";
-  auto* expect = R"(
-struct tint_padded_array_element {
-  @size(8)
-  el : i32;
-}
-
-struct S {
-  rta : array<tint_padded_array_element>;
-}
-)";
-
-  auto got = Run<PadArrayElements>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, ArrayFunctionVar) {
-  auto* src = R"(
-fn f() {
-  var arr : @stride(16) array<i32, 4>;
-  arr = @stride(16) array<i32, 4>();
-  arr = @stride(16) array<i32, 4>(1, 2, 3, 4);
-  let x = arr[3];
-}
-)";
-  auto* expect = R"(
-struct tint_padded_array_element {
-  @size(16)
-  el : i32;
-}
-
-fn f() {
-  var arr : array<tint_padded_array_element, 4u>;
-  arr = array<tint_padded_array_element, 4u>();
-  arr = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
-  let x = arr[3].el;
-}
-)";
-
-  auto got = Run<PadArrayElements>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, ArrayAsParam) {
-  auto* src = R"(
-fn f(a : @stride(12) array<i32, 4>) -> i32 {
-  return a[2];
-}
-)";
-  auto* expect = R"(
-struct tint_padded_array_element {
-  @size(12)
-  el : i32;
-}
-
-fn f(a : array<tint_padded_array_element, 4u>) -> i32 {
-  return a[2].el;
-}
-)";
-
-  auto got = Run<PadArrayElements>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
-// TODO(crbug.com/tint/781): Cannot parse the stride on the return array type.
-TEST_F(PadArrayElementsTest, DISABLED_ArrayAsReturn) {
-  auto* src = R"(
-fn f() -> @stride(8) array<i32, 4> {
-  return array<i32, 4>(1, 2, 3, 4);
-}
-)";
-  auto* expect = R"(
-struct tint_padded_array_element {
-  el : i32;
-  @size(4)
-  padding : u32;
-};
-
-fn f() -> array<tint_padded_array_element, 4> {
-  return array<tint_padded_array_element, 4>(tint_padded_array_element(1, 0u), tint_padded_array_element(2, 0u), tint_padded_array_element(3, 0u), tint_padded_array_element(4, 0u));
-}
-)";
-
-  auto got = Run<PadArrayElements>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, ArrayAlias) {
-  auto* src = R"(
-type Array = @stride(16) array<i32, 4>;
-
-fn f() {
-  var arr : Array;
-  arr = Array();
-  arr = Array(1, 2, 3, 4);
-  let vals : Array = Array(1, 2, 3, 4);
-  arr = vals;
-  let x = arr[3];
-}
-)";
-  auto* expect = R"(
-struct tint_padded_array_element {
-  @size(16)
-  el : i32;
-}
-
-type Array = array<tint_padded_array_element, 4u>;
-
-fn f() {
-  var arr : array<tint_padded_array_element, 4u>;
-  arr = array<tint_padded_array_element, 4u>();
-  arr = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
-  let vals : array<tint_padded_array_element, 4u> = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
-  arr = vals;
-  let x = arr[3].el;
-}
-)";
-
-  auto got = Run<PadArrayElements>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, ArrayAlias_OutOfOrder) {
-  auto* src = R"(
-fn f() {
-  var arr : Array;
-  arr = Array();
-  arr = Array(1, 2, 3, 4);
-  let vals : Array = Array(1, 2, 3, 4);
-  arr = vals;
-  let x = arr[3];
-}
-
-type Array = @stride(16) array<i32, 4>;
-)";
-  auto* expect = R"(
-struct tint_padded_array_element {
-  @size(16)
-  el : i32;
-}
-
-fn f() {
-  var arr : array<tint_padded_array_element, 4u>;
-  arr = array<tint_padded_array_element, 4u>();
-  arr = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
-  let vals : array<tint_padded_array_element, 4u> = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
-  arr = vals;
-  let x = arr[3].el;
-}
-
-type Array = array<tint_padded_array_element, 4u>;
-)";
-
-  auto got = Run<PadArrayElements>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, ArraysInStruct) {
-  auto* src = R"(
-struct S {
-  a : @stride(8) array<i32, 4>;
-  b : @stride(8) array<i32, 8>;
-  c : @stride(8) array<i32, 4>;
-  d : @stride(12) array<i32, 8>;
-};
-)";
-  auto* expect = R"(
-struct tint_padded_array_element {
-  @size(8)
-  el : i32;
-}
-
-struct tint_padded_array_element_1 {
-  @size(8)
-  el : i32;
-}
-
-struct tint_padded_array_element_2 {
-  @size(12)
-  el : i32;
-}
-
-struct S {
-  a : array<tint_padded_array_element, 4u>;
-  b : array<tint_padded_array_element_1, 8u>;
-  c : array<tint_padded_array_element, 4u>;
-  d : array<tint_padded_array_element_2, 8u>;
-}
-)";
-
-  auto got = Run<PadArrayElements>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, ArraysOfArraysInStruct) {
-  auto* src = R"(
-struct S {
-  a : @stride(512) array<i32, 4>;
-  b : @stride(512) array<@stride(32) array<i32, 4>, 4>;
-  c : @stride(512) array<@stride(64) array<@stride(8) array<i32, 4>, 4>, 4>;
-};
-)";
-  auto* expect = R"(
-struct tint_padded_array_element {
-  @size(512)
-  el : i32;
-}
-
-struct tint_padded_array_element_2 {
-  @size(32)
-  el : i32;
-}
-
-struct tint_padded_array_element_1 {
-  @size(512)
-  el : array<tint_padded_array_element_2, 4u>;
-}
-
-struct tint_padded_array_element_5 {
-  @size(8)
-  el : i32;
-}
-
-struct tint_padded_array_element_4 {
-  @size(64)
-  el : array<tint_padded_array_element_5, 4u>;
-}
-
-struct tint_padded_array_element_3 {
-  @size(512)
-  el : array<tint_padded_array_element_4, 4u>;
-}
-
-struct S {
-  a : array<tint_padded_array_element, 4u>;
-  b : array<tint_padded_array_element_1, 4u>;
-  c : array<tint_padded_array_element_3, 4u>;
-}
-)";
-
-  auto got = Run<PadArrayElements>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, AccessArraysOfArraysInStruct) {
-  auto* src = R"(
-struct S {
-  a : @stride(512) array<i32, 4>;
-  b : @stride(512) array<@stride(32) array<i32, 4>, 4>;
-  c : @stride(512) array<@stride(64) array<@stride(8) array<i32, 4>, 4>, 4>;
-};
-
-fn f(s : S) -> i32 {
-  return s.a[2] + s.b[1][2] + s.c[3][1][2];
-}
-)";
-  auto* expect = R"(
-struct tint_padded_array_element {
-  @size(512)
-  el : i32;
-}
-
-struct tint_padded_array_element_2 {
-  @size(32)
-  el : i32;
-}
-
-struct tint_padded_array_element_1 {
-  @size(512)
-  el : array<tint_padded_array_element_2, 4u>;
-}
-
-struct tint_padded_array_element_5 {
-  @size(8)
-  el : i32;
-}
-
-struct tint_padded_array_element_4 {
-  @size(64)
-  el : array<tint_padded_array_element_5, 4u>;
-}
-
-struct tint_padded_array_element_3 {
-  @size(512)
-  el : array<tint_padded_array_element_4, 4u>;
-}
-
-struct S {
-  a : array<tint_padded_array_element, 4u>;
-  b : array<tint_padded_array_element_1, 4u>;
-  c : array<tint_padded_array_element_3, 4u>;
-}
-
-fn f(s : S) -> i32 {
-  return ((s.a[2].el + s.b[1].el[2].el) + s.c[3].el[1].el[2].el);
-}
-)";
-
-  auto got = Run<PadArrayElements>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, AccessArraysOfArraysInStruct_OutOfOrder) {
-  auto* src = R"(
-fn f(s : S) -> i32 {
-  return s.a[2] + s.b[1][2] + s.c[3][1][2];
-}
-
-struct S {
-  a : @stride(512) array<i32, 4>;
-  b : @stride(512) array<@stride(32) array<i32, 4>, 4>;
-  c : @stride(512) array<@stride(64) array<@stride(8) array<i32, 4>, 4>, 4>;
-};
-)";
-  auto* expect = R"(
-struct tint_padded_array_element {
-  @size(512)
-  el : i32;
-}
-
-struct tint_padded_array_element_1 {
-  @size(32)
-  el : i32;
-}
-
-struct tint_padded_array_element_2 {
-  @size(512)
-  el : array<tint_padded_array_element_1, 4u>;
-}
-
-struct tint_padded_array_element_3 {
-  @size(8)
-  el : i32;
-}
-
-struct tint_padded_array_element_4 {
-  @size(64)
-  el : array<tint_padded_array_element_3, 4u>;
-}
-
-struct tint_padded_array_element_5 {
-  @size(512)
-  el : array<tint_padded_array_element_4, 4u>;
-}
-
-fn f(s : S) -> i32 {
-  return ((s.a[2].el + s.b[1].el[2].el) + s.c[3].el[1].el[2].el);
-}
-
-struct S {
-  a : array<tint_padded_array_element, 4u>;
-  b : array<tint_padded_array_element_2, 4u>;
-  c : array<tint_padded_array_element_5, 4u>;
-}
-)";
-
-  auto got = Run<PadArrayElements>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
-TEST_F(PadArrayElementsTest, DeclarationOrder) {
-  auto* src = R"(
-type T0 = i32;
-
-type T1 = @stride(8) array<i32, 1>;
-
-type T2 = i32;
-
-fn f1(a : @stride(8) array<i32, 2>) {
-}
-
-type T3 = i32;
-
-fn f2() {
-  var v : @stride(8) array<i32, 3>;
-}
-)";
-  auto* expect = R"(
-type T0 = i32;
-
-struct tint_padded_array_element {
-  @size(8)
-  el : i32;
-}
-
-type T1 = array<tint_padded_array_element, 1u>;
-
-type T2 = i32;
-
-struct tint_padded_array_element_1 {
-  @size(8)
-  el : i32;
-}
-
-fn f1(a : array<tint_padded_array_element_1, 2u>) {
-}
-
-type T3 = i32;
-
-struct tint_padded_array_element_2 {
-  @size(8)
-  el : i32;
-}
-
-fn f2() {
-  var v : array<tint_padded_array_element_2, 3u>;
-}
-)";
-
-  auto got = Run<PadArrayElements>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
-}  // namespace
-}  // namespace transform
-}  // namespace tint
diff --git a/src/tint/transform/robustness_test.cc b/src/tint/transform/robustness_test.cc
index edc33f5..1d19987 100644
--- a/src/tint/transform/robustness_test.cc
+++ b/src/tint/transform/robustness_test.cc
@@ -1431,7 +1431,7 @@
 };
 @group(0) @binding(0) var<storage, read> s : S;
 
-type UArr = @stride(16) array<f32, 4>;
+type UArr = array<vec4<f32>, 4>;
 struct U {
   a : UArr;
 };
@@ -1451,11 +1451,11 @@
   var i32_sb4 : f32 = s.b[-1];
   var i32_sb5 : f32 = s.b[-4];
 
-  var i32_ua1 : f32 = u.a[4];
-  var i32_ua2 : f32 = u.a[1];
-  var i32_ua3 : f32 = u.a[0];
-  var i32_ua4 : f32 = u.a[-1];
-  var i32_ua5 : f32 = u.a[-4];
+  var i32_ua1 : f32 = u.a[4].x;
+  var i32_ua2 : f32 = u.a[1].x;
+  var i32_ua3 : f32 = u.a[0].x;
+  var i32_ua4 : f32 = u.a[-1].x;
+  var i32_ua5 : f32 = u.a[-4].x;
 
   // Unsigned
   var u32_sa1 : f32 = s.a[0u];
@@ -1472,12 +1472,12 @@
   var u32_sb5 : f32 = s.b[10u];
   var u32_sb6 : f32 = s.b[100u];
 
-  var u32_ua1 : f32 = u.a[0u];
-  var u32_ua2 : f32 = u.a[1u];
-  var u32_ua3 : f32 = u.a[3u];
-  var u32_ua4 : f32 = u.a[4u];
-  var u32_ua5 : f32 = u.a[10u];
-  var u32_ua6 : f32 = u.a[100u];
+  var u32_ua1 : f32 = u.a[0u].x;
+  var u32_ua2 : f32 = u.a[1u].x;
+  var u32_ua3 : f32 = u.a[3u].x;
+  var u32_ua4 : f32 = u.a[4u].x;
+  var u32_ua5 : f32 = u.a[10u].x;
+  var u32_ua6 : f32 = u.a[100u].x;
 }
 )";
 
@@ -1490,7 +1490,7 @@
 
 @group(0) @binding(0) var<storage, read> s : S;
 
-type UArr = @stride(16) array<f32, 4>;
+type UArr = array<vec4<f32>, 4>;
 
 struct U {
   a : UArr;
@@ -1509,11 +1509,11 @@
   var i32_sb3 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
   var i32_sb4 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
   var i32_sb5 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
-  var i32_ua1 : f32 = u.a[3];
-  var i32_ua2 : f32 = u.a[1];
-  var i32_ua3 : f32 = u.a[0];
-  var i32_ua4 : f32 = u.a[0];
-  var i32_ua5 : f32 = u.a[0];
+  var i32_ua1 : f32 = u.a[3].x;
+  var i32_ua2 : f32 = u.a[1].x;
+  var i32_ua3 : f32 = u.a[0].x;
+  var i32_ua4 : f32 = u.a[0].x;
+  var i32_ua5 : f32 = u.a[0].x;
   var u32_sa1 : f32 = s.a[0u];
   var u32_sa2 : f32 = s.a[1u];
   var u32_sa3 : f32 = s.a[3u];
@@ -1526,12 +1526,12 @@
   var u32_sb4 : f32 = s.b[min(4u, (arrayLength(&(s.b)) - 1u))];
   var u32_sb5 : f32 = s.b[min(10u, (arrayLength(&(s.b)) - 1u))];
   var u32_sb6 : f32 = s.b[min(100u, (arrayLength(&(s.b)) - 1u))];
-  var u32_ua1 : f32 = u.a[0u];
-  var u32_ua2 : f32 = u.a[1u];
-  var u32_ua3 : f32 = u.a[3u];
-  var u32_ua4 : f32 = u.a[3u];
-  var u32_ua5 : f32 = u.a[3u];
-  var u32_ua6 : f32 = u.a[3u];
+  var u32_ua1 : f32 = u.a[0u].x;
+  var u32_ua2 : f32 = u.a[1u].x;
+  var u32_ua3 : f32 = u.a[3u].x;
+  var u32_ua4 : f32 = u.a[3u].x;
+  var u32_ua5 : f32 = u.a[3u].x;
+  var u32_ua6 : f32 = u.a[3u].x;
 }
 )";
 
@@ -1553,7 +1553,7 @@
 
 @group(0) @binding(0) var<storage, read> s : S;
 
-type UArr = @stride(16) array<f32, 4>;
+type UArr = array<vec4<f32>, 4>;
 
 struct U {
   a : UArr;
@@ -1572,11 +1572,11 @@
   var i32_sb3 : f32 = s.b[0];
   var i32_sb4 : f32 = s.b[-1];
   var i32_sb5 : f32 = s.b[-4];
-  var i32_ua1 : f32 = u.a[3];
-  var i32_ua2 : f32 = u.a[1];
-  var i32_ua3 : f32 = u.a[0];
-  var i32_ua4 : f32 = u.a[0];
-  var i32_ua5 : f32 = u.a[0];
+  var i32_ua1 : f32 = u.a[3].x;
+  var i32_ua2 : f32 = u.a[1].x;
+  var i32_ua3 : f32 = u.a[0].x;
+  var i32_ua4 : f32 = u.a[0].x;
+  var i32_ua5 : f32 = u.a[0].x;
   var u32_sa1 : f32 = s.a[0u];
   var u32_sa2 : f32 = s.a[1u];
   var u32_sa3 : f32 = s.a[3u];
@@ -1589,12 +1589,12 @@
   var u32_sb4 : f32 = s.b[4u];
   var u32_sb5 : f32 = s.b[10u];
   var u32_sb6 : f32 = s.b[100u];
-  var u32_ua1 : f32 = u.a[0u];
-  var u32_ua2 : f32 = u.a[1u];
-  var u32_ua3 : f32 = u.a[3u];
-  var u32_ua4 : f32 = u.a[3u];
-  var u32_ua5 : f32 = u.a[3u];
-  var u32_ua6 : f32 = u.a[3u];
+  var u32_ua1 : f32 = u.a[0u].x;
+  var u32_ua2 : f32 = u.a[1u].x;
+  var u32_ua3 : f32 = u.a[3u].x;
+  var u32_ua4 : f32 = u.a[3u].x;
+  var u32_ua5 : f32 = u.a[3u].x;
+  var u32_ua6 : f32 = u.a[3u].x;
 }
 )";
 
@@ -1618,7 +1618,7 @@
 
 @group(0) @binding(0) var<storage, read> s : S;
 
-type UArr = @stride(16) array<f32, 4>;
+type UArr = array<vec4<f32>, 4>;
 
 struct U {
   a : UArr;
@@ -1637,11 +1637,11 @@
   var i32_sb3 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
   var i32_sb4 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
   var i32_sb5 : f32 = s.b[min(0u, (arrayLength(&(s.b)) - 1u))];
-  var i32_ua1 : f32 = u.a[4];
-  var i32_ua2 : f32 = u.a[1];
-  var i32_ua3 : f32 = u.a[0];
-  var i32_ua4 : f32 = u.a[-1];
-  var i32_ua5 : f32 = u.a[-4];
+  var i32_ua1 : f32 = u.a[4].x;
+  var i32_ua2 : f32 = u.a[1].x;
+  var i32_ua3 : f32 = u.a[0].x;
+  var i32_ua4 : f32 = u.a[-1].x;
+  var i32_ua5 : f32 = u.a[-4].x;
   var u32_sa1 : f32 = s.a[0u];
   var u32_sa2 : f32 = s.a[1u];
   var u32_sa3 : f32 = s.a[3u];
@@ -1654,12 +1654,12 @@
   var u32_sb4 : f32 = s.b[min(4u, (arrayLength(&(s.b)) - 1u))];
   var u32_sb5 : f32 = s.b[min(10u, (arrayLength(&(s.b)) - 1u))];
   var u32_sb6 : f32 = s.b[min(100u, (arrayLength(&(s.b)) - 1u))];
-  var u32_ua1 : f32 = u.a[0u];
-  var u32_ua2 : f32 = u.a[1u];
-  var u32_ua3 : f32 = u.a[3u];
-  var u32_ua4 : f32 = u.a[4u];
-  var u32_ua5 : f32 = u.a[10u];
-  var u32_ua6 : f32 = u.a[100u];
+  var u32_ua1 : f32 = u.a[0u].x;
+  var u32_ua2 : f32 = u.a[1u].x;
+  var u32_ua3 : f32 = u.a[3u].x;
+  var u32_ua4 : f32 = u.a[4u].x;
+  var u32_ua5 : f32 = u.a[10u].x;
+  var u32_ua6 : f32 = u.a[100u].x;
 }
 )";
 
@@ -1683,7 +1683,7 @@
 
 @group(0) @binding(0) var<storage, read> s : S;
 
-type UArr = @stride(16) array<f32, 4>;
+type UArr = array<vec4<f32>, 4>;
 
 struct U {
   a : UArr;
@@ -1702,11 +1702,11 @@
   var i32_sb3 : f32 = s.b[0];
   var i32_sb4 : f32 = s.b[-1];
   var i32_sb5 : f32 = s.b[-4];
-  var i32_ua1 : f32 = u.a[4];
-  var i32_ua2 : f32 = u.a[1];
-  var i32_ua3 : f32 = u.a[0];
-  var i32_ua4 : f32 = u.a[-1];
-  var i32_ua5 : f32 = u.a[-4];
+  var i32_ua1 : f32 = u.a[4].x;
+  var i32_ua2 : f32 = u.a[1].x;
+  var i32_ua3 : f32 = u.a[0].x;
+  var i32_ua4 : f32 = u.a[-1].x;
+  var i32_ua5 : f32 = u.a[-4].x;
   var u32_sa1 : f32 = s.a[0u];
   var u32_sa2 : f32 = s.a[1u];
   var u32_sa3 : f32 = s.a[3u];
@@ -1719,12 +1719,12 @@
   var u32_sb4 : f32 = s.b[4u];
   var u32_sb5 : f32 = s.b[10u];
   var u32_sb6 : f32 = s.b[100u];
-  var u32_ua1 : f32 = u.a[0u];
-  var u32_ua2 : f32 = u.a[1u];
-  var u32_ua3 : f32 = u.a[3u];
-  var u32_ua4 : f32 = u.a[4u];
-  var u32_ua5 : f32 = u.a[10u];
-  var u32_ua6 : f32 = u.a[100u];
+  var u32_ua1 : f32 = u.a[0u].x;
+  var u32_ua2 : f32 = u.a[1u].x;
+  var u32_ua3 : f32 = u.a[3u].x;
+  var u32_ua4 : f32 = u.a[4u].x;
+  var u32_ua5 : f32 = u.a[10u].x;
+  var u32_ua6 : f32 = u.a[100u].x;
 }
 )";
 
diff --git a/src/tint/transform/vertex_pulling.cc b/src/tint/transform/vertex_pulling.cc
index e515f04..e0d3128 100644
--- a/src/tint/transform/vertex_pulling.cc
+++ b/src/tint/transform/vertex_pulling.cc
@@ -258,7 +258,7 @@
         ctx.dst->Symbols().New(kStructName),
         {
             ctx.dst->Member(GetStructBufferName(),
-                            ctx.dst->ty.array<ProgramBuilder::u32>(4)),
+                            ctx.dst->ty.array<ProgramBuilder::u32>()),
         });
     for (uint32_t i = 0; i < cfg.vertex_state.size(); ++i) {
       // The decorated variable with struct type
diff --git a/src/tint/transform/vertex_pulling_test.cc b/src/tint/transform/vertex_pulling_test.cc
index c1e63a9..725b7de 100644
--- a/src/tint/transform/vertex_pulling_test.cc
+++ b/src/tint/transform/vertex_pulling_test.cc
@@ -108,7 +108,7 @@
 
   auto* expect = R"(
 struct TintVertexData {
-  tint_vertex_data : @stride(4) array<u32>;
+  tint_vertex_data : array<u32>;
 }
 
 @stage(vertex)
@@ -137,7 +137,7 @@
 
   auto* expect = R"(
 struct TintVertexData {
-  tint_vertex_data : @stride(4) array<u32>;
+  tint_vertex_data : array<u32>;
 }
 
 @binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -175,7 +175,7 @@
 
   auto* expect = R"(
 struct TintVertexData {
-  tint_vertex_data : @stride(4) array<u32>;
+  tint_vertex_data : array<u32>;
 }
 
 @binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -213,7 +213,7 @@
 
   auto* expect = R"(
 struct TintVertexData {
-  tint_vertex_data : @stride(4) array<u32>;
+  tint_vertex_data : array<u32>;
 }
 
 @binding(0) @group(5) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -256,7 +256,7 @@
 
   auto* expect = R"(
 struct TintVertexData {
-  tint_vertex_data : @stride(4) array<u32>;
+  tint_vertex_data : array<u32>;
 }
 
 @binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -304,7 +304,7 @@
 
   auto* expect = R"(
 struct TintVertexData {
-  tint_vertex_data : @stride(4) array<u32>;
+  tint_vertex_data : array<u32>;
 }
 
 @binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -364,7 +364,7 @@
 
   auto* expect = R"(
 struct TintVertexData {
-  tint_vertex_data : @stride(4) array<u32>;
+  tint_vertex_data : array<u32>;
 }
 
 @binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -444,7 +444,7 @@
 
   auto* expect = R"(
 struct TintVertexData {
-  tint_vertex_data : @stride(4) array<u32>;
+  tint_vertex_data : array<u32>;
 }
 
 @binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -526,7 +526,7 @@
 
   auto* expect = R"(
 struct TintVertexData {
-  tint_vertex_data : @stride(4) array<u32>;
+  tint_vertex_data : array<u32>;
 }
 
 @binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -603,7 +603,7 @@
 
   auto* expect = R"(
 struct TintVertexData {
-  tint_vertex_data : @stride(4) array<u32>;
+  tint_vertex_data : array<u32>;
 }
 
 @binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -670,7 +670,7 @@
 
   auto* expect = R"(
 struct TintVertexData {
-  tint_vertex_data : @stride(4) array<u32>;
+  tint_vertex_data : array<u32>;
 }
 
 @binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -715,7 +715,7 @@
 
   auto* expect = R"(
 struct TintVertexData {
-  tint_vertex_data : @stride(4) array<u32>;
+  tint_vertex_data : array<u32>;
 }
 
 @binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -771,7 +771,7 @@
 
   auto* expect = R"(
 struct TintVertexData {
-  tint_vertex_data_1 : @stride(4) array<u32>;
+  tint_vertex_data_1 : array<u32>;
 }
 
 @binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0_1 : TintVertexData;
@@ -848,7 +848,7 @@
 
   auto* expect = R"(
 struct TintVertexData {
-  tint_vertex_data : @stride(4) array<u32>;
+  tint_vertex_data : array<u32>;
 }
 
 @binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -1008,7 +1008,7 @@
   auto* expect =
       R"(
 struct TintVertexData {
-  tint_vertex_data : @stride(4) array<u32>;
+  tint_vertex_data : array<u32>;
 }
 
 @binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
@@ -1167,7 +1167,7 @@
 
   auto* expect = R"(
 struct TintVertexData {
-  tint_vertex_data : @stride(4) array<u32>;
+  tint_vertex_data : array<u32>;
 }
 
 @binding(0) @group(4) var<storage, read> tint_pulling_vertex_buffer_0 : TintVertexData;
diff --git a/src/tint/writer/glsl/generator_impl_member_accessor_test.cc b/src/tint/writer/glsl/generator_impl_member_accessor_test.cc
index 4bf817f..e297050 100644
--- a/src/tint/writer/glsl/generator_impl_member_accessor_test.cc
+++ b/src/tint/writer/glsl/generator_impl_member_accessor_test.cc
@@ -363,7 +363,7 @@
 TEST_F(GlslGeneratorImplTest_MemberAccessor,
        EmitExpression_IndexAccessor_StorageBuffer_Load_Int_FromArray) {
   // struct Data {
-  //   a : @stride(4) array<i32, 5>;
+  //   a : array<i32, 5>;
   // };
   // var<storage> data : Data;
   // data.a[2];
@@ -409,7 +409,7 @@
 TEST_F(GlslGeneratorImplTest_MemberAccessor,
        EmitExpression_IndexAccessor_StorageBuffer_Load_Int_FromArray_ExprIdx) {
   // struct Data {
-  //   a : @stride(4) array<i32, 5>;
+  //   a : array<i32, 5>;
   // };
   // var<storage> data : Data;
   // data.a[(2 + 4) - 3];
@@ -455,7 +455,7 @@
 
 TEST_F(GlslGeneratorImplTest_MemberAccessor, StorageBuffer_Store_ToArray) {
   // struct Data {
-  //   a : @stride(4) array<i32, 5>;
+  //   a : array<i32, 5>;
   // };
   // var<storage> data : Data;
   // data.a[2] = 2;
@@ -503,7 +503,7 @@
   //   b : vec3<f32>;
   // };
   // struct Data {
-  //   var c : @stride(32) array<Inner, 4>;
+  //   var c : array<Inner, 4>;
   // };
   //
   // var<storage> data : Pre;
@@ -562,7 +562,7 @@
   //   b : vec3<f32>;
   // };
   // struct Data {
-  //   var c : @stride(32) array<Inner, 4>;
+  //   var c : array<Inner, 4>;
   // };
   //
   // var<storage> data : Pre;
@@ -623,7 +623,7 @@
   //   b : vec3<f32>;
   // };
   // struct Data {
-  //   var c : @stride(32) array<Inner, 4>;
+  //   var c : array<Inner, 4>;
   // };
   //
   // var<storage> data : Pre;
@@ -684,7 +684,7 @@
   //   b : vec3<f32>;
   // };
   // struct Data {
-  //   var c : @stride(32) array<Inner, 4>;
+  //   var c : array<Inner, 4>;
   // };
   //
   // var<storage> data : Pre;
@@ -744,7 +744,7 @@
   //   b : vec3<f32>;
   // };
   // struct Data {
-  //   var c : @stride(32) array<Inner, 4>;
+  //   var c : array<Inner, 4>;
   // };
   //
   // var<storage> data : Pre;
@@ -802,7 +802,7 @@
   //   b : vec3<f32>;
   // };
   // struct Data {
-  //   var c : @stride(32) array<Inner, 4>;
+  //   var c : array<Inner, 4>;
   // };
   //
   // var<storage> data : Pre;
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index cf0a320..909fcff 100644
--- a/src/tint/writer/hlsl/generator_impl.cc
+++ b/src/tint/writer/hlsl/generator_impl.cc
@@ -57,7 +57,6 @@
 #include "src/tint/transform/loop_to_for_loop.h"
 #include "src/tint/transform/manager.h"
 #include "src/tint/transform/num_workgroups_from_uniform.h"
-#include "src/tint/transform/pad_array_elements.h"
 #include "src/tint/transform/promote_initializers_to_const_var.h"
 #include "src/tint/transform/promote_side_effects_to_decl.h"
 #include "src/tint/transform/remove_phonies.h"
@@ -209,7 +208,6 @@
   manager.Add<transform::CalculateArrayLength>();
   manager.Add<transform::PromoteInitializersToConstVar>();
 
-  manager.Add<transform::PadArrayElements>();
   manager.Add<transform::AddEmptyEntryPoint>();
 
   data.Add<transform::CanonicalizeEntryPointIO::Config>(
diff --git a/src/tint/writer/hlsl/generator_impl_member_accessor_test.cc b/src/tint/writer/hlsl/generator_impl_member_accessor_test.cc
index 7942ba4..79be125 100644
--- a/src/tint/writer/hlsl/generator_impl_member_accessor_test.cc
+++ b/src/tint/writer/hlsl/generator_impl_member_accessor_test.cc
@@ -391,7 +391,7 @@
 TEST_F(HlslGeneratorImplTest_MemberAccessor,
        EmitExpression_IndexAccessor_StorageBuffer_Load_Int_FromArray) {
   // struct Data {
-  //   a : @stride(4) array<i32, 5>;
+  //   a : array<i32, 5>;
   // };
   // var<storage> data : Data;
   // data.a[2];
@@ -423,7 +423,7 @@
 TEST_F(HlslGeneratorImplTest_MemberAccessor,
        EmitExpression_IndexAccessor_StorageBuffer_Load_Int_FromArray_ExprIdx) {
   // struct Data {
-  //   a : @stride(4) array<i32, 5>;
+  //   a : array<i32, 5>;
   // };
   // var<storage> data : Data;
   // data.a[(2 + 4) - 3];
@@ -455,7 +455,7 @@
 
 TEST_F(HlslGeneratorImplTest_MemberAccessor, StorageBuffer_Store_ToArray) {
   // struct Data {
-  //   a : @stride(4) array<i32, 5>;
+  //   a : array<i32, 5>;
   // };
   // var<storage> data : Data;
   // data.a[2] = 2;
@@ -489,7 +489,7 @@
   //   b : vec3<f32>;
   // };
   // struct Data {
-  //   var c : @stride(32) array<Inner, 4>;
+  //   var c : array<Inner, 4>;
   // };
   //
   // var<storage> data : Pre;
@@ -531,7 +531,7 @@
   //   b : vec3<f32>;
   // };
   // struct Data {
-  //   var c : @stride(32) array<Inner, 4>;
+  //   var c : array<Inner, 4>;
   // };
   //
   // var<storage> data : Pre;
@@ -575,7 +575,7 @@
   //   b : vec3<f32>;
   // };
   // struct Data {
-  //   var c : @stride(32) array<Inner, 4>;
+  //   var c : array<Inner, 4>;
   // };
   //
   // var<storage> data : Pre;
@@ -619,7 +619,7 @@
   //   b : vec3<f32>;
   // };
   // struct Data {
-  //   var c : @stride(32) array<Inner, 4>;
+  //   var c : array<Inner, 4>;
   // };
   //
   // var<storage> data : Pre;
@@ -662,7 +662,7 @@
   //   b : vec3<f32>;
   // };
   // struct Data {
-  //   var c : @stride(32) array<Inner, 4>;
+  //   var c : array<Inner, 4>;
   // };
   //
   // var<storage> data : Pre;
@@ -703,7 +703,7 @@
   //   b : vec3<f32>;
   // };
   // struct Data {
-  //   var c : @stride(32) array<Inner, 4>;
+  //   var c : array<Inner, 4>;
   // };
   //
   // var<storage> data : Pre;
diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc
index fad7968..78388dd 100644
--- a/src/tint/writer/msl/generator_impl.cc
+++ b/src/tint/writer/msl/generator_impl.cc
@@ -63,7 +63,6 @@
 #include "src/tint/transform/canonicalize_entry_point_io.h"
 #include "src/tint/transform/manager.h"
 #include "src/tint/transform/module_scope_var_to_entry_point_param.h"
-#include "src/tint/transform/pad_array_elements.h"
 #include "src/tint/transform/promote_initializers_to_const_var.h"
 #include "src/tint/transform/promote_side_effects_to_decl.h"
 #include "src/tint/transform/remove_phonies.h"
@@ -179,7 +178,6 @@
 
   manager.Add<transform::VectorizeScalarMatrixConstructors>();
   manager.Add<transform::WrapArraysInStructs>();
-  manager.Add<transform::PadArrayElements>();
   manager.Add<transform::RemovePhonies>();
   manager.Add<transform::SimplifyPointers>();
   // ArrayLengthFromUniform must come after SimplifyPointers, as
@@ -2945,9 +2943,8 @@
 
       [&](const sem::Array* arr) {
         if (!arr->IsStrideImplicit()) {
-          TINT_ICE(Writer, diagnostics_)
-              << "arrays with explicit strides should have "
-                 "removed with the PadArrayElements transform";
+          TINT_ICE(Writer, diagnostics_) << "arrays with explicit strides not "
+                                            "exist past the SPIR-V reader";
           return SizeAndAlign{};
         }
         auto num_els = std::max<uint32_t>(arr->Count(), 1);
diff --git a/src/tint/writer/msl/generator_impl_type_test.cc b/src/tint/writer/msl/generator_impl_type_test.cc
index 08f49a2..b59ca32 100644
--- a/src/tint/writer/msl/generator_impl_type_test.cc
+++ b/src/tint/writer/msl/generator_impl_type_test.cc
@@ -119,30 +119,6 @@
   EXPECT_EQ(out.str(), "bool ary[1]");
 }
 
-TEST_F(MslGeneratorImplTest, EmitType_ArrayWithStride) {
-  auto* s = Structure("s", {Member("arr", ty.array<f32, 4>(64))});
-  auto* ubo = Global("ubo", ty.Of(s), ast::StorageClass::kUniform,
-                     ast::AttributeList{
-                         create<ast::GroupAttribute>(0),
-                         create<ast::BindingAttribute>(1),
-                     });
-  WrapInFunction(MemberAccessor(ubo, "arr"));
-
-  GeneratorImpl& gen = SanitizeAndBuild();
-
-  ASSERT_TRUE(gen.Generate()) << gen.error();
-  EXPECT_THAT(gen.result(), HasSubstr(R"(struct tint_padded_array_element {
-  /* 0x0000 */ float el;
-  /* 0x0004 */ int8_t tint_pad[60];
-};)"));
-  EXPECT_THAT(gen.result(), HasSubstr(R"(struct tint_array_wrapper {
-  /* 0x0000 */ tint_padded_array_element arr[4];
-};)"));
-  EXPECT_THAT(gen.result(), HasSubstr(R"(struct s {
-  /* 0x0000 */ tint_array_wrapper arr;
-};)"));
-}
-
 TEST_F(MslGeneratorImplTest, EmitType_Bool) {
   auto* bool_ = create<sem::Bool>();
 
diff --git a/src/tint/writer/spirv/builder_builtin_test.cc b/src/tint/writer/spirv/builder_builtin_test.cc
index b7f6a6a..6738546 100644
--- a/src/tint/writer/spirv/builder_builtin_test.cc
+++ b/src/tint/writer/spirv/builder_builtin_test.cc
@@ -1622,7 +1622,7 @@
 
 TEST_F(BuiltinBuilderTest, Call_ArrayLength_ViaLets_WithPtrNoise) {
   // struct my_struct {
-  //   a : @stride(4) array<f32>;
+  //   a : array<f32>;
   // };
   // @binding(1) @group(2) var<storage, read> b : my_struct;
   //
diff --git a/test/tint/BUILD.gn b/test/tint/BUILD.gn
index 302fb65..834d948 100644
--- a/test/tint/BUILD.gn
+++ b/test/tint/BUILD.gn
@@ -327,7 +327,6 @@
     "../../src/tint/transform/module_scope_var_to_entry_point_param_test.cc",
     "../../src/tint/transform/multiplanar_external_texture_test.cc",
     "../../src/tint/transform/num_workgroups_from_uniform_test.cc",
-    "../../src/tint/transform/pad_array_elements_test.cc",
     "../../src/tint/transform/promote_initializers_to_const_var_test.cc",
     "../../src/tint/transform/promote_side_effects_to_decl_test.cc",
     "../../src/tint/transform/remove_phonies_test.cc",
diff --git a/test/tint/benchmark/animometer.wgsl.expected.glsl b/test/tint/benchmark/animometer.wgsl.expected.glsl
deleted file mode 100644
index 2a44a68..0000000
--- a/test/tint/benchmark/animometer.wgsl.expected.glsl
+++ /dev/null
@@ -1,98 +0,0 @@
-#version 310 es
-
-float tint_float_modulo(float lhs, float rhs) {
-  return (lhs - rhs * trunc(lhs / rhs));
-}
-
-
-layout(location = 0) in vec4 position_1;
-layout(location = 1) in vec4 color_1;
-layout(location = 0) out vec4 v_color_1;
-struct Time {
-  float value;
-};
-
-struct Uniforms {
-  float scale;
-  float offsetX;
-  float offsetY;
-  float scalar;
-  float scalarOffset;
-};
-
-layout(binding = 0) uniform Time_1 {
-  float value;
-} time;
-
-layout(binding = 1) uniform Uniforms_1 {
-  float scale;
-  float offsetX;
-  float offsetY;
-  float scalar;
-  float scalarOffset;
-} uniforms;
-
-struct VertexOutput {
-  vec4 Position;
-  vec4 v_color;
-};
-
-VertexOutput vert_main(vec4 position, vec4 color) {
-  float fade = tint_float_modulo((uniforms.scalarOffset + ((time.value * uniforms.scalar) / 10.0f)), 1.0f);
-  if ((fade < 0.5f)) {
-    fade = (fade * 2.0f);
-  } else {
-    fade = ((1.0f - fade) * 2.0f);
-  }
-  float xpos = (position.x * uniforms.scale);
-  float ypos = (position.y * uniforms.scale);
-  float angle = ((3.141590118f * 2.0f) * fade);
-  float xrot = ((xpos * cos(angle)) - (ypos * sin(angle)));
-  float yrot = ((xpos * sin(angle)) + (ypos * cos(angle)));
-  xpos = (xrot + uniforms.offsetX);
-  ypos = (yrot + uniforms.offsetY);
-  VertexOutput tint_symbol = VertexOutput(vec4(0.0f, 0.0f, 0.0f, 0.0f), vec4(0.0f, 0.0f, 0.0f, 0.0f));
-  tint_symbol.v_color = (vec4(fade, (1.0f - fade), 0.0f, 1.0f) + color);
-  tint_symbol.Position = vec4(xpos, ypos, 0.0f, 1.0f);
-  return tint_symbol;
-}
-
-void main() {
-  VertexOutput inner_result = vert_main(position_1, color_1);
-  gl_Position = inner_result.Position;
-  v_color_1 = inner_result.v_color;
-  gl_Position.y = -(gl_Position.y);
-  gl_Position.z = ((2.0f * gl_Position.z) - gl_Position.w);
-  return;
-}
-#version 310 es
-precision mediump float;
-
-layout(location = 0) in vec4 v_color_1;
-layout(location = 0) out vec4 value_1;
-struct Time {
-  float value;
-};
-
-struct Uniforms {
-  float scale;
-  float offsetX;
-  float offsetY;
-  float scalar;
-  float scalarOffset;
-};
-
-struct VertexOutput {
-  vec4 Position;
-  vec4 v_color;
-};
-
-vec4 frag_main(vec4 v_color) {
-  return v_color;
-}
-
-void main() {
-  vec4 inner_result = frag_main(v_color_1);
-  value_1 = inner_result;
-  return;
-}
diff --git a/test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl b/test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl
index 4021327..0397c89 100644
--- a/test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl
+++ b/test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl
@@ -48,7 +48,7 @@
   dirIntensity : f32;
   dirDirection : vec3<f32>;
   lightCount : u32;
-  lights : @stride(32) array<Light>;
+  lights : array<Light>;
 }
 
 @binding(2) @group(0) var<storage, read> globalLights : GlobalLights;
diff --git a/test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl.expected.glsl b/test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl.expected.glsl
deleted file mode 100644
index de73d59..0000000
--- a/test/tint/benchmark/skinned-shadowed-pbr-fragment.wgsl.expected.glsl
+++ /dev/null
@@ -1,381 +0,0 @@
-benchmark/skinned-shadowed-pbr-fragment.wgsl:51:13 warning: use of deprecated language feature: the @stride attribute is deprecated; use a larger type if necessary
-  lights : @stride(32) array<Light>;
-            ^^^^^^
-
-#version 310 es
-precision mediump float;
-
-layout(location = 0) in vec3 worldPos_1;
-layout(location = 1) in vec3 view_1;
-layout(location = 2) in vec2 texcoord_1;
-layout(location = 3) in vec2 texcoord2_1;
-layout(location = 4) in vec4 color_1;
-layout(location = 5) in vec4 instanceColor_1;
-layout(location = 6) in vec3 normal_1;
-layout(location = 7) in vec3 tangent_1;
-layout(location = 8) in vec3 bitangent_1;
-layout(location = 0) out vec4 color_2;
-layout(location = 1) out vec4 emissive_1;
-const float GAMMA = 2.200000048f;
-vec3 linearTosRGB(vec3 linear) {
-  float INV_GAMMA = (1.0f / GAMMA);
-  return pow(linear, vec3(INV_GAMMA));
-}
-
-struct Camera {
-  mat4 projection;
-  mat4 inverseProjection;
-  mat4 view;
-  vec3 position;
-  float time;
-  vec2 outputSize;
-  float zNear;
-  float zFar;
-};
-
-layout(binding = 0) uniform Camera_1 {
-  mat4 projection;
-  mat4 inverseProjection;
-  mat4 view;
-  vec3 position;
-  float time;
-  vec2 outputSize;
-  float zNear;
-  float zFar;
-} camera;
-
-struct ClusterLights {
-  uint offset;
-  uint count;
-};
-
-struct ClusterLightGroup {
-  uint offset;
-  ClusterLights lights[27648];
-  uint indices[1769472];
-};
-
-layout(binding = 1, std430) buffer ClusterLightGroup_1 {
-  uint offset;
-  ClusterLights lights[27648];
-  uint indices[1769472];
-} clusterLights;
-struct Light {
-  vec3 position;
-  float range;
-  vec3 color;
-  float intensity;
-};
-
-layout(binding = 2, std430) buffer GlobalLights_1 {
-  vec3 ambient;
-  vec3 dirColor;
-  float dirIntensity;
-  vec3 dirDirection;
-  uint lightCount;
-  Light lights[];
-} globalLights;
-const uvec3 tileCount = uvec3(32u, 18u, 48u);
-float linearDepth(float depthSample) {
-  return ((camera.zFar * camera.zNear) / ((depthSample) * ((camera.zNear - camera.zFar)) + (camera.zFar)));
-}
-
-uvec3 getTile(vec4 fragCoord) {
-  float sliceScale = (float(tileCount.z) / log2((camera.zFar / camera.zNear)));
-  float sliceBias = -(((float(tileCount.z) * log2(camera.zNear)) / log2((camera.zFar / camera.zNear))));
-  float tint_symbol_3 = linearDepth(fragCoord.z);
-  float tint_symbol_4 = log2(tint_symbol_3);
-  float tint_symbol_5 = max(((tint_symbol_4 * sliceScale) + sliceBias), 0.0f);
-  uint zTile = uint(tint_symbol_5);
-  return uvec3(uint((fragCoord.x / (camera.outputSize.x / float(tileCount.x)))), uint((fragCoord.y / (camera.outputSize.y / float(tileCount.y)))), zTile);
-}
-
-uint getClusterIndex(vec4 fragCoord) {
-  uvec3 tile = getTile(fragCoord);
-  return ((tile.x + (tile.y * tileCount.x)) + ((tile.z * tileCount.x) * tileCount.y));
-}
-
-layout(binding = 6, std430) buffer LightShadowTable_1 {
-  int light[];
-} lightShadowTable;
-vec2 shadowSampleOffsets[16] = vec2[16](vec2(-1.5f, -1.5f), vec2(-1.5f, -0.5f), vec2(-1.5f, 0.5f), vec2(-1.5f, 1.5f), vec2(-0.5f, -1.5f), vec2(-0.5f, -0.5f), vec2(-0.5f, 0.5f), vec2(-0.5f, 1.5f), vec2(0.5f, -1.5f), vec2(0.5f, -0.5f), vec2(0.5f, 0.5f), vec2(0.5f, 1.5f), vec2(1.5f, -1.5f), vec2(1.5f, -0.5f), vec2(1.5f, 0.5f), vec2(1.5f, 1.5f));
-const uint shadowSampleCount = 16u;
-struct ShadowProperties {
-  vec4 viewport;
-  mat4 viewProj;
-};
-
-layout(binding = 7, std430) buffer LightShadows_1 {
-  ShadowProperties properties[];
-} shadow;
-uniform highp sampler2D shadowTexture_1;
-uniform highp sampler2DShadow shadowTexture_shadowSampler;
-
-float dirLightVisibility(vec3 worldPos) {
-  int shadowIndex = lightShadowTable.light[0u];
-  if ((shadowIndex == -1)) {
-    return 1.0f;
-  }
-  vec4 viewport = shadow.properties[shadowIndex].viewport;
-  vec4 lightPos = (shadow.properties[shadowIndex].viewProj * vec4(worldPos, 1.0f));
-  vec3 shadowPos = vec3((((lightPos.xy / lightPos.w) * vec2(0.5f, -0.5f)) + vec2(0.5f, 0.5f)), (lightPos.z / lightPos.w));
-  vec2 viewportPos = vec2((viewport.xy + (shadowPos.xy * viewport.zw)));
-  vec2 texelSize = (1.0f / vec2(textureSize(shadowTexture_1, 0)));
-  vec4 clampRect = vec4((viewport.xy - texelSize), ((viewport.xy + viewport.zw) + texelSize));
-  float visibility = 0.0f;
-  {
-    for(uint i = 0u; (i < shadowSampleCount); i = (i + 1u)) {
-      visibility = (visibility + texture(shadowTexture_shadowSampler, vec3(clamp((viewportPos + (shadowSampleOffsets[i] * texelSize)), clampRect.xy, clampRect.zw), (shadowPos.z - 0.003f))));
-    }
-  }
-  return (visibility / float(shadowSampleCount));
-}
-
-int getCubeFace(vec3 v) {
-  vec3 vAbs = abs(v);
-  bool tint_tmp = (vAbs.z >= vAbs.x);
-  if (tint_tmp) {
-    tint_tmp = (vAbs.z >= vAbs.y);
-  }
-  if ((tint_tmp)) {
-    if ((v.z < 0.0f)) {
-      return 5;
-    }
-    return 4;
-  }
-  if ((vAbs.y >= vAbs.x)) {
-    if ((v.y < 0.0f)) {
-      return 3;
-    }
-    return 2;
-  }
-  if ((v.x < 0.0f)) {
-    return 1;
-  }
-  return 0;
-}
-
-float pointLightVisibility(uint lightIndex, vec3 worldPos, vec3 pointToLight) {
-  int shadowIndex = lightShadowTable.light[(lightIndex + 1u)];
-  if ((shadowIndex == -1)) {
-    return 1.0f;
-  }
-  int tint_symbol_6 = shadowIndex;
-  int tint_symbol_7 = getCubeFace((pointToLight * -1.0f));
-  shadowIndex = (tint_symbol_6 + tint_symbol_7);
-  vec4 viewport = shadow.properties[shadowIndex].viewport;
-  vec4 lightPos = (shadow.properties[shadowIndex].viewProj * vec4(worldPos, 1.0f));
-  vec3 shadowPos = vec3((((lightPos.xy / lightPos.w) * vec2(0.5f, -0.5f)) + vec2(0.5f, 0.5f)), (lightPos.z / lightPos.w));
-  vec2 viewportPos = vec2((viewport.xy + (shadowPos.xy * viewport.zw)));
-  vec2 texelSize = (1.0f / vec2(textureSize(shadowTexture_1, 0)));
-  vec4 clampRect = vec4(viewport.xy, (viewport.xy + viewport.zw));
-  float visibility = 0.0f;
-  {
-    for(uint i = 0u; (i < shadowSampleCount); i = (i + 1u)) {
-      visibility = (visibility + texture(shadowTexture_shadowSampler, vec3(clamp((viewportPos + (shadowSampleOffsets[i] * texelSize)), clampRect.xy, clampRect.zw), (shadowPos.z - 0.01f))));
-    }
-  }
-  return (visibility / float(shadowSampleCount));
-}
-
-struct VertexOutput {
-  vec4 position;
-  vec3 worldPos;
-  vec3 view;
-  vec2 texcoord;
-  vec2 texcoord2;
-  vec4 color;
-  vec4 instanceColor;
-  vec3 normal;
-  vec3 tangent;
-  vec3 bitangent;
-};
-
-struct Material {
-  vec4 baseColorFactor;
-  vec3 emissiveFactor;
-  float occlusionStrength;
-  vec2 metallicRoughnessFactor;
-  float alphaCutoff;
-};
-
-layout(binding = 8) uniform Material_1 {
-  vec4 baseColorFactor;
-  vec3 emissiveFactor;
-  float occlusionStrength;
-  vec2 metallicRoughnessFactor;
-  float alphaCutoff;
-} material;
-
-struct SurfaceInfo {
-  vec4 baseColor;
-  vec3 albedo;
-  float metallic;
-  float roughness;
-  vec3 normal;
-  vec3 f0;
-  float ao;
-  vec3 emissive;
-  vec3 v;
-};
-
-uniform highp sampler2D normalTexture_normalSampler;
-uniform highp sampler2D baseColorTexture_baseColorSampler;
-uniform highp sampler2D metallicRoughnessTexture_metallicRoughnessSampler;
-uniform highp sampler2D occlusionTexture_occlusionSampler;
-uniform highp sampler2D emissiveTexture_emissiveSampler;
-
-SurfaceInfo GetSurfaceInfo(VertexOutput tint_symbol) {
-  SurfaceInfo surface = SurfaceInfo(vec4(0.0f, 0.0f, 0.0f, 0.0f), vec3(0.0f, 0.0f, 0.0f), 0.0f, 0.0f, vec3(0.0f, 0.0f, 0.0f), vec3(0.0f, 0.0f, 0.0f), 0.0f, vec3(0.0f, 0.0f, 0.0f), vec3(0.0f, 0.0f, 0.0f));
-  surface.v = normalize(tint_symbol.view);
-  mat3 tbn = mat3(tint_symbol.tangent, tint_symbol.bitangent, tint_symbol.normal);
-  vec3 normalMap = texture(normalTexture_normalSampler, tint_symbol.texcoord).rgb;
-  surface.normal = normalize((tbn * ((2.0f * normalMap) - vec3(1.0f))));
-  vec4 baseColorMap = texture(baseColorTexture_baseColorSampler, tint_symbol.texcoord);
-  surface.baseColor = ((tint_symbol.color * material.baseColorFactor) * baseColorMap);
-  if ((surface.baseColor.a < material.alphaCutoff)) {
-    discard;
-  }
-  surface.albedo = surface.baseColor.rgb;
-  vec4 metallicRoughnessMap = texture(metallicRoughnessTexture_metallicRoughnessSampler, tint_symbol.texcoord);
-  surface.metallic = (material.metallicRoughnessFactor.x * metallicRoughnessMap.b);
-  surface.roughness = (material.metallicRoughnessFactor.y * metallicRoughnessMap.g);
-  vec3 dielectricSpec = vec3(0.039999999f);
-  surface.f0 = mix(dielectricSpec, surface.albedo, vec3(surface.metallic));
-  vec4 occlusionMap = texture(occlusionTexture_occlusionSampler, tint_symbol.texcoord);
-  surface.ao = (material.occlusionStrength * occlusionMap.r);
-  vec4 emissiveMap = texture(emissiveTexture_emissiveSampler, tint_symbol.texcoord);
-  surface.emissive = (material.emissiveFactor * emissiveMap.rgb);
-  if ((tint_symbol.instanceColor.a == 0.0f)) {
-    surface.albedo = (surface.albedo + tint_symbol.instanceColor.rgb);
-  } else {
-    surface.albedo = (surface.albedo * tint_symbol.instanceColor.rgb);
-  }
-  return surface;
-}
-
-const float PI = 3.141592741f;
-const uint LightType_Point = 0u;
-const uint LightType_Directional = 2u;
-struct PuctualLight {
-  uint lightType;
-  vec3 pointToLight;
-  float range;
-  vec3 color;
-  float intensity;
-};
-
-vec3 FresnelSchlick(float cosTheta, vec3 F0) {
-  return (F0 + ((vec3(1.0f) - F0) * pow((1.0f - cosTheta), 5.0f)));
-}
-
-float DistributionGGX(vec3 N, vec3 H, float roughness) {
-  float a_1 = (roughness * roughness);
-  float a2 = (a_1 * a_1);
-  float NdotH = max(dot(N, H), 0.0f);
-  float NdotH2 = (NdotH * NdotH);
-  float num = a2;
-  float denom = ((NdotH2 * (a2 - 1.0f)) + 1.0f);
-  return (num / ((PI * denom) * denom));
-}
-
-float GeometrySchlickGGX(float NdotV, float roughness) {
-  float r_1 = (roughness + 1.0f);
-  float k = ((r_1 * r_1) / 8.0f);
-  float num = NdotV;
-  float denom = ((NdotV * (1.0f - k)) + k);
-  return (num / denom);
-}
-
-float GeometrySmith(vec3 N, vec3 V, vec3 L, float roughness) {
-  float NdotV = max(dot(N, V), 0.0f);
-  float NdotL = max(dot(N, L), 0.0f);
-  float ggx2 = GeometrySchlickGGX(NdotV, roughness);
-  float ggx1 = GeometrySchlickGGX(NdotL, roughness);
-  return (ggx1 * ggx2);
-}
-
-float lightAttenuation(PuctualLight light) {
-  if ((light.lightType == LightType_Directional)) {
-    return 1.0f;
-  }
-  float tint_symbol_1 = length(light.pointToLight);
-  if ((light.range <= 0.0f)) {
-    return (1.0f / pow(tint_symbol_1, 2.0f));
-  }
-  return (clamp((1.0f - pow((tint_symbol_1 / light.range), 4.0f)), 0.0f, 1.0f) / pow(tint_symbol_1, 2.0f));
-}
-
-vec3 lightRadiance(PuctualLight light, SurfaceInfo surface) {
-  vec3 L = normalize(light.pointToLight);
-  vec3 H = normalize((surface.v + L));
-  float NDF = DistributionGGX(surface.normal, H, surface.roughness);
-  float G = GeometrySmith(surface.normal, surface.v, L, surface.roughness);
-  vec3 F = FresnelSchlick(max(dot(H, surface.v), 0.0f), surface.f0);
-  vec3 kD = ((vec3(1.0f) - F) * (1.0f - surface.metallic));
-  float NdotL = max(dot(surface.normal, L), 0.0f);
-  vec3 numerator = ((NDF * G) * F);
-  float denominator = max(((4.0f * max(dot(surface.normal, surface.v), 0.0f)) * NdotL), 0.001f);
-  vec3 specular = (numerator / vec3(denominator));
-  vec3 tint_symbol_8 = (light.color * light.intensity);
-  float tint_symbol_9 = lightAttenuation(light);
-  vec3 radiance = (tint_symbol_8 * tint_symbol_9);
-  return (((((kD * surface.albedo) / vec3(PI)) + specular) * radiance) * NdotL);
-}
-
-struct FragmentOutput {
-  vec4 color;
-  vec4 emissive;
-};
-
-uniform highp sampler2D ssaoTexture_1;
-uniform highp sampler2D ssaoTexture_defaultSampler;
-FragmentOutput fragmentMain(VertexOutput tint_symbol) {
-  SurfaceInfo surface = GetSurfaceInfo(tint_symbol);
-  vec3 Lo = vec3(0.0f, 0.0f, 0.0f);
-  if ((globalLights.dirIntensity > 0.0f)) {
-    PuctualLight light = PuctualLight(0u, vec3(0.0f, 0.0f, 0.0f), 0.0f, vec3(0.0f, 0.0f, 0.0f), 0.0f);
-    light.lightType = LightType_Directional;
-    light.pointToLight = globalLights.dirDirection;
-    light.color = globalLights.dirColor;
-    light.intensity = globalLights.dirIntensity;
-    float lightVis = dirLightVisibility(tint_symbol.worldPos);
-    vec3 tint_symbol_10 = Lo;
-    vec3 tint_symbol_11 = lightRadiance(light, surface);
-    Lo = (tint_symbol_10 + (tint_symbol_11 * lightVis));
-  }
-  uint clusterIndex = getClusterIndex(tint_symbol.position);
-  uint lightOffset = clusterLights.lights[clusterIndex].offset;
-  uint lightCount = clusterLights.lights[clusterIndex].count;
-  {
-    for(uint lightIndex = 0u; (lightIndex < lightCount); lightIndex = (lightIndex + 1u)) {
-      uint i = clusterLights.indices[(lightOffset + lightIndex)];
-      PuctualLight light = PuctualLight(0u, vec3(0.0f, 0.0f, 0.0f), 0.0f, vec3(0.0f, 0.0f, 0.0f), 0.0f);
-      light.lightType = LightType_Point;
-      light.pointToLight = (globalLights.lights[i].position.xyz - tint_symbol.worldPos);
-      light.range = globalLights.lights[i].range;
-      light.color = globalLights.lights[i].color;
-      light.intensity = globalLights.lights[i].intensity;
-      float lightVis = pointLightVisibility(i, tint_symbol.worldPos, light.pointToLight);
-      vec3 tint_symbol_12 = Lo;
-      vec3 tint_symbol_13 = lightRadiance(light, surface);
-      Lo = (tint_symbol_12 + (tint_symbol_13 * lightVis));
-    }
-  }
-  vec2 ssaoCoord = (tint_symbol.position.xy / vec2(textureSize(ssaoTexture_1, 0).xy));
-  float ssaoFactor = texture(ssaoTexture_defaultSampler, ssaoCoord).r;
-  vec3 ambient = (((globalLights.ambient * surface.albedo) * surface.ao) * ssaoFactor);
-  vec3 color = linearTosRGB(((Lo + ambient) + surface.emissive));
-  FragmentOutput tint_symbol_2 = FragmentOutput(vec4(0.0f, 0.0f, 0.0f, 0.0f), vec4(0.0f, 0.0f, 0.0f, 0.0f));
-  tint_symbol_2.color = vec4(color, surface.baseColor.a);
-  tint_symbol_2.emissive = vec4(surface.emissive, surface.baseColor.a);
-  return tint_symbol_2;
-}
-
-void main() {
-  VertexOutput tint_symbol_14 = VertexOutput(gl_FragCoord, worldPos_1, view_1, texcoord_1, texcoord2_1, color_1, instanceColor_1, normal_1, tangent_1, bitangent_1);
-  FragmentOutput inner_result = fragmentMain(tint_symbol_14);
-  color_2 = inner_result.color;
-  emissive_1 = inner_result.emissive;
-  return;
-}