Revert "Validate storage class constraints"
This reverts commit fd5829e5ead28012956b214940392407aa455530.
Reason for revert: Temporarily reverting as this is preventing a tint->dawn roll, which is needed to fix the dawn->chrome roll.
Original change's description:
> 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>
TBR=bclayton@google.com,amaiorano@google.com,noreply+kokoro@google.com,tint-scoped@luci-project-accounts.iam.gserviceaccount.com
Change-Id: I8dbd0e46b3e3291ef08797a196d0d9abd2a78845
No-Presubmit: true
No-Tree-Checks: true
No-Try: true
Bug: tint:643
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/57704
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Ben Clayton <bclayton@google.com>
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index c545a74..dd3f2d1 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -637,7 +637,6 @@
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 54078d5..2aa1c49 100644
--- a/src/inspector/inspector_test.cc
+++ b/src/inspector/inspector_test.cc
@@ -2239,14 +2239,11 @@
}
TEST_F(InspectorGetUniformBufferResourceBindingsTest, ContainingArray) {
- // 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<ast::StructBlockDecoration>()});
-
+ // 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<u32, 4>()});
AddUniformBuffer("foo_ub", ty.Of(foo_struct_type), 0, 0);
MakeStructVariableReferenceBodyFunction("ub_func", "foo_ub", {{0, ty.i32()}});
@@ -2266,8 +2263,8 @@
result[0].resource_type);
EXPECT_EQ(0u, result[0].bind_group);
EXPECT_EQ(0u, result[0].binding);
- EXPECT_EQ(80u, result[0].size);
- EXPECT_EQ(80u, result[0].size_no_padding);
+ EXPECT_EQ(20u, result[0].size);
+ EXPECT_EQ(20u, result[0].size_no_padding);
}
TEST_F(InspectorGetStorageBufferResourceBindingsTest, Simple) {
diff --git a/src/program_builder.h b/src/program_builder.h
index 70b356d..f0b0e05 100644
--- a/src/program_builder.h
+++ b/src/program_builder.h
@@ -25,7 +25,6 @@
#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"
@@ -55,7 +54,6 @@
#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"
@@ -1381,9 +1379,7 @@
/// value.
/// @returns a new `ast::Variable`, which is automatically registered as a
/// global variable with the ast::Module.
- template <typename NAME,
- typename... OPTIONAL,
- traits::EnableIfIsNotType<traits::Decay<NAME>, Source>* = nullptr>
+ template <typename NAME, typename... OPTIONAL>
ast::Variable* Global(NAME&& name,
const ast::Type* type,
OPTIONAL&&... optional) {
@@ -1680,35 +1676,6 @@
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
diff --git a/src/resolver/resolver.cc b/src/resolver/resolver.cc
index 3f1e6f3..fe559f3 100644
--- a/src/resolver/resolver.cc
+++ b/src/resolver/resolver.cc
@@ -15,8 +15,6 @@
#include "src/resolver/resolver.h"
#include <algorithm>
-#include <cmath>
-#include <iomanip>
#include <utility>
#include "src/ast/alias.h"
@@ -687,228 +685,6 @@
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;
}
@@ -3860,6 +3636,7 @@
// 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 bc3f58f..04b7326 100644
--- a/src/resolver/resolver.h
+++ b/src/resolver/resolver.h
@@ -16,11 +16,9 @@
#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"
@@ -124,9 +122,6 @@
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 {
@@ -317,10 +312,6 @@
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
deleted file mode 100644
index 5339f13..0000000
--- a/src/resolver/storage_class_layout_validation_test.cc
+++ /dev/null
@@ -1,390 +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/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
diff --git a/src/resolver/storage_class_validation_test.cc b/src/resolver/storage_class_validation_test.cc
index 2d75c1f..c9a1ea5 100644
--- a/src/resolver/storage_class_validation_test.cc
+++ b/src/resolver/storage_class_validation_test.cc
@@ -263,7 +263,7 @@
create<ast::GroupDecoration>(0),
});
- ASSERT_TRUE(r()->Resolve()) << r()->error();
+ ASSERT_TRUE(r()->Resolve());
}
TEST_F(ResolverStorageClassValidationTest, UniformBufferNoError_Aliases) {
@@ -279,7 +279,7 @@
create<ast::GroupDecoration>(0),
});
- ASSERT_TRUE(r()->Resolve()) << r()->error();
+ ASSERT_TRUE(r()->Resolve());
}
} // namespace
diff --git a/src/transform/array_length_from_uniform.cc b/src/transform/array_length_from_uniform.cc
index 35fbcb9..bb182c9 100644
--- a/src/transform/array_length_from_uniform.cc
+++ b/src/transform/array_length_from_uniform.cc
@@ -67,16 +67,11 @@
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.vec4(ctx.dst->ty.u32()),
- (max_buffer_size_index / 4) + 1))},
-
+ ctx.dst->ty.array(ctx.dst->ty.u32(), max_buffer_size_index + 1))},
ast::DecorationList{ctx.dst->create<ast::StructBlockDecoration>()});
buffer_size_ubo = ctx.dst->Global(
ctx.dst->Sym(), ctx.dst->ty.Of(buffer_size_struct),
@@ -104,20 +99,18 @@
// 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<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();
@@ -125,8 +118,7 @@
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;
}
@@ -143,13 +135,9 @@
}
// Load the total storage buffer size from the UBO.
- uint32_t array_index = idx_itr->second / 4;
- auto* vec_expr = ctx.dst->IndexAccessor(
+ auto* total_storage_buffer_size = ctx.dst->IndexAccessor(
ctx.dst->MemberAccessor(get_ubo()->symbol(), kBufferSizeMemberName),
- array_index);
- uint32_t vec_index = idx_itr->second % 4;
- auto* total_storage_buffer_size =
- ctx.dst->IndexAccessor(vec_expr, vec_index);
+ idx_itr->second);
// 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 6ab39ee..a173de8 100644
--- a/src/transform/array_length_from_uniform_test.cc
+++ b/src/transform/array_length_from_uniform_test.cc
@@ -81,7 +81,7 @@
auto* expect = R"(
[[block]]
struct tint_symbol {
- buffer_size : array<vec4<u32>, 1>;
+ buffer_size : array<u32, 1>;
};
[[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
@@ -96,7 +96,7 @@
[[stage(compute), workgroup_size(1)]]
fn main() {
- var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
+ var len : u32 = ((tint_symbol_1.buffer_size[0u] - 4u) / 4u);
}
)";
@@ -134,7 +134,7 @@
auto* expect = R"(
[[block]]
struct tint_symbol {
- buffer_size : array<vec4<u32>, 1>;
+ buffer_size : array<u32, 1>;
};
[[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
@@ -150,7 +150,7 @@
[[stage(compute), workgroup_size(1)]]
fn main() {
- var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 8u) / 64u);
+ var len : u32 = ((tint_symbol_1.buffer_size[0u] - 8u) / 64u);
}
)";
@@ -175,48 +175,29 @@
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 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);
+ var x : u32 = (len1 + len2);
}
)";
auto* expect = R"(
[[block]]
struct tint_symbol {
- buffer_size : array<vec4<u32>, 2>;
+ buffer_size : array<u32, 2>;
};
[[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
@@ -233,51 +214,21 @@
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][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);
+ 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);
}
)";
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));
diff --git a/test/BUILD.gn b/test/BUILD.gn
index 2eebbff..7439a56 100644
--- a/test/BUILD.gn
+++ b/test/BUILD.gn
@@ -246,7 +246,6 @@
"../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 9c7e44c..d690d24 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 */ uint4 buffer_size[1];
+ /* 0x0000 */ uint 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] - 0u) / 4u);
+ uint const l1 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u);
return;
}
diff --git a/test/intrinsics/arrayLength/deprecated.wgsl.expected.msl b/test/intrinsics/arrayLength/deprecated.wgsl.expected.msl
index 40f5a45..c86fc59 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 */ uint4 buffer_size[1];
+ /* 0x0000 */ uint 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] - 0u) / 4u);
- uint const l2 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u);
+ uint const l1 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u);
+ uint const l2 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u);
return;
}
diff --git a/test/intrinsics/arrayLength/simple.wgsl.expected.msl b/test/intrinsics/arrayLength/simple.wgsl.expected.msl
index 9c7e44c..d690d24 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 */ uint4 buffer_size[1];
+ /* 0x0000 */ uint 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] - 0u) / 4u);
+ uint const l1 = ((tint_symbol_2.buffer_size[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 9c7e44c..d690d24 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 */ uint4 buffer_size[1];
+ /* 0x0000 */ uint 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] - 0u) / 4u);
+ uint const l1 = ((tint_symbol_2.buffer_size[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 9c7e44c..d690d24 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 */ uint4 buffer_size[1];
+ /* 0x0000 */ uint 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] - 0u) / 4u);
+ uint const l1 = ((tint_symbol_2.buffer_size[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 3553097..4f53a60 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 */ uint4 buffer_size[1];
+ /* 0x0000 */ uint buffer_size[2];
};
struct SB_RO {
/* 0x0000 */ int arg_0[1];
@@ -12,7 +12,7 @@
};
void arrayLength_1588cd(constant tint_symbol_2& tint_symbol_3) {
- uint res = ((tint_symbol_3.buffer_size[0u][1u] - 0u) / 4u);
+ uint res = ((tint_symbol_3.buffer_size[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 97d2857..71c7a83 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 */ uint4 buffer_size[1];
+ /* 0x0000 */ uint buffer_size[1];
};
struct SB_RW {
/* 0x0000 */ int arg_0[1];
@@ -12,7 +12,7 @@
};
void arrayLength_61b1c7(constant tint_symbol_2& tint_symbol_3) {
- uint res = ((tint_symbol_3.buffer_size[0u][0u] - 0u) / 4u);
+ uint res = ((tint_symbol_3.buffer_size[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 ed94999..ad28d92 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 */ uint4 buffer_size[1];
+ /* 0x0000 */ uint buffer_size[2];
};
struct SB_RO {
/* 0x0000 */ float arg_0[1];
@@ -12,7 +12,7 @@
};
void arrayLength_a0f5ca(constant tint_symbol_2& tint_symbol_3) {
- uint res = ((tint_symbol_3.buffer_size[0u][1u] - 0u) / 4u);
+ uint res = ((tint_symbol_3.buffer_size[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 525920c..fdbb89d 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 */ uint4 buffer_size[1];
+ /* 0x0000 */ uint buffer_size[1];
};
struct SB_RW {
/* 0x0000 */ float arg_0[1];
@@ -12,7 +12,7 @@
};
void arrayLength_cdd123(constant tint_symbol_2& tint_symbol_3) {
- uint res = ((tint_symbol_3.buffer_size[0u][0u] - 0u) / 4u);
+ uint res = ((tint_symbol_3.buffer_size[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 97cbb4f..14068e5 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 */ uint4 buffer_size[1];
+ /* 0x0000 */ uint buffer_size[2];
};
struct SB_RO {
/* 0x0000 */ uint arg_0[1];
@@ -12,7 +12,7 @@
};
void arrayLength_cfca0a(constant tint_symbol_2& tint_symbol_3) {
- uint res = ((tint_symbol_3.buffer_size[0u][1u] - 0u) / 4u);
+ uint res = ((tint_symbol_3.buffer_size[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 d345929..04f8508 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 */ uint4 buffer_size[1];
+ /* 0x0000 */ uint buffer_size[1];
};
struct SB_RW {
/* 0x0000 */ uint arg_0[1];
@@ -12,7 +12,7 @@
};
void arrayLength_eb510f(constant tint_symbol_2& tint_symbol_3) {
- uint res = ((tint_symbol_3.buffer_size[0u][0u] - 0u) / 4u);
+ uint res = ((tint_symbol_3.buffer_size[0u] - 0u) / 4u);
}
vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
