encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

Revert "Validate storage class constraints" 

This reverts commit fd5829e5ea.

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>
This commit is contained in:
Ben Clayton 2021-07-12 20:15:43 +00:00 committed by Tint LUCI CQ
parent 66a88a2643
commit aa48b1ad8d
21 changed files with 53 additions and 774 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
src/CMakeLists.txt
View File

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

17
src/inspector/inspector_test.cc
View File

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

35
src/program_builder.h
View File

@ -25,7 +25,6 @@
#include "src/ast/assignment_statement.h" #include "src/ast/assignment_statement.h"
#include "src/ast/atomic.h" #include "src/ast/atomic.h"
#include "src/ast/binary_expression.h" #include "src/ast/binary_expression.h"
#include "src/ast/binding_decoration.h"
#include "src/ast/bitcast_expression.h" #include "src/ast/bitcast_expression.h"
#include "src/ast/bool.h" #include "src/ast/bool.h"
#include "src/ast/bool_literal.h" #include "src/ast/bool_literal.h"
@ -55,7 +54,6 @@
#include "src/ast/stage_decoration.h" #include "src/ast/stage_decoration.h"
#include "src/ast/storage_texture.h" #include "src/ast/storage_texture.h"
#include "src/ast/stride_decoration.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_align_decoration.h"
#include "src/ast/struct_member_offset_decoration.h" #include "src/ast/struct_member_offset_decoration.h"
#include "src/ast/struct_member_size_decoration.h" #include "src/ast/struct_member_size_decoration.h"
@ -1381,9 +1379,7 @@ class ProgramBuilder {
/// value. /// value.
/// @returns a new `ast::Variable`, which is automatically registered as a /// @returns a new `ast::Variable`, which is automatically registered as a
/// global variable with the ast::Module. /// global variable with the ast::Module.
template <typename NAME, template <typename NAME, typename... OPTIONAL>
typename... OPTIONAL,
traits::EnableIfIsNotType<traits::Decay<NAME>, Source>* = nullptr>
ast::Variable* Global(NAME&& name, ast::Variable* Global(NAME&& name,
const ast::Type* type, const ast::Type* type,
OPTIONAL&&... optional) { OPTIONAL&&... optional) {
@ -1680,35 +1676,6 @@ class ProgramBuilder {
return create<ast::StructMemberAlignDecoration>(source_, val); 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. /// Creates an ast::Function and registers it with the ast::Module.
/// @param source the source information /// @param source the source information
/// @param name the function name /// @param name the function name

225
src/resolver/resolver.cc
View File

@ -15,8 +15,6 @@
#include "src/resolver/resolver.h" #include "src/resolver/resolver.h"
#include <algorithm> #include <algorithm>
#include <cmath>
#include <iomanip>
#include <utility> #include <utility>
#include "src/ast/alias.h" #include "src/ast/alias.h"
@ -687,228 +685,6 @@ bool Resolver::GlobalVariable(ast::Variable* var) {
return false; 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; return true;
} }
@ -3860,6 +3636,7 @@ sem::Struct* Resolver::Structure(const ast::Struct* str) {
// Validation of storage-class rules requires analysing the actual variable // Validation of storage-class rules requires analysing the actual variable
// usage of the structure, and so is performed as part of the variable // usage of the structure, and so is performed as part of the variable
// validation. // validation.
// TODO(crbug.com/tint/628): Actually implement storage-class validation.
uint32_t struct_size = 0; uint32_t struct_size = 0;
uint32_t struct_align = 1; uint32_t struct_align = 1;

9
src/resolver/resolver.h
View File

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

390
src/resolver/storage_class_layout_validation_test.cc
View File

@ -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

4
src/resolver/storage_class_validation_test.cc
View File

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

28
src/transform/array_length_from_uniform.cc
View File

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

71
src/transform/array_length_from_uniform_test.cc
View File

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

1
test/BUILD.gn
View File

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

4
test/intrinsics/arrayLength/complex_via_let.wgsl.expected.msl
View File

@ -2,14 +2,14 @@
using namespace metal; using namespace metal;
struct tint_symbol_1 { struct tint_symbol_1 {
/* 0x0000 */ uint4 buffer_size[1]; /* 0x0000 */ uint buffer_size[1];
}; };
struct S { struct S {
/* 0x0000 */ int a[1]; /* 0x0000 */ int a[1];
}; };
kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) { 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; return;
} }

6
test/intrinsics/arrayLength/deprecated.wgsl.expected.msl
View File

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

4
test/intrinsics/arrayLength/simple.wgsl.expected.msl
View File

@ -2,14 +2,14 @@
using namespace metal; using namespace metal;
struct tint_symbol_1 { struct tint_symbol_1 {
/* 0x0000 */ uint4 buffer_size[1]; /* 0x0000 */ uint buffer_size[1];
}; };
struct S { struct S {
/* 0x0000 */ int a[1]; /* 0x0000 */ int a[1];
}; };
kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) { 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; return;
} }

4
test/intrinsics/arrayLength/via_let.wgsl.expected.msl
View File

@ -2,14 +2,14 @@
using namespace metal; using namespace metal;
struct tint_symbol_1 { struct tint_symbol_1 {
/* 0x0000 */ uint4 buffer_size[1]; /* 0x0000 */ uint buffer_size[1];
}; };
struct S { struct S {
/* 0x0000 */ int a[1]; /* 0x0000 */ int a[1];
}; };
kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) { 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; return;
} }

4
test/intrinsics/arrayLength/via_let_complex.wgsl.expected.msl
View File

@ -2,14 +2,14 @@
using namespace metal; using namespace metal;
struct tint_symbol_1 { struct tint_symbol_1 {
/* 0x0000 */ uint4 buffer_size[1]; /* 0x0000 */ uint buffer_size[1];
}; };
struct S { struct S {
/* 0x0000 */ int a[1]; /* 0x0000 */ int a[1];
}; };
kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) { 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; return;
} }

4
test/intrinsics/gen/arrayLength/1588cd.wgsl.expected.msl
View File

@ -2,7 +2,7 @@
using namespace metal; using namespace metal;
struct tint_symbol_2 { struct tint_symbol_2 {
/* 0x0000 */ uint4 buffer_size[1]; /* 0x0000 */ uint buffer_size[2];
}; };
struct SB_RO { struct SB_RO {
/* 0x0000 */ int arg_0[1]; /* 0x0000 */ int arg_0[1];
@ -12,7 +12,7 @@ struct tint_symbol {
}; };
void arrayLength_1588cd(constant tint_symbol_2& tint_symbol_3) { 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)]]) { vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {

4
test/intrinsics/gen/arrayLength/61b1c7.wgsl.expected.msl
View File

@ -2,7 +2,7 @@
using namespace metal; using namespace metal;
struct tint_symbol_2 { struct tint_symbol_2 {
/* 0x0000 */ uint4 buffer_size[1]; /* 0x0000 */ uint buffer_size[1];
}; };
struct SB_RW { struct SB_RW {
/* 0x0000 */ int arg_0[1]; /* 0x0000 */ int arg_0[1];
@ -12,7 +12,7 @@ struct tint_symbol {
}; };
void arrayLength_61b1c7(constant tint_symbol_2& tint_symbol_3) { 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)]]) { vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {

4
test/intrinsics/gen/arrayLength/a0f5ca.wgsl.expected.msl
View File

@ -2,7 +2,7 @@
using namespace metal; using namespace metal;
struct tint_symbol_2 { struct tint_symbol_2 {
/* 0x0000 */ uint4 buffer_size[1]; /* 0x0000 */ uint buffer_size[2];
}; };
struct SB_RO { struct SB_RO {
/* 0x0000 */ float arg_0[1]; /* 0x0000 */ float arg_0[1];
@ -12,7 +12,7 @@ struct tint_symbol {
}; };
void arrayLength_a0f5ca(constant tint_symbol_2& tint_symbol_3) { 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)]]) { vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {

4
test/intrinsics/gen/arrayLength/cdd123.wgsl.expected.msl
View File

@ -2,7 +2,7 @@
using namespace metal; using namespace metal;
struct tint_symbol_2 { struct tint_symbol_2 {
/* 0x0000 */ uint4 buffer_size[1]; /* 0x0000 */ uint buffer_size[1];
}; };
struct SB_RW { struct SB_RW {
/* 0x0000 */ float arg_0[1]; /* 0x0000 */ float arg_0[1];
@ -12,7 +12,7 @@ struct tint_symbol {
}; };
void arrayLength_cdd123(constant tint_symbol_2& tint_symbol_3) { 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)]]) { vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {

4
test/intrinsics/gen/arrayLength/cfca0a.wgsl.expected.msl
View File

@ -2,7 +2,7 @@
using namespace metal; using namespace metal;
struct tint_symbol_2 { struct tint_symbol_2 {
/* 0x0000 */ uint4 buffer_size[1]; /* 0x0000 */ uint buffer_size[2];
}; };
struct SB_RO { struct SB_RO {
/* 0x0000 */ uint arg_0[1]; /* 0x0000 */ uint arg_0[1];
@ -12,7 +12,7 @@ struct tint_symbol {
}; };
void arrayLength_cfca0a(constant tint_symbol_2& tint_symbol_3) { 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)]]) { vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {

4
test/intrinsics/gen/arrayLength/eb510f.wgsl.expected.msl
View File

@ -2,7 +2,7 @@
using namespace metal; using namespace metal;
struct tint_symbol_2 { struct tint_symbol_2 {
/* 0x0000 */ uint4 buffer_size[1]; /* 0x0000 */ uint buffer_size[1];
}; };
struct SB_RW { struct SB_RW {
/* 0x0000 */ uint arg_0[1]; /* 0x0000 */ uint arg_0[1];
@ -12,7 +12,7 @@ struct tint_symbol {
}; };
void arrayLength_eb510f(constant tint_symbol_2& tint_symbol_3) { 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)]]) { vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {