Validate storage class constraints

As defined by https://gpuweb.github.io/gpuweb/wgsl/#storage-class-layout-constraints

Bug: tint:643
Change-Id: I9c78ba69a792a80c263a17b0a6e9b4810fdb7f30
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56780
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Antonio Maiorano <amaiorano@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>

src/CMakeLists.txt

@@ -637,6 +637,7 @@ if(${TINT_BUILD_TESTS})
     resolver/resolver_test_helper.cc
     resolver/resolver_test_helper.h
     resolver/resolver_test.cc
+    resolver/storage_class_layout_validation_test.cc
     resolver/storage_class_validation_test.cc
     resolver/struct_layout_test.cc
     resolver/struct_pipeline_stage_use_test.cc

src/inspector/inspector_test.cc

@@ -2239,11 +2239,14 @@ TEST_F(InspectorGetUniformBufferResourceBindingsTest, MultipleUniformBuffers) {
 }
 
 TEST_F(InspectorGetUniformBufferResourceBindingsTest, ContainingArray) {
-  // TODO(bclayton) - This is not a legal structure layout for uniform buffer
+  // Manually create uniform buffer to make sure it had a valid layout (array
-  // usage. Once crbug.com/tint/628 is implemented, this will fail validation
+  // with elem stride of 16, and that is 16-byte aligned within the struct)
-  // and will need to be fixed.
+  ast::Struct* foo_struct_type = Structure(
-  ast::Struct* foo_struct_type =
+      "foo_type",
-      MakeUniformBufferType("foo_type", {ty.i32(), ty.array<u32, 4>()});
+      {Member("0__i32", ty.i32()),
+       Member("b", ty.array(ty.u32(), 4, /*stride*/ 16), {MemberAlign(16)})},
+      {create<ast::StructBlockDecoration>()});
+
   AddUniformBuffer("foo_ub", ty.Of(foo_struct_type), 0, 0);
 
   MakeStructVariableReferenceBodyFunction("ub_func", "foo_ub", {{0, ty.i32()}});
@@ -2263,8 +2266,8 @@ TEST_F(InspectorGetUniformBufferResourceBindingsTest, ContainingArray) {
             result[0].resource_type);
   EXPECT_EQ(0u, result[0].bind_group);
   EXPECT_EQ(0u, result[0].binding);
-  EXPECT_EQ(20u, result[0].size);
+  EXPECT_EQ(80u, result[0].size);
-  EXPECT_EQ(20u, result[0].size_no_padding);
+  EXPECT_EQ(80u, result[0].size_no_padding);
 }
 
 TEST_F(InspectorGetStorageBufferResourceBindingsTest, Simple) {

src/program_builder.h

@@ -25,6 +25,7 @@
 #include "src/ast/assignment_statement.h"
 #include "src/ast/atomic.h"
 #include "src/ast/binary_expression.h"
+#include "src/ast/binding_decoration.h"
 #include "src/ast/bitcast_expression.h"
 #include "src/ast/bool.h"
 #include "src/ast/bool_literal.h"
@@ -54,6 +55,7 @@
 #include "src/ast/stage_decoration.h"
 #include "src/ast/storage_texture.h"
 #include "src/ast/stride_decoration.h"
+#include "src/ast/struct_block_decoration.h"
 #include "src/ast/struct_member_align_decoration.h"
 #include "src/ast/struct_member_offset_decoration.h"
 #include "src/ast/struct_member_size_decoration.h"
@@ -1379,7 +1381,9 @@ class ProgramBuilder {
   /// value.
   /// @returns a new `ast::Variable`, which is automatically registered as a
   /// global variable with the ast::Module.
-  template <typename NAME, typename... OPTIONAL>
+  template <typename NAME,
+            typename... OPTIONAL,
+            traits::EnableIfIsNotType<traits::Decay<NAME>, Source>* = nullptr>
   ast::Variable* Global(NAME&& name,
                         const ast::Type* type,
                         OPTIONAL&&... optional) {
@@ -1676,6 +1680,35 @@ class ProgramBuilder {
     return create<ast::StructMemberAlignDecoration>(source_, val);
   }
 
+  /// Creates a ast::StructBlockDecoration
+  /// @returns the struct block decoration pointer
+  ast::StructBlockDecoration* StructBlock() {
+    return create<ast::StructBlockDecoration>();
+  }
+
+  /// Creates the ast::GroupDecoration
+  /// @param value group decoration index
+  /// @returns the group decoration pointer
+  ast::GroupDecoration* Group(uint32_t value) {
+    return create<ast::GroupDecoration>(value);
+  }
+
+  /// Creates the ast::BindingDecoration
+  /// @param value the binding index
+  /// @returns the binding deocration pointer
+  ast::BindingDecoration* Binding(uint32_t value) {
+    return create<ast::BindingDecoration>(value);
+  }
+
+  /// Convenience function to create both a ast::GroupDecoration and
+  /// ast::BindingDecoration
+  /// @param group the group index
+  /// @param binding the binding index
+  /// @returns a decoration list with both the group and binding decorations
+  ast::DecorationList GroupAndBinding(uint32_t group, uint32_t binding) {
+    return {Group(group), Binding(binding)};
+  }
+
   /// Creates an ast::Function and registers it with the ast::Module.
   /// @param source the source information
   /// @param name the function name

src/resolver/resolver.cc

@@ -15,6 +15,8 @@
 #include "src/resolver/resolver.h"
 
 #include <algorithm>
+#include <cmath>
+#include <iomanip>
 #include <utility>
 
 #include "src/ast/alias.h"
@@ -685,6 +687,228 @@ bool Resolver::GlobalVariable(ast::Variable* var) {
     return false;
   }
 
+  // TODO(bclayton): Call this at the end of resolve on all uniform and storage
+  // referenced structs
+  if (!ValidateStorageClassLayout(info)) {
+    return false;
+  }
+
+  return true;
+}
+
+bool Resolver::ValidateStorageClassLayout(const sem::Struct* str,
+                                          ast::StorageClass sc) {
+  // https://gpuweb.github.io/gpuweb/wgsl/#storage-class-layout-constraints
+
+  auto is_uniform_struct_or_array = [sc](const sem::Type* ty) {
+    return sc == ast::StorageClass::kUniform &&
+           ty->IsAnyOf<sem::Array, sem::Struct>();
+  };
+
+  auto is_uniform_struct = [sc](const sem::Type* ty) {
+    return sc == ast::StorageClass::kUniform && ty->Is<sem::Struct>();
+  };
+
+  auto required_alignment_of = [&](const sem::Type* ty) {
+    uint32_t actual_align = 0;
+    uint32_t actual_size = 0;
+    DefaultAlignAndSize(ty, actual_align, actual_size);
+    uint32_t required_align = actual_align;
+    if (is_uniform_struct_or_array(ty)) {
+      required_align = utils::RoundUp(16u, actual_align);
+    }
+    return required_align;
+  };
+
+  auto member_name_of = [this](const sem::StructMember* sm) {
+    return builder_->Symbols().NameFor(sm->Declaration()->symbol());
+  };
+
+  auto type_name_of = [this](const sem::StructMember* sm) {
+    return sm->Declaration()->type()->FriendlyName(builder_->Symbols());
+  };
+
+  // TODO(amaiorano): Output struct and member decorations so that this output
+  // can be copied verbatim back into source
+  auto get_struct_layout_string = [&](const sem::Struct* st) -> std::string {
+    std::stringstream ss;
+
+    if (st->Members().empty()) {
+      TINT_ICE(Resolver, diagnostics_) << "Validation should have ensured that "
+                                          "structs have at least one member";
+      return {};
+    }
+    const auto* const last_member = st->Members().back();
+    const uint32_t last_member_struct_padding_offset =
+        last_member->Offset() + last_member->Size();
+
+    // Compute max widths to align output
+    const auto offset_w =
+        static_cast<int>(::log10(last_member_struct_padding_offset)) + 1;
+    const auto size_w = static_cast<int>(::log10(st->Size())) + 1;
+    const auto align_w = static_cast<int>(::log10(st->Align())) + 1;
+
+    auto print_struct_begin_line = [&](size_t align, size_t size,
+                                       std::string struct_name) {
+      ss << "/*          " << std::setw(offset_w) << " "
+         << "align(" << std::setw(align_w) << align << ") size("
+         << std::setw(size_w) << size << ") */ struct " << struct_name
+         << " {\n";
+    };
+
+    auto print_struct_end_line = [&]() {
+      ss << "/*                         "
+         << std::setw(offset_w + size_w + align_w) << " "
+         << "*/ };";
+    };
+
+    auto print_member_line = [&](size_t offset, size_t align, size_t size,
+                                 std::string s) {
+      ss << "/* offset(" << std::setw(offset_w) << offset << ") align("
+         << std::setw(align_w) << align << ") size(" << std::setw(size_w)
+         << size << ") */   " << s << ";\n";
+    };
+
+    print_struct_begin_line(st->Align(), st->Size(),
+                            st->FriendlyName(builder_->Symbols()));
+
+    for (size_t i = 0; i < st->Members().size(); ++i) {
+      auto* const m = st->Members()[i];
+
+      // Output field alignment padding, if any
+      auto* const prev_member = (i == 0) ? nullptr : str->Members()[i - 1];
+      if (prev_member) {
+        uint32_t padding =
+            m->Offset() - (prev_member->Offset() + prev_member->Size());
+        if (padding > 0) {
+          size_t padding_offset = m->Offset() - padding;
+          print_member_line(padding_offset, 1, padding,
+                            "// -- implicit field alignment padding --");
+        }
+      }
+
+      // Output member
+      std::string member_name = member_name_of(m);
+      print_member_line(m->Offset(), m->Align(), m->Size(),
+                        member_name_of(m) + " : " + type_name_of(m));
+    }
+
+    // Output struct size padding, if any
+    uint32_t struct_padding = st->Size() - last_member_struct_padding_offset;
+    if (struct_padding > 0) {
+      print_member_line(last_member_struct_padding_offset, 1, struct_padding,
+                        "// -- implicit struct size padding --");
+    }
+
+    print_struct_end_line();
+
+    return ss.str();
+  };
+
+  if (!ast::IsHostShareable(sc)) {
+    return true;
+  }
+
+  for (size_t i = 0; i < str->Members().size(); ++i) {
+    auto* const m = str->Members()[i];
+    uint32_t required_align = required_alignment_of(m->Type());
+
+    // Validate that member is at a valid byte offset
+    if (m->Offset() % required_align != 0) {
+      AddError("the offset of a struct member of type '" + type_name_of(m) +
+                   "' in storage class '" + ast::str(sc) +
+                   "' must be a multiple of " + std::to_string(required_align) +
+                   " bytes, but '" + member_name_of(m) +
+                   "' is currently at offset " + std::to_string(m->Offset()) +
+                   ". Consider setting [[align(" +
+                   std::to_string(required_align) + ")]] on this member",
+               m->Declaration()->source());
+
+      AddNote("see layout of struct:\n" + get_struct_layout_string(str),
+              str->Declaration()->source());
+
+      if (auto* member_str = m->Type()->As<sem::Struct>()) {
+        AddNote("and layout of struct member:\n" +
+                    get_struct_layout_string(member_str),
+                member_str->Declaration()->source());
+      }
+
+      return false;
+    }
+
+    // For uniform buffers, validate that the number of bytes between the
+    // previous member of type struct and the current is a multiple of 16 bytes.
+    auto* const prev_member = (i == 0) ? nullptr : str->Members()[i - 1];
+    if (prev_member && is_uniform_struct(prev_member->Type())) {
+      const uint32_t prev_to_curr_offset = m->Offset() - prev_member->Offset();
+      if (prev_to_curr_offset % 16 != 0) {
+        AddError(
+            "uniform storage requires that the number of bytes between the "
+            "start of the previous member of type struct and the current "
+            "member be a multiple of 16 bytes, but there are currently " +
+                std::to_string(prev_to_curr_offset) + " bytes between '" +
+                member_name_of(prev_member) + "' and '" + member_name_of(m) +
+                "'. Consider setting [[align(16)]] on this member",
+            m->Declaration()->source());
+
+        AddNote("see layout of struct:\n" + get_struct_layout_string(str),
+                str->Declaration()->source());
+
+        auto* prev_member_str = prev_member->Type()->As<sem::Struct>();
+        AddNote("and layout of previous member struct:\n" +
+                    get_struct_layout_string(prev_member_str),
+                prev_member_str->Declaration()->source());
+        return false;
+      }
+    }
+
+    // For uniform buffer array members, validate that array elements are
+    // aligned to 16 bytes
+    if (auto* arr = m->Type()->As<sem::Array>()) {
+      if (sc == ast::StorageClass::kUniform) {
+        // We already validated that this array member is itself aligned to 16
+        // bytes above, so we only need to validate that stride is a multiple of
+        // 16 bytes.
+        if (arr->Stride() % 16 != 0) {
+          AddError(
+              "uniform storage requires that array elements be aligned to 16 "
+              "bytes, but array stride 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",
+              m->Declaration()->type()->source());
+          AddNote("see layout of struct:\n" + get_struct_layout_string(str),
+                  str->Declaration()->source());
+          return false;
+        }
+      }
+    }
+
+    // If member is struct, recurse
+    if (auto* str_member = m->Type()->As<sem::Struct>()) {
+      // Cache result of struct + storage class pair
+      if (valid_struct_storage_layouts_.emplace(str_member, sc).second) {
+        if (!ValidateStorageClassLayout(str_member, sc)) {
+          return false;
+        }
+      }
+    }
+  }
+
+  return true;
+}
+
+bool Resolver::ValidateStorageClassLayout(const VariableInfo* info) {
+  if (auto* str = info->type->UnwrapRef()->As<sem::Struct>()) {
+    if (!ValidateStorageClassLayout(str, info->storage_class)) {
+      AddNote("see declaration of variable", info->declaration->source());
+      return false;
+    }
+  }
+
   return true;
 }
 
@@ -3636,7 +3860,6 @@ sem::Struct* Resolver::Structure(const ast::Struct* str) {
   // Validation of storage-class rules requires analysing the actual variable
   // usage of the structure, and so is performed as part of the variable
   // validation.
-  // TODO(crbug.com/tint/628): Actually implement storage-class validation.
   uint32_t struct_size = 0;
   uint32_t struct_align = 1;
 

src/resolver/resolver.h

@@ -16,9 +16,11 @@
 #define SRC_RESOLVER_RESOLVER_H_
 
 #include <memory>
+#include <set>
 #include <string>
 #include <unordered_map>
 #include <unordered_set>
+#include <utility>
 #include <vector>
 
 #include "src/intrinsic_table.h"
@@ -122,6 +124,9 @@ class Resolver {
     const sem::Intrinsic* intrinsic;
   };
 
+  std::set<std::pair<const sem::Struct*, ast::StorageClass>>
+      valid_struct_storage_layouts_;
+
   /// Structure holding semantic information about a function.
   /// Used to build the sem::Function nodes at the end of resolving.
   struct FunctionInfo {
@@ -312,6 +317,10 @@ class Resolver {
                                 const sem::Array* arr_type);
   bool ValidateTypeDecl(const ast::TypeDecl* named_type) const;
   bool ValidateNoDuplicateDecorations(const ast::DecorationList& decorations);
+  // sem::Struct is assumed to have at least one member
+  bool ValidateStorageClassLayout(const sem::Struct* type,
+                                  ast::StorageClass sc);
+  bool ValidateStorageClassLayout(const VariableInfo* info);
 
   /// @returns the sem::Type for the ast::Type `ty`, building it if it
   /// hasn't been constructed already. If an error is raised, nullptr is

src/resolver/storage_class_layout_validation_test.cc

@@ -0,0 +1,390 @@
+// 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/resolver/resolver.h"
+
+#include "gmock/gmock.h"
+#include "src/resolver/resolver_test_helper.h"
+
+namespace tint {
+namespace resolver {
+namespace {
+
+using ResolverStorageClassLayoutValidationTest = ResolverTest;
+
+// Detect unaligned member for storage buffers
+TEST_F(ResolverStorageClassLayoutValidationTest,
+       StorageBuffer_UnalignedMember) {
+  // [[block]]
+  // struct S {
+  //     [[size(5)]] a : f32;
+  //     [[align(1)]] b : f32;
+  // };
+  // [[group(0), binding(0)]]
+  // var<storage> a : S;
+
+  Structure(Source{{12, 34}}, "S",
+            {Member("a", ty.f32(), {MemberSize(5)}),
+             Member(Source{{34, 56}}, "b", ty.f32(), {MemberAlign(1)})},
+            {StructBlock()});
+
+  Global(Source{{78, 90}}, "a", ty.type_name("S"), ast::StorageClass::kStorage,
+         GroupAndBinding(0, 0));
+
+  ASSERT_FALSE(r()->Resolve());
+  EXPECT_EQ(
+      r()->error(),
+      R"(34:56 error: the offset of a struct member of type 'f32' in storage class 'storage' must be a multiple of 4 bytes, but 'b' is currently at offset 5. Consider setting [[align(4)]] on this member
+12:34 note: see layout of struct:
+/*           align(4) size(12) */ struct S {
+/* offset(0) align(4) size( 5) */   a : f32;
+/* offset(5) align(1) size( 4) */   b : f32;
+/* offset(9) align(1) size( 3) */   // -- implicit struct size padding --;
+/*                             */ };
+78:90 note: see declaration of variable)");
+}
+
+TEST_F(ResolverStorageClassLayoutValidationTest,
+       StorageBuffer_UnalignedMember_SuggestedFix) {
+  // [[block]]
+  // struct S {
+  //     [[size(5)]] a : f32;
+  //     [[align(4)]] b : f32;
+  // };
+  // [[group(0), binding(0)]]
+  // var<storage> a : S;
+
+  Structure(Source{{12, 34}}, "S",
+            {Member("a", ty.f32(), {MemberSize(5)}),
+             Member(Source{{34, 56}}, "b", ty.f32(), {MemberAlign(4)})},
+            {StructBlock()});
+
+  Global(Source{{78, 90}}, "a", ty.type_name("S"), ast::StorageClass::kStorage,
+         GroupAndBinding(0, 0));
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
+// Detect unaligned struct member for uniform buffers
+TEST_F(ResolverStorageClassLayoutValidationTest,
+       UniformBuffer_UnalignedMember_Struct) {
+  // struct Inner {
+  //   scalar : i32;
+  // };
+  //
+  // [[block]]
+  // struct Outer {
+  //   scalar : f32;
+  //   inner : Inner;
+  // };
+  //
+  // [[group(0), binding(0)]]
+  // var<uniform> a : Outer;
+
+  Structure(Source{{12, 34}}, "Inner", {Member("scalar", ty.i32())});
+
+  Structure(Source{{34, 56}}, "Outer",
+            {
+                Member("scalar", ty.f32()),
+                Member(Source{{56, 78}}, "inner", ty.type_name("Inner")),
+            },
+            {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"(56:78 error: the offset of a struct member of type 'Inner' in storage class 'uniform' must be a multiple of 16 bytes, but 'inner' is currently at offset 4. Consider setting [[align(16)]] on this member
+34:56 note: see layout of struct:
+/*           align(4) size(8) */ struct Outer {
+/* offset(0) align(4) size(4) */   scalar : f32;
+/* offset(4) align(4) size(4) */   inner : Inner;
+/*                            */ };
+12:34 note: and layout of struct member:
+/*           align(4) size(4) */ struct Inner {
+/* offset(0) align(4) size(4) */   scalar : i32;
+/*                            */ };
+78:90 note: see declaration of variable)");
+}
+
+TEST_F(ResolverStorageClassLayoutValidationTest,
+       UniformBuffer_UnalignedMember_Struct_SuggestedFix) {
+  // struct Inner {
+  //   scalar : i32;
+  // };
+  //
+  // [[block]]
+  // struct Outer {
+  //   scalar : f32;
+  //   [[align(16)]] inner : Inner;
+  // };
+  //
+  // [[group(0), binding(0)]]
+  // var<uniform> a : Outer;
+
+  Structure(Source{{12, 34}}, "Inner", {Member("scalar", ty.i32())});
+
+  Structure(Source{{34, 56}}, "Outer",
+            {
+                Member("scalar", ty.f32()),
+                Member(Source{{56, 78}}, "inner", ty.type_name("Inner"),
+                       {MemberAlign(16)}),
+            },
+            {StructBlock()});
+
+  Global(Source{{78, 90}}, "a", ty.type_name("Outer"),
+         ast::StorageClass::kUniform, GroupAndBinding(0, 0));
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
+// Detect unaligned array member for uniform buffers
+TEST_F(ResolverStorageClassLayoutValidationTest,
+       UniformBuffer_UnalignedMember_Array) {
+  // type Inner = [[stride(16)]] array<f32, 10>;
+  //
+  // [[block]]
+  // struct Outer {
+  //   scalar : f32;
+  //   inner : Inner;
+  // };
+  //
+  // [[group(0), binding(0)]]
+  // var<uniform> a : Outer;
+  Alias("Inner", ty.array(ty.f32(), 10, 16));
+
+  Structure(Source{{12, 34}}, "Outer",
+            {
+                Member("scalar", ty.f32()),
+                Member(Source{{56, 78}}, "inner", ty.type_name("Inner")),
+            },
+            {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"(56:78 error: the offset of a struct member of type 'Inner' in storage class 'uniform' must be a multiple of 16 bytes, but 'inner' is currently at offset 4. Consider setting [[align(16)]] on this member
+12:34 note: see layout of struct:
+/*             align(4) size(164) */ struct Outer {
+/* offset(  0) align(4) size(  4) */   scalar : f32;
+/* offset(  4) align(4) size(160) */   inner : Inner;
+/*                                */ };
+78:90 note: see declaration of variable)");
+}
+
+TEST_F(ResolverStorageClassLayoutValidationTest,
+       UniformBuffer_UnalignedMember_Array_SuggestedFix) {
+  // type Inner = [[stride(16)]] array<f32, 10>;
+  //
+  // [[block]]
+  // struct Outer {
+  //   scalar : f32;
+  //   [[align(16)]] inner : Inner;
+  // };
+  //
+  // [[group(0), binding(0)]]
+  // var<uniform> a : Outer;
+  Alias("Inner", ty.array(ty.f32(), 10, 16));
+
+  Structure(Source{{12, 34}}, "Outer",
+            {
+                Member("scalar", ty.f32()),
+                Member(Source{{34, 56}}, "inner", ty.type_name("Inner"),
+                       {MemberAlign(16)}),
+            },
+            {StructBlock()});
+
+  Global(Source{{78, 90}}, "a", ty.type_name("Outer"),
+         ast::StorageClass::kUniform, GroupAndBinding(0, 0));
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
+// Detect uniform buffers with byte offset between 2 members that is not a
+// multiple of 16 bytes
+TEST_F(ResolverStorageClassLayoutValidationTest,
+       UniformBuffer_MembersOffsetNotMultipleOf16) {
+  // struct Inner {
+  //   [[align(1), size(5)]] scalar : i32;
+  // };
+  //
+  // [[block]]
+  // struct Outer {
+  //   inner : Inner;
+  //   scalar : i32;
+  // };
+  //
+  // [[group(0), binding(0)]]
+  // var<uniform> a : Outer;
+
+  Structure(Source{{12, 34}}, "Inner",
+            {Member("scalar", ty.i32(), {MemberAlign(1), MemberSize(5)})});
+
+  Structure(Source{{34, 56}}, "Outer",
+            {
+                Member(Source{{56, 78}}, "inner", ty.type_name("Inner")),
+                Member(Source{{78, 90}}, "scalar", ty.i32()),
+            },
+            {StructBlock()});
+
+  Global(Source{{22, 24}}, "a", ty.type_name("Outer"),
+         ast::StorageClass::kUniform, GroupAndBinding(0, 0));
+
+  ASSERT_FALSE(r()->Resolve());
+  EXPECT_EQ(
+      r()->error(),
+      R"(78:90 error: uniform storage requires that the number of bytes between the start of the previous member of type struct and the current member be a multiple of 16 bytes, but there are currently 8 bytes between 'inner' and 'scalar'. Consider setting [[align(16)]] on this member
+34:56 note: see layout of struct:
+/*            align(4) size(12) */ struct Outer {
+/* offset( 0) align(1) size( 5) */   inner : Inner;
+/* offset( 5) align(1) size( 3) */   // -- implicit field alignment padding --;
+/* offset( 8) align(4) size( 4) */   scalar : i32;
+/*                              */ };
+12:34 note: and layout of previous member struct:
+/*           align(1) size(5) */ struct Inner {
+/* offset(0) align(1) size(5) */   scalar : i32;
+/*                            */ };
+22:24 note: see declaration of variable)");
+}
+
+TEST_F(ResolverStorageClassLayoutValidationTest,
+       UniformBuffer_MembersOffsetNotMultipleOf16_SuggestedFix) {
+  // struct Inner {
+  //   [[align(1), size(5)]] scalar : i32;
+  // };
+  //
+  // [[block]]
+  // struct Outer {
+  //   [[align(16)]] inner : Inner;
+  //   scalar : i32;
+  // };
+  //
+  // [[group(0), binding(0)]]
+  // var<uniform> a : Outer;
+
+  Structure(Source{{12, 34}}, "Inner",
+            {Member("scalar", ty.i32(), {MemberAlign(1), MemberSize(5)})});
+
+  Structure(Source{{34, 56}}, "Outer",
+            {
+                Member(Source{{56, 78}}, "inner", ty.type_name("Inner")),
+                Member(Source{{78, 90}}, "scalar", ty.i32(), {MemberAlign(16)}),
+            },
+            {StructBlock()});
+
+  Global(Source{{22, 34}}, "a", ty.type_name("Outer"),
+         ast::StorageClass::kUniform, GroupAndBinding(0, 0));
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
+// Make sure that this doesn't fail validation because vec3's align is 16, but
+// size is 12. 's' should be at offset 12, which is okay here.
+TEST_F(ResolverStorageClassLayoutValidationTest,
+       UniformBuffer_Vec3MemberOffset_NoFail) {
+  // [[block]]
+  // struct ScalarPackedAtEndOfVec3 {
+  //     v : vec3<f32>;
+  //     s : f32;
+  // };
+  // [[group(0), binding(0)]]
+  // var<uniform> a : ScalarPackedAtEndOfVec3;
+
+  Structure("ScalarPackedAtEndOfVec3",
+            {
+                Member("v", ty.vec3(ty.f32())),
+                Member("s", ty.f32()),
+            },
+            {StructBlock()});
+
+  Global(Source{{78, 90}}, "a", ty.type_name("ScalarPackedAtEndOfVec3"),
+         ast::StorageClass::kUniform, GroupAndBinding(0, 0));
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
+// 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>;
+  //
+  // [[block]]
+  // struct Outer {
+  //   inner : Inner;
+  //   scalar : i32;
+  // };
+  //
+  // [[group(0), binding(0)]]
+  // var<uniform> a : Outer;
+
+  Alias("Inner", ty.array(ty.f32(), 10, 8));
+
+  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 stride of 'inner' is currently 8. Consider setting [[stride(16)]] on the array type
+12:34 note: see layout of struct:
+/*            align(4) size(84) */ struct Outer {
+/* offset( 0) align(4) size(80) */   inner : Inner;
+/* offset(80) align(4) size( 4) */   scalar : i32;
+/*                              */ };
+78:90 note: see declaration of variable)");
+}
+
+TEST_F(ResolverStorageClassLayoutValidationTest,
+       UniformBuffer_InvalidArrayStride_SuggestedFix) {
+  // type Inner = [[stride(16)]] array<f32, 10>;
+  //
+  // [[block]]
+  // struct Outer {
+  //   inner : Inner;
+  //   scalar : i32;
+  // };
+  //
+  // [[group(0), binding(0)]]
+  // var<uniform> a : Outer;
+
+  Alias("Inner", ty.array(ty.f32(), 10, 16));
+
+  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_TRUE(r()->Resolve()) << r()->error();
+}
+
+}  // namespace
+}  // namespace resolver
+}  // namespace tint

src/resolver/storage_class_validation_test.cc

@@ -263,7 +263,7 @@ TEST_F(ResolverStorageClassValidationTest, UniformBufferNoError_Basic) {
              create<ast::GroupDecoration>(0),
          });
 
-  ASSERT_TRUE(r()->Resolve());
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverStorageClassValidationTest, UniformBufferNoError_Aliases) {
@@ -279,7 +279,7 @@ TEST_F(ResolverStorageClassValidationTest, UniformBufferNoError_Aliases) {
              create<ast::GroupDecoration>(0),
          });
 
-  ASSERT_TRUE(r()->Resolve());
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 }  // namespace

src/transform/array_length_from_uniform.cc

@@ -67,11 +67,16 @@ void ArrayLengthFromUniform::Run(CloneContext& ctx,
   ast::Variable* buffer_size_ubo = nullptr;
   auto get_ubo = [&]() {
     if (!buffer_size_ubo) {
+      // Emit an array<vec4<u32>, N>, where N is 1/4 number of elements.
+      // We do this because UBOs require an element stride that is 16-byte
+      // aligned.
       auto* buffer_size_struct = ctx.dst->Structure(
           ctx.dst->Sym(),
           {ctx.dst->Member(
               kBufferSizeMemberName,
-              ctx.dst->ty.array(ctx.dst->ty.u32(), max_buffer_size_index + 1))},
+              ctx.dst->ty.array(ctx.dst->ty.vec4(ctx.dst->ty.u32()),
+                                (max_buffer_size_index / 4) + 1))},
+
           ast::DecorationList{ctx.dst->create<ast::StructBlockDecoration>()});
       buffer_size_ubo = ctx.dst->Global(
           ctx.dst->Sym(), ctx.dst->ty.Of(buffer_size_struct),
@@ -99,18 +104,20 @@ void ArrayLengthFromUniform::Run(CloneContext& ctx,
 
     // Get the storage buffer that contains the runtime array.
     // We assume that the argument to `arrayLength` has the form
-    // `&resource.array`, which requires that `InlinePointerLets` and `Simplify`
+    // `&resource.array`, which requires that `InlinePointerLets` and
-    // have been run before this transform.
+    // `Simplify` have been run before this transform.
     auto* param = call_expr->params()[0]->As<ast::UnaryOpExpression>();
     if (!param || param->op() != ast::UnaryOp::kAddressOf) {
       TINT_ICE(Transform, ctx.dst->Diagnostics())
-          << "expected form of arrayLength argument to be &resource.array";
+          << "expected form of arrayLength argument to be "
+             "&resource.array";
       break;
     }
     auto* accessor = param->expr()->As<ast::MemberAccessorExpression>();
     if (!accessor) {
       TINT_ICE(Transform, ctx.dst->Diagnostics())
-          << "expected form of arrayLength argument to be &resource.array";
+          << "expected form of arrayLength argument to be "
+             "&resource.array";
       break;
     }
     auto* storage_buffer_expr = accessor->structure();
@@ -118,7 +125,8 @@ void ArrayLengthFromUniform::Run(CloneContext& ctx,
         sem.Get(storage_buffer_expr)->As<sem::VariableUser>();
     if (!storage_buffer_sem) {
       TINT_ICE(Transform, ctx.dst->Diagnostics())
-          << "expected form of arrayLength argument to be &resource.array";
+          << "expected form of arrayLength argument to be "
+             "&resource.array";
       break;
     }
 
@@ -135,9 +143,13 @@ void ArrayLengthFromUniform::Run(CloneContext& ctx,
     }
 
     // Load the total storage buffer size from the UBO.
-    auto* total_storage_buffer_size = ctx.dst->IndexAccessor(
+    uint32_t array_index = idx_itr->second / 4;
+    auto* vec_expr = ctx.dst->IndexAccessor(
         ctx.dst->MemberAccessor(get_ubo()->symbol(), kBufferSizeMemberName),
-        idx_itr->second);
+        array_index);
+    uint32_t vec_index = idx_itr->second % 4;
+    auto* total_storage_buffer_size =
+        ctx.dst->IndexAccessor(vec_expr, vec_index);
 
     // Calculate actual array length
     //                total_storage_buffer_size - array_offset

src/transform/array_length_from_uniform_test.cc

@@ -81,7 +81,7 @@ fn main() {
   auto* expect = R"(
 [[block]]
 struct tint_symbol {
-  buffer_size : array<u32, 1>;
+  buffer_size : array<vec4<u32>, 1>;
 };
 
 [[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
@@ -96,7 +96,7 @@ struct SB {
 
 [[stage(compute), workgroup_size(1)]]
 fn main() {
-  var len : u32 = ((tint_symbol_1.buffer_size[0u] - 4u) / 4u);
+  var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
 }
 )";
 
@@ -134,7 +134,7 @@ fn main() {
   auto* expect = R"(
 [[block]]
 struct tint_symbol {
-  buffer_size : array<u32, 1>;
+  buffer_size : array<vec4<u32>, 1>;
 };
 
 [[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
@@ -150,7 +150,7 @@ struct SB {
 
 [[stage(compute), workgroup_size(1)]]
 fn main() {
-  var len : u32 = ((tint_symbol_1.buffer_size[0u] - 8u) / 64u);
+  var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 8u) / 64u);
 }
 )";
 
@@ -175,29 +175,48 @@ struct SB1 {
   x : i32;
   arr1 : array<i32>;
 };
-
 [[block]]
 struct SB2 {
   x : i32;
   arr2 : array<vec4<f32>>;
 };
+[[block]]
+struct SB3 {
+  x : i32;
+  arr3 : array<vec4<f32>>;
+};
+[[block]]
+struct SB4 {
+  x : i32;
+  arr4 : array<vec4<f32>>;
+};
+[[block]]
+struct SB5 {
+  x : i32;
+  arr5 : array<vec4<f32>>;
+};
 
 [[group(0), binding(2)]] var<storage, read> sb1 : SB1;
-
 [[group(1), binding(2)]] var<storage, read> sb2 : SB2;
+[[group(2), binding(2)]] var<storage, read> sb3 : SB3;
+[[group(3), binding(2)]] var<storage, read> sb4 : SB4;
+[[group(4), binding(2)]] var<storage, read> sb5 : SB5;
 
 [[stage(compute), workgroup_size(1)]]
 fn main() {
   var len1 : u32 = arrayLength(&(sb1.arr1));
   var len2 : u32 = arrayLength(&(sb2.arr2));
-  var x : u32 = (len1 + len2);
+  var len3 : u32 = arrayLength(&(sb3.arr3));
+  var len4 : u32 = arrayLength(&(sb4.arr4));
+  var len5 : u32 = arrayLength(&(sb5.arr5));
+  var x : u32 = (len1 + len2 + len3 + len4 + len5);
 }
 )";
 
   auto* expect = R"(
 [[block]]
 struct tint_symbol {
-  buffer_size : array<u32, 2>;
+  buffer_size : array<vec4<u32>, 2>;
 };
 
 [[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
@@ -214,21 +233,51 @@ struct SB2 {
   arr2 : array<vec4<f32>>;
 };
 
+[[block]]
+struct SB3 {
+  x : i32;
+  arr3 : array<vec4<f32>>;
+};
+
+[[block]]
+struct SB4 {
+  x : i32;
+  arr4 : array<vec4<f32>>;
+};
+
+[[block]]
+struct SB5 {
+  x : i32;
+  arr5 : array<vec4<f32>>;
+};
+
 [[group(0), binding(2)]] var<storage, read> sb1 : SB1;
 
 [[group(1), binding(2)]] var<storage, read> sb2 : SB2;
 
+[[group(2), binding(2)]] var<storage, read> sb3 : SB3;
+
+[[group(3), binding(2)]] var<storage, read> sb4 : SB4;
+
+[[group(4), binding(2)]] var<storage, read> sb5 : SB5;
+
 [[stage(compute), workgroup_size(1)]]
 fn main() {
-  var len1 : u32 = ((tint_symbol_1.buffer_size[0u] - 4u) / 4u);
+  var len1 : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
-  var len2 : u32 = ((tint_symbol_1.buffer_size[1u] - 16u) / 16u);
+  var len2 : u32 = ((tint_symbol_1.buffer_size[0u][1u] - 16u) / 16u);
-  var x : u32 = (len1 + len2);
+  var len3 : u32 = ((tint_symbol_1.buffer_size[0u][2u] - 16u) / 16u);
+  var len4 : u32 = ((tint_symbol_1.buffer_size[0u][3u] - 16u) / 16u);
+  var len5 : u32 = ((tint_symbol_1.buffer_size[1u][0u] - 16u) / 16u);
+  var x : u32 = ((((len1 + len2) + len3) + len4) + len5);
 }
 )";
 
   ArrayLengthFromUniform::Config cfg({0, 30u});
   cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 2u}, 0);
   cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{1u, 2u}, 1);
+  cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{2u, 2u}, 2);
+  cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{3u, 2u}, 3);
+  cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{4u, 2u}, 4);
 
   DataMap data;
   data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));

test/BUILD.gn

@@ -246,6 +246,7 @@ tint_unittests_source_set("tint_unittests_core_src") {
     "../src/resolver/resolver_test.cc",
     "../src/resolver/resolver_test_helper.cc",
     "../src/resolver/resolver_test_helper.h",
+    "../src/resolver/storage_class_layout_validation_test.cc",
     "../src/resolver/storage_class_validation_test.cc",
     "../src/resolver/struct_layout_test.cc",
     "../src/resolver/struct_pipeline_stage_use_test.cc",

test/intrinsics/arrayLength/complex_via_let.wgsl.expected.msl

@@ -2,14 +2,14 @@
 
 using namespace metal;
 struct tint_symbol_1 {
-  /* 0x0000 */ uint buffer_size[1];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct S {
   /* 0x0000 */ int a[1];
 };
 
 kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) {
-  uint const l1 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u);
+  uint const l1 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u);
   return;
 }
 

test/intrinsics/arrayLength/deprecated.wgsl.expected.msl

@@ -2,15 +2,15 @@
 
 using namespace metal;
 struct tint_symbol_1 {
-  /* 0x0000 */ uint buffer_size[1];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct S {
   /* 0x0000 */ int a[1];
 };
 
 kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) {
-  uint const l1 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u);
+  uint const l1 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u);
-  uint const l2 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u);
+  uint const l2 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u);
   return;
 }
 

test/intrinsics/arrayLength/simple.wgsl.expected.msl

@@ -2,14 +2,14 @@
 
 using namespace metal;
 struct tint_symbol_1 {
-  /* 0x0000 */ uint buffer_size[1];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct S {
   /* 0x0000 */ int a[1];
 };
 
 kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) {
-  uint const l1 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u);
+  uint const l1 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u);
   return;
 }
 

test/intrinsics/arrayLength/via_let.wgsl.expected.msl

@@ -2,14 +2,14 @@
 
 using namespace metal;
 struct tint_symbol_1 {
-  /* 0x0000 */ uint buffer_size[1];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct S {
   /* 0x0000 */ int a[1];
 };
 
 kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) {
-  uint const l1 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u);
+  uint const l1 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u);
   return;
 }
 

test/intrinsics/arrayLength/via_let_complex.wgsl.expected.msl

@@ -2,14 +2,14 @@
 
 using namespace metal;
 struct tint_symbol_1 {
-  /* 0x0000 */ uint buffer_size[1];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct S {
   /* 0x0000 */ int a[1];
 };
 
 kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) {
-  uint const l1 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u);
+  uint const l1 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u);
   return;
 }
 

test/intrinsics/gen/arrayLength/1588cd.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 struct tint_symbol_2 {
-  /* 0x0000 */ uint buffer_size[2];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct SB_RO {
   /* 0x0000 */ int arg_0[1];
@@ -12,7 +12,7 @@ struct tint_symbol {
 };
 
 void arrayLength_1588cd(constant tint_symbol_2& tint_symbol_3) {
-  uint res = ((tint_symbol_3.buffer_size[1u] - 0u) / 4u);
+  uint res = ((tint_symbol_3.buffer_size[0u][1u] - 0u) / 4u);
 }
 
 vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {

test/intrinsics/gen/arrayLength/61b1c7.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 struct tint_symbol_2 {
-  /* 0x0000 */ uint buffer_size[1];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct SB_RW {
   /* 0x0000 */ int arg_0[1];
@@ -12,7 +12,7 @@ struct tint_symbol {
 };
 
 void arrayLength_61b1c7(constant tint_symbol_2& tint_symbol_3) {
-  uint res = ((tint_symbol_3.buffer_size[0u] - 0u) / 4u);
+  uint res = ((tint_symbol_3.buffer_size[0u][0u] - 0u) / 4u);
 }
 
 vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {

test/intrinsics/gen/arrayLength/a0f5ca.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 struct tint_symbol_2 {
-  /* 0x0000 */ uint buffer_size[2];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct SB_RO {
   /* 0x0000 */ float arg_0[1];
@@ -12,7 +12,7 @@ struct tint_symbol {
 };
 
 void arrayLength_a0f5ca(constant tint_symbol_2& tint_symbol_3) {
-  uint res = ((tint_symbol_3.buffer_size[1u] - 0u) / 4u);
+  uint res = ((tint_symbol_3.buffer_size[0u][1u] - 0u) / 4u);
 }
 
 vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {

test/intrinsics/gen/arrayLength/cdd123.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 struct tint_symbol_2 {
-  /* 0x0000 */ uint buffer_size[1];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct SB_RW {
   /* 0x0000 */ float arg_0[1];
@@ -12,7 +12,7 @@ struct tint_symbol {
 };
 
 void arrayLength_cdd123(constant tint_symbol_2& tint_symbol_3) {
-  uint res = ((tint_symbol_3.buffer_size[0u] - 0u) / 4u);
+  uint res = ((tint_symbol_3.buffer_size[0u][0u] - 0u) / 4u);
 }
 
 vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {

test/intrinsics/gen/arrayLength/cfca0a.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 struct tint_symbol_2 {
-  /* 0x0000 */ uint buffer_size[2];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct SB_RO {
   /* 0x0000 */ uint arg_0[1];
@@ -12,7 +12,7 @@ struct tint_symbol {
 };
 
 void arrayLength_cfca0a(constant tint_symbol_2& tint_symbol_3) {
-  uint res = ((tint_symbol_3.buffer_size[1u] - 0u) / 4u);
+  uint res = ((tint_symbol_3.buffer_size[0u][1u] - 0u) / 4u);
 }
 
 vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {

test/intrinsics/gen/arrayLength/eb510f.wgsl.expected.msl

@@ -2,7 +2,7 @@
 
 using namespace metal;
 struct tint_symbol_2 {
-  /* 0x0000 */ uint buffer_size[1];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct SB_RW {
   /* 0x0000 */ uint arg_0[1];
@@ -12,7 +12,7 @@ struct tint_symbol {
 };
 
 void arrayLength_eb510f(constant tint_symbol_2& tint_symbol_3) {
-  uint res = ((tint_symbol_3.buffer_size[0u] - 0u) / 4u);
+  uint res = ((tint_symbol_3.buffer_size[0u][0u] - 0u) / 4u);
 }
 
 vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {