From fd5829e5ead28012956b214940392407aa455530 Mon Sep 17 00:00:00 2001 From: Antonio Maiorano Date: Mon, 12 Jul 2021 15:25:49 +0000 Subject: [PATCH] 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 Commit-Queue: Antonio Maiorano Reviewed-by: Ben Clayton --- src/CMakeLists.txt | 1 + src/inspector/inspector_test.cc | 17 +- src/program_builder.h | 35 +- src/resolver/resolver.cc | 225 +++++++++- src/resolver/resolver.h | 9 + .../storage_class_layout_validation_test.cc | 390 ++++++++++++++++++ src/resolver/storage_class_validation_test.cc | 4 +- src/transform/array_length_from_uniform.cc | 28 +- .../array_length_from_uniform_test.cc | 71 +++- test/BUILD.gn | 1 + .../complex_via_let.wgsl.expected.msl | 4 +- .../arrayLength/deprecated.wgsl.expected.msl | 6 +- .../arrayLength/simple.wgsl.expected.msl | 4 +- .../arrayLength/via_let.wgsl.expected.msl | 4 +- .../via_let_complex.wgsl.expected.msl | 4 +- .../gen/arrayLength/1588cd.wgsl.expected.msl | 4 +- .../gen/arrayLength/61b1c7.wgsl.expected.msl | 4 +- .../gen/arrayLength/a0f5ca.wgsl.expected.msl | 4 +- .../gen/arrayLength/cdd123.wgsl.expected.msl | 4 +- .../gen/arrayLength/cfca0a.wgsl.expected.msl | 4 +- .../gen/arrayLength/eb510f.wgsl.expected.msl | 4 +- 21 files changed, 774 insertions(+), 53 deletions(-) create mode 100644 src/resolver/storage_class_layout_validation_test.cc diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index dd3f2d127f..c545a740ea 100644 --- a/src/CMakeLists.txt +++ b/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 diff --git a/src/inspector/inspector_test.cc b/src/inspector/inspector_test.cc index 2aa1c49d02..54078d52e6 100644 --- a/src/inspector/inspector_test.cc +++ b/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 - // usage. Once crbug.com/tint/628 is implemented, this will fail validation - // and will need to be fixed. - ast::Struct* foo_struct_type = - MakeUniformBufferType("foo_type", {ty.i32(), ty.array()}); + // Manually create uniform buffer to make sure it had a valid layout (array + // with elem stride of 16, and that is 16-byte aligned within the struct) + ast::Struct* foo_struct_type = Structure( + "foo_type", + {Member("0__i32", ty.i32()), + Member("b", ty.array(ty.u32(), 4, /*stride*/ 16), {MemberAlign(16)})}, + {create()}); + 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(20u, result[0].size_no_padding); + EXPECT_EQ(80u, result[0].size); + EXPECT_EQ(80u, result[0].size_no_padding); } TEST_F(InspectorGetStorageBufferResourceBindingsTest, Simple) { diff --git a/src/program_builder.h b/src/program_builder.h index f0b0e059c8..70b356d948 100644 --- a/src/program_builder.h +++ b/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 + template , Source>* = nullptr> ast::Variable* Global(NAME&& name, const ast::Type* type, OPTIONAL&&... optional) { @@ -1676,6 +1680,35 @@ class ProgramBuilder { return create(source_, val); } + /// Creates a ast::StructBlockDecoration + /// @returns the struct block decoration pointer + ast::StructBlockDecoration* StructBlock() { + return create(); + } + + /// Creates the ast::GroupDecoration + /// @param value group decoration index + /// @returns the group decoration pointer + ast::GroupDecoration* Group(uint32_t value) { + return create(value); + } + + /// Creates the ast::BindingDecoration + /// @param value the binding index + /// @returns the binding deocration pointer + ast::BindingDecoration* Binding(uint32_t value) { + return create(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 diff --git a/src/resolver/resolver.cc b/src/resolver/resolver.cc index fe559f36c0..3f1e6f3bce 100644 --- a/src/resolver/resolver.cc +++ b/src/resolver/resolver.cc @@ -15,6 +15,8 @@ #include "src/resolver/resolver.h" #include +#include +#include #include #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(); + }; + + auto is_uniform_struct = [sc](const sem::Type* ty) { + return sc == ast::StorageClass::kUniform && ty->Is(); + }; + + 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(::log10(last_member_struct_padding_offset)) + 1; + const auto size_w = static_cast(::log10(st->Size())) + 1; + const auto align_w = static_cast(::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()) { + 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(); + 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()) { + 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()) { + // 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()) { + 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; diff --git a/src/resolver/resolver.h b/src/resolver/resolver.h index 04b7326088..bc3f58fa45 100644 --- a/src/resolver/resolver.h +++ b/src/resolver/resolver.h @@ -16,9 +16,11 @@ #define SRC_RESOLVER_RESOLVER_H_ #include +#include #include #include #include +#include #include #include "src/intrinsic_table.h" @@ -122,6 +124,9 @@ class Resolver { const sem::Intrinsic* intrinsic; }; + std::set> + 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 diff --git a/src/resolver/storage_class_layout_validation_test.cc b/src/resolver/storage_class_layout_validation_test.cc new file mode 100644 index 0000000000..5339f13cf3 --- /dev/null +++ b/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 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 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 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 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; + // + // [[block]] + // struct Outer { + // scalar : f32; + // inner : Inner; + // }; + // + // [[group(0), binding(0)]] + // var 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; + // + // [[block]] + // struct Outer { + // scalar : f32; + // [[align(16)]] inner : Inner; + // }; + // + // [[group(0), binding(0)]] + // var 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 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 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; + // s : f32; + // }; + // [[group(0), binding(0)]] + // var 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; + // + // [[block]] + // struct Outer { + // inner : Inner; + // scalar : i32; + // }; + // + // [[group(0), binding(0)]] + // var 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; + // + // [[block]] + // struct Outer { + // inner : Inner; + // scalar : i32; + // }; + // + // [[group(0), binding(0)]] + // var 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 diff --git a/src/resolver/storage_class_validation_test.cc b/src/resolver/storage_class_validation_test.cc index c9a1ea5e7e..2d75c1f581 100644 --- a/src/resolver/storage_class_validation_test.cc +++ b/src/resolver/storage_class_validation_test.cc @@ -263,7 +263,7 @@ TEST_F(ResolverStorageClassValidationTest, UniformBufferNoError_Basic) { create(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(0), }); - ASSERT_TRUE(r()->Resolve()); + ASSERT_TRUE(r()->Resolve()) << r()->error(); } } // namespace diff --git a/src/transform/array_length_from_uniform.cc b/src/transform/array_length_from_uniform.cc index bb182c9380..35fbcb999f 100644 --- a/src/transform/array_length_from_uniform.cc +++ b/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, 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()}); 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` - // have been run before this transform. + // `&resource.array`, which requires that `InlinePointerLets` and + // `Simplify` have been run before this transform. auto* param = call_expr->params()[0]->As(); 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(); 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(); 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 diff --git a/src/transform/array_length_from_uniform_test.cc b/src/transform/array_length_from_uniform_test.cc index a173de8ab4..6ab39ee055 100644 --- a/src/transform/array_length_from_uniform_test.cc +++ b/src/transform/array_length_from_uniform_test.cc @@ -81,7 +81,7 @@ fn main() { auto* expect = R"( [[block]] struct tint_symbol { - buffer_size : array; + buffer_size : array, 1>; }; [[group(0), binding(30)]] var 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; + buffer_size : array, 1>; }; [[group(0), binding(30)]] var 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; }; - [[block]] struct SB2 { x : i32; arr2 : array>; }; +[[block]] +struct SB3 { + x : i32; + arr3 : array>; +}; +[[block]] +struct SB4 { + x : i32; + arr4 : array>; +}; +[[block]] +struct SB5 { + x : i32; + arr5 : array>; +}; [[group(0), binding(2)]] var sb1 : SB1; - [[group(1), binding(2)]] var sb2 : SB2; +[[group(2), binding(2)]] var sb3 : SB3; +[[group(3), binding(2)]] var sb4 : SB4; +[[group(4), binding(2)]] var 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; + buffer_size : array, 2>; }; [[group(0), binding(30)]] var tint_symbol_1 : tint_symbol; @@ -214,21 +233,51 @@ struct SB2 { arr2 : array>; }; +[[block]] +struct SB3 { + x : i32; + arr3 : array>; +}; + +[[block]] +struct SB4 { + x : i32; + arr4 : array>; +}; + +[[block]] +struct SB5 { + x : i32; + arr5 : array>; +}; + [[group(0), binding(2)]] var sb1 : SB1; [[group(1), binding(2)]] var sb2 : SB2; +[[group(2), binding(2)]] var sb3 : SB3; + +[[group(3), binding(2)]] var sb4 : SB4; + +[[group(4), binding(2)]] var sb5 : SB5; + [[stage(compute), workgroup_size(1)]] fn main() { - var len1 : u32 = ((tint_symbol_1.buffer_size[0u] - 4u) / 4u); - var len2 : u32 = ((tint_symbol_1.buffer_size[1u] - 16u) / 16u); - var x : u32 = (len1 + len2); + var len1 : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u); + var len2 : u32 = ((tint_symbol_1.buffer_size[0u][1u] - 16u) / 16u); + 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(std::move(cfg)); diff --git a/test/BUILD.gn b/test/BUILD.gn index 7439a5640c..2eebbffdf4 100644 --- a/test/BUILD.gn +++ b/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", diff --git a/test/intrinsics/arrayLength/complex_via_let.wgsl.expected.msl b/test/intrinsics/arrayLength/complex_via_let.wgsl.expected.msl index d690d2497e..9c7e44c11b 100644 --- a/test/intrinsics/arrayLength/complex_via_let.wgsl.expected.msl +++ b/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; } diff --git a/test/intrinsics/arrayLength/deprecated.wgsl.expected.msl b/test/intrinsics/arrayLength/deprecated.wgsl.expected.msl index c86fc59bd4..40f5a450e5 100644 --- a/test/intrinsics/arrayLength/deprecated.wgsl.expected.msl +++ b/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 l2 = ((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] - 0u) / 4u); return; } diff --git a/test/intrinsics/arrayLength/simple.wgsl.expected.msl b/test/intrinsics/arrayLength/simple.wgsl.expected.msl index d690d2497e..9c7e44c11b 100644 --- a/test/intrinsics/arrayLength/simple.wgsl.expected.msl +++ b/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; } diff --git a/test/intrinsics/arrayLength/via_let.wgsl.expected.msl b/test/intrinsics/arrayLength/via_let.wgsl.expected.msl index d690d2497e..9c7e44c11b 100644 --- a/test/intrinsics/arrayLength/via_let.wgsl.expected.msl +++ b/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; } diff --git a/test/intrinsics/arrayLength/via_let_complex.wgsl.expected.msl b/test/intrinsics/arrayLength/via_let_complex.wgsl.expected.msl index d690d2497e..9c7e44c11b 100644 --- a/test/intrinsics/arrayLength/via_let_complex.wgsl.expected.msl +++ b/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; } diff --git a/test/intrinsics/gen/arrayLength/1588cd.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/1588cd.wgsl.expected.msl index 4f53a604fc..3553097e11 100644 --- a/test/intrinsics/gen/arrayLength/1588cd.wgsl.expected.msl +++ b/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)]]) { diff --git a/test/intrinsics/gen/arrayLength/61b1c7.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/61b1c7.wgsl.expected.msl index 71c7a837af..97d28572b0 100644 --- a/test/intrinsics/gen/arrayLength/61b1c7.wgsl.expected.msl +++ b/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)]]) { diff --git a/test/intrinsics/gen/arrayLength/a0f5ca.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/a0f5ca.wgsl.expected.msl index ad28d92e06..ed949991fa 100644 --- a/test/intrinsics/gen/arrayLength/a0f5ca.wgsl.expected.msl +++ b/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)]]) { diff --git a/test/intrinsics/gen/arrayLength/cdd123.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/cdd123.wgsl.expected.msl index fdbb89de82..525920c8df 100644 --- a/test/intrinsics/gen/arrayLength/cdd123.wgsl.expected.msl +++ b/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)]]) { diff --git a/test/intrinsics/gen/arrayLength/cfca0a.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/cfca0a.wgsl.expected.msl index 14068e5729..97cbb4f133 100644 --- a/test/intrinsics/gen/arrayLength/cfca0a.wgsl.expected.msl +++ b/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)]]) { diff --git a/test/intrinsics/gen/arrayLength/eb510f.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/eb510f.wgsl.expected.msl index 04f85089e2..d345929025 100644 --- a/test/intrinsics/gen/arrayLength/eb510f.wgsl.expected.msl +++ b/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)]]) {