transform: Handle arrayLength for non-struct buffers
These two transforms previously assumed that the argument to
arrayLength had the form `&struct_var.array_member`. We now also need
to handle the case where it is just `&array_var`.
Bug: tint:1372
Change-Id: I173a84bd32c324445573a295b281a51e291c2ae2
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/76163
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/transform/array_length_from_uniform.cc b/src/transform/array_length_from_uniform.cc
index 5084a54..32289fe 100644
--- a/src/transform/array_length_from_uniform.cc
+++ b/src/transform/array_length_from_uniform.cc
@@ -59,30 +59,26 @@
}
// Get the storage buffer that contains the runtime array.
- // We assume that the argument to `arrayLength` has the form
- // `&resource.array`, which requires that `SimplifyPointers` have been run
- // before this transform.
+ // Since we require SimplifyPointers, we can assume that the arrayLength()
+ // call has one of two forms:
+ // arrayLength(&struct_var.array_member)
+ // arrayLength(&array_var)
auto* param = call_expr->args[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 &array_var or "
+ "&struct_var.array_member";
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";
- break;
+ auto* storage_buffer_expr = param->expr;
+ if (auto* accessor = param->expr->As<ast::MemberAccessorExpression>()) {
+ storage_buffer_expr = accessor->structure;
}
- auto* storage_buffer_expr = accessor->structure;
- auto* storage_buffer_sem =
- sem.Get(storage_buffer_expr)->As<sem::VariableUser>();
+ auto* storage_buffer_sem = sem.Get<sem::VariableUser>(storage_buffer_expr);
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 &array_var or "
+ "&struct_var.array_member";
break;
}
@@ -183,14 +179,25 @@
// total_storage_buffer_size - array_offset
// array_length = ----------------------------------------
// array_stride
- auto* storage_buffer_type =
- storage_buffer_sem->Type()->UnwrapRef()->As<sem::Struct>();
- auto* array_member_sem = storage_buffer_type->Members().back();
- uint32_t array_offset = array_member_sem->Offset();
- uint32_t array_stride = array_member_sem->Size();
- auto* array_length =
- ctx.dst->Div(ctx.dst->Sub(total_storage_buffer_size, array_offset),
- array_stride);
+ const ast::Expression* total_size = total_storage_buffer_size;
+ auto* storage_buffer_type = storage_buffer_sem->Type()->UnwrapRef();
+ const sem::Array* array_type = nullptr;
+ if (auto* str = storage_buffer_type->As<sem::Struct>()) {
+ // The variable is a struct, so subtract the byte offset of the array
+ // member.
+ auto* array_member_sem = str->Members().back();
+ array_type = array_member_sem->Type()->As<sem::Array>();
+ total_size = ctx.dst->Sub(total_storage_buffer_size,
+ array_member_sem->Offset());
+ } else if (auto* arr = storage_buffer_type->As<sem::Array>()) {
+ array_type = arr;
+ } else {
+ TINT_ICE(Transform, ctx.dst->Diagnostics())
+ << "expected form of arrayLength argument to be &array_var or "
+ "&struct_var.array_member";
+ return;
+ }
+ auto* array_length = ctx.dst->Div(total_size, array_type->Stride());
ctx.Replace(call_expr, array_length);
});
diff --git a/src/transform/array_length_from_uniform_test.cc b/src/transform/array_length_from_uniform_test.cc
index 59f5a4d..77fb0c3 100644
--- a/src/transform/array_length_from_uniform_test.cc
+++ b/src/transform/array_length_from_uniform_test.cc
@@ -52,6 +52,44 @@
TEST_F(ArrayLengthFromUniformTest, Basic) {
auto* src = R"(
+[[group(0), binding(0)]] var<storage, read> sb : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ var len : u32 = arrayLength(&sb);
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol {
+ buffer_size : array<vec4<u32>, 1u>;
+};
+
+[[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
+
+[[group(0), binding(0)]] var<storage, read> sb : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ var len : u32 = (tint_symbol_1.buffer_size[0u][0u] / 4u);
+}
+)";
+
+ ArrayLengthFromUniform::Config cfg({0, 30u});
+ cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 0}, 0);
+
+ DataMap data;
+ data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
+
+ auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+ EXPECT_EQ(std::unordered_set<uint32_t>({0}),
+ got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
+}
+
+TEST_F(ArrayLengthFromUniformTest, BasicInStruct) {
+ auto* src = R"(
struct SB {
x : i32;
arr : array<i32>;
@@ -100,6 +138,44 @@
TEST_F(ArrayLengthFromUniformTest, WithStride) {
auto* src = R"(
+[[group(0), binding(0)]] var<storage, read> sb : [[stride(64)]] array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ var len : u32 = arrayLength(&sb);
+}
+)";
+
+ auto* expect = R"(
+struct tint_symbol {
+ buffer_size : array<vec4<u32>, 1u>;
+};
+
+[[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
+
+[[group(0), binding(0)]] var<storage, read> sb : [[stride(64)]] array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ var len : u32 = (tint_symbol_1.buffer_size[0u][0u] / 64u);
+}
+)";
+
+ ArrayLengthFromUniform::Config cfg({0, 30u});
+ cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 0}, 0);
+
+ DataMap data;
+ data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
+
+ auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
+
+ EXPECT_EQ(expect, str(got));
+ EXPECT_EQ(std::unordered_set<uint32_t>({0}),
+ got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
+}
+
+TEST_F(ArrayLengthFromUniformTest, WithStride_InStruct) {
+ auto* src = R"(
struct SB {
x : i32;
y : f32;
@@ -158,32 +234,24 @@
x : i32;
arr2 : array<vec4<f32>>;
};
-struct SB3 {
- x : i32;
- arr3 : array<vec4<f32>>;
-};
struct SB4 {
x : i32;
arr4 : array<vec4<f32>>;
};
-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(2), binding(2)]] var<storage, read> sb3 : array<vec4<f32>>;
[[group(3), binding(2)]] var<storage, read> sb4 : SB4;
-[[group(4), binding(2)]] var<storage, read> sb5 : SB5;
+[[group(4), binding(2)]] var<storage, read> sb5 : array<vec4<f32>>;
[[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 len3 : u32 = arrayLength(&sb3);
var len4 : u32 = arrayLength(&(sb4.arr4));
- var len5 : u32 = arrayLength(&(sb5.arr5));
+ var len5 : u32 = arrayLength(&sb5);
var x : u32 = (len1 + len2 + len3 + len4 + len5);
}
)";
@@ -205,38 +273,28 @@
arr2 : array<vec4<f32>>;
};
-struct SB3 {
- x : i32;
- arr3 : array<vec4<f32>>;
-};
-
struct SB4 {
x : i32;
arr4 : array<vec4<f32>>;
};
-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(2), binding(2)]] var<storage, read> sb3 : array<vec4<f32>>;
[[group(3), binding(2)]] var<storage, read> sb4 : SB4;
-[[group(4), binding(2)]] var<storage, read> sb5 : SB5;
+[[group(4), binding(2)]] var<storage, read> sb5 : array<vec4<f32>>;
[[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 len3 : u32 = (tint_symbol_1.buffer_size[0u][2u] / 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 len5 : u32 = (tint_symbol_1.buffer_size[1u][0u] / 16u);
var x : u32 = ((((len1 + len2) + len3) + len4) + len5);
}
)";
@@ -268,29 +326,21 @@
x : i32;
arr2 : array<vec4<f32>>;
};
-struct SB3 {
- x : i32;
- arr3 : array<vec4<f32>>;
-};
struct SB4 {
x : i32;
arr4 : array<vec4<f32>>;
};
-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(2), binding(2)]] var<storage, read> sb3 : array<vec4<f32>>;
[[group(3), binding(2)]] var<storage, read> sb4 : SB4;
-[[group(4), binding(2)]] var<storage, read> sb5 : SB5;
+[[group(4), binding(2)]] var<storage, read> sb5 : array<vec4<f32>>;
[[stage(compute), workgroup_size(1)]]
fn main() {
var len1 : u32 = arrayLength(&(sb1.arr1));
- var len3 : u32 = arrayLength(&(sb3.arr3));
+ var len3 : u32 = arrayLength(&sb3);
var x : u32 = (len1 + len3);
}
)";
@@ -312,35 +362,25 @@
arr2 : array<vec4<f32>>;
};
-struct SB3 {
- x : i32;
- arr3 : array<vec4<f32>>;
-};
-
struct SB4 {
x : i32;
arr4 : array<vec4<f32>>;
};
-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(2), binding(2)]] var<storage, read> sb3 : array<vec4<f32>>;
[[group(3), binding(2)]] var<storage, read> sb4 : SB4;
-[[group(4), binding(2)]] var<storage, read> sb5 : SB5;
+[[group(4), binding(2)]] var<storage, read> sb5 : array<vec4<f32>>;
[[stage(compute), workgroup_size(1)]]
fn main() {
var len1 : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
- var len3 : u32 = ((tint_symbol_1.buffer_size[0u][2u] - 16u) / 16u);
+ var len3 : u32 = (tint_symbol_1.buffer_size[0u][2u] / 16u);
var x : u32 = (len1 + len3);
}
)";
diff --git a/src/transform/calculate_array_length.cc b/src/transform/calculate_array_length.cc
index e06b959..6c6c6ff 100644
--- a/src/transform/calculate_array_length.cc
+++ b/src/transform/calculate_array_length.cc
@@ -42,7 +42,7 @@
/// It is used as a key by the array_length_by_usage map.
struct ArrayUsage {
ast::BlockStatement const* const block;
- sem::Node const* const buffer;
+ sem::Variable const* const buffer;
bool operator==(const ArrayUsage& rhs) const {
return block == rhs.block && buffer == rhs.buffer;
}
@@ -80,12 +80,11 @@
// get_buffer_size_intrinsic() emits the function decorated with
// BufferSizeIntrinsic that is transformed by the HLSL writer into a call to
// [RW]ByteAddressBuffer.GetDimensions().
- std::unordered_map<const sem::Struct*, Symbol> buffer_size_intrinsics;
- auto get_buffer_size_intrinsic = [&](const sem::Struct* buffer_type) {
+ std::unordered_map<const sem::Type*, Symbol> buffer_size_intrinsics;
+ auto get_buffer_size_intrinsic = [&](const sem::Type* buffer_type) {
return utils::GetOrCreate(buffer_size_intrinsics, buffer_type, [&] {
auto name = ctx.dst->Sym();
- auto* buffer_typename =
- ctx.dst->ty.type_name(ctx.Clone(buffer_type->Declaration()->name));
+ auto* type = CreateASTTypeFor(ctx, buffer_type);
auto* disable_validation = ctx.dst->Disable(
ast::DisabledValidation::kIgnoreConstructibleFunctionParameter);
auto* func = ctx.dst->create<ast::Function>(
@@ -95,7 +94,7 @@
// in order for HLSL to emit this as a ByteAddressBuffer.
ctx.dst->create<ast::Variable>(
ctx.dst->Sym("buffer"), ast::StorageClass::kStorage,
- ast::Access::kUndefined, buffer_typename, true, nullptr,
+ ast::Access::kUndefined, type, true, nullptr,
ast::DecorationList{disable_validation}),
ctx.dst->Param("result",
ctx.dst->ty.pointer(ctx.dst->ty.u32(),
@@ -106,8 +105,12 @@
ctx.dst->ASTNodes().Create<BufferSizeIntrinsic>(ctx.dst->ID()),
},
ast::DecorationList{});
- ctx.InsertAfter(ctx.src->AST().GlobalDeclarations(),
- buffer_type->Declaration(), func);
+ if (auto* str = buffer_type->As<sem::Struct>()) {
+ ctx.InsertAfter(ctx.src->AST().GlobalDeclarations(), str->Declaration(),
+ func);
+ } else {
+ ctx.InsertFront(ctx.src->AST().GlobalDeclarations(), func);
+ }
return name;
});
};
@@ -123,71 +126,47 @@
if (intrinsic->Type() == sem::IntrinsicType::kArrayLength) {
// We're dealing with an arrayLength() call
- // https://gpuweb.github.io/gpuweb/wgsl/#array-types states:
- //
- // * The last member of the structure type defining the store type for
- // a variable in the storage storage class may be a runtime-sized
- // array.
- // * A runtime-sized array must not be used as the store type or
- // contained within a store type in any other cases.
- // * An expression must not evaluate to a runtime-sized array type.
- //
- // We can assume that the arrayLength() call has a single argument of
- // the form: arrayLength(&X.Y) where X is an expression that resolves
- // to the storage buffer structure, and Y is the runtime sized array.
+ // A runtime-sized array can only appear as the store type of a
+ // variable, or the last element of a structure (which cannot itself
+ // be nested). Given that we require SimplifyPointers, we can assume
+ // that the arrayLength() call has one of two forms:
+ // arrayLength(&struct_var.array_member)
+ // arrayLength(&array_var)
auto* arg = call_expr->args[0];
auto* address_of = arg->As<ast::UnaryOpExpression>();
if (!address_of || address_of->op != ast::UnaryOp::kAddressOf) {
TINT_ICE(Transform, ctx.dst->Diagnostics())
- << "arrayLength() expected pointer to member access, got "
- << address_of->TypeInfo().name;
+ << "arrayLength() expected address-of, got "
+ << arg->TypeInfo().name;
}
- auto* array_expr = address_of->expr;
-
- auto* accessor = array_expr->As<ast::MemberAccessorExpression>();
- if (!accessor) {
+ auto* storage_buffer_expr = address_of->expr;
+ if (auto* accessor =
+ storage_buffer_expr->As<ast::MemberAccessorExpression>()) {
+ storage_buffer_expr = accessor->structure;
+ }
+ auto* storage_buffer_sem =
+ sem.Get<sem::VariableUser>(storage_buffer_expr);
+ if (!storage_buffer_sem) {
TINT_ICE(Transform, ctx.dst->Diagnostics())
- << "arrayLength() expected pointer to member access, got "
- "pointer to "
- << array_expr->TypeInfo().name;
+ << "expected form of arrayLength argument to be &array_var or "
+ "&struct_var.array_member";
break;
}
- auto* storage_buffer_expr = accessor->structure;
- auto* storage_buffer_sem = sem.Get(storage_buffer_expr);
- auto* storage_buffer_type =
- storage_buffer_sem->Type()->UnwrapRef()->As<sem::Struct>();
+ auto* storage_buffer_var = storage_buffer_sem->Variable();
+ auto* storage_buffer_type = storage_buffer_sem->Type()->UnwrapRef();
// Generate BufferSizeIntrinsic for this storage type if we haven't
// already
auto buffer_size = get_buffer_size_intrinsic(storage_buffer_type);
- if (!storage_buffer_type) {
- TINT_ICE(Transform, ctx.dst->Diagnostics())
- << "arrayLength(X.Y) expected X to be sem::Struct, got "
- << storage_buffer_type->FriendlyName(ctx.src->Symbols());
- break;
- }
-
// Find the current statement block
auto* block = call->Stmt()->Block()->Declaration();
- // If the storage_buffer_expr is resolves to a variable (typically
- // true) then key the array_length from the variable. If not, key off
- // the expression semantic node, which will be unique per call to
- // arrayLength().
- const sem::Node* storage_buffer_usage = storage_buffer_sem;
- if (auto* user = storage_buffer_sem->As<sem::VariableUser>()) {
- storage_buffer_usage = user->Variable();
- }
-
auto array_length = utils::GetOrCreate(
- array_length_by_usage, {block, storage_buffer_usage}, [&] {
+ array_length_by_usage, {block, storage_buffer_var}, [&] {
// First time this array length is used for this block.
// Let's calculate it.
- // Semantic info for the runtime array structure member
- auto* array_member_sem = storage_buffer_type->Members().back();
-
// Construct the variable that'll hold the result of
// RWByteAddressBuffer.GetDimensions()
auto* buffer_size_result = ctx.dst->Decl(
@@ -208,14 +187,28 @@
// array_length = ----------------------------------------
// array_stride
auto name = ctx.dst->Sym();
- uint32_t array_offset = array_member_sem->Offset();
- uint32_t array_stride = array_member_sem->Size();
- auto* array_length_var = ctx.dst->Decl(ctx.dst->Const(
- name, ctx.dst->ty.u32(),
- ctx.dst->Div(
- ctx.dst->Sub(buffer_size_result->variable->symbol,
- array_offset),
- array_stride)));
+ const ast::Expression* total_size =
+ ctx.dst->Expr(buffer_size_result->variable);
+ const sem::Array* array_type = nullptr;
+ if (auto* str = storage_buffer_type->As<sem::Struct>()) {
+ // The variable is a struct, so subtract the byte offset of
+ // the array member.
+ auto* array_member_sem = str->Members().back();
+ array_type = array_member_sem->Type()->As<sem::Array>();
+ total_size =
+ ctx.dst->Sub(total_size, array_member_sem->Offset());
+ } else if (auto* arr = storage_buffer_type->As<sem::Array>()) {
+ array_type = arr;
+ } else {
+ TINT_ICE(Transform, ctx.dst->Diagnostics())
+ << "expected form of arrayLength argument to be "
+ "&array_var or &struct_var.array_member";
+ return name;
+ }
+ uint32_t array_stride = array_type->Size();
+ auto* array_length_var = ctx.dst->Decl(
+ ctx.dst->Const(name, ctx.dst->ty.u32(),
+ ctx.dst->Div(total_size, array_stride)));
// Insert the array length calculations at the top of the block
ctx.InsertBefore(block->statements, block->statements[0],
diff --git a/src/transform/calculate_array_length_test.cc b/src/transform/calculate_array_length_test.cc
index 4522768..4b42bf9 100644
--- a/src/transform/calculate_array_length_test.cc
+++ b/src/transform/calculate_array_length_test.cc
@@ -38,6 +38,36 @@
TEST_F(CalculateArrayLengthTest, Basic) {
auto* src = R"(
+[[group(0), binding(0)]] var<storage, read> sb : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ var len : u32 = arrayLength(&sb);
+}
+)";
+
+ auto* expect = R"(
+[[internal(intrinsic_buffer_size)]]
+fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : array<i32>, result : ptr<function, u32>)
+
+[[group(0), binding(0)]] var<storage, read> sb : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ var tint_symbol_1 : u32 = 0u;
+ tint_symbol(sb, &(tint_symbol_1));
+ let tint_symbol_2 : u32 = (tint_symbol_1 / 4u);
+ var len : u32 = tint_symbol_2;
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(CalculateArrayLengthTest, BasicInStruct) {
+ auto* src = R"(
struct SB {
x : i32;
arr : array<i32>;
@@ -78,6 +108,40 @@
TEST_F(CalculateArrayLengthTest, InSameBlock) {
auto* src = R"(
+[[group(0), binding(0)]] var<storage, read> sb : array<i32>;;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ var a : u32 = arrayLength(&sb);
+ var b : u32 = arrayLength(&sb);
+ var c : u32 = arrayLength(&sb);
+}
+)";
+
+ auto* expect = R"(
+[[internal(intrinsic_buffer_size)]]
+fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : array<i32>, result : ptr<function, u32>)
+
+[[group(0), binding(0)]] var<storage, read> sb : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ var tint_symbol_1 : u32 = 0u;
+ tint_symbol(sb, &(tint_symbol_1));
+ let tint_symbol_2 : u32 = (tint_symbol_1 / 4u);
+ var a : u32 = tint_symbol_2;
+ var b : u32 = tint_symbol_2;
+ var c : u32 = tint_symbol_2;
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(CalculateArrayLengthTest, InSameBlock_Struct) {
+ auto* src = R"(
struct SB {
x : i32;
arr : array<i32>;
@@ -122,6 +186,36 @@
TEST_F(CalculateArrayLengthTest, WithStride) {
auto* src = R"(
+[[group(0), binding(0)]] var<storage, read> sb : [[stride(64)]] array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ var len : u32 = arrayLength(&sb);
+}
+)";
+
+ auto* expect = R"(
+[[internal(intrinsic_buffer_size)]]
+fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : [[stride(64)]] array<i32>, result : ptr<function, u32>)
+
+[[group(0), binding(0)]] var<storage, read> sb : [[stride(64)]] array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ var tint_symbol_1 : u32 = 0u;
+ tint_symbol(sb, &(tint_symbol_1));
+ let tint_symbol_2 : u32 = (tint_symbol_1 / 64u);
+ var len : u32 = tint_symbol_2;
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(CalculateArrayLengthTest, WithStride_InStruct) {
+ auto* src = R"(
struct SB {
x : i32;
y : f32;
@@ -233,15 +327,21 @@
[[group(0), binding(1)]] var<storage, read> sb2 : SB2;
+[[group(0), binding(2)]] var<storage, read> sb3 : array<i32>;
+
[[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);
+ var x : u32 = (len1 + len2 + len3);
}
)";
auto* expect = R"(
+[[internal(intrinsic_buffer_size)]]
+fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : array<i32>, result : ptr<function, u32>)
+
struct SB1 {
x : i32;
arr1 : array<i32>;
@@ -262,6 +362,8 @@
[[group(0), binding(1)]] var<storage, read> sb2 : SB2;
+[[group(0), binding(2)]] var<storage, read> sb3 : array<i32>;
+
[[stage(compute), workgroup_size(1)]]
fn main() {
var tint_symbol_1 : u32 = 0u;
@@ -270,9 +372,13 @@
var tint_symbol_4 : u32 = 0u;
tint_symbol_3(sb2, &(tint_symbol_4));
let tint_symbol_5 : u32 = ((tint_symbol_4 - 16u) / 16u);
+ var tint_symbol_7 : u32 = 0u;
+ tint_symbol_6(sb3, &(tint_symbol_7));
+ let tint_symbol_8 : u32 = (tint_symbol_7 / 4u);
var len1 : u32 = tint_symbol_2;
var len2 : u32 = tint_symbol_5;
- var x : u32 = (len1 + len2);
+ var len3 : u32 = tint_symbol_8;
+ var x : u32 = ((len1 + len2) + len3);
}
)";
diff --git a/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl b/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl
new file mode 100644
index 0000000..2e2966d
--- /dev/null
+++ b/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> G : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let p = &G;
+ let p2 = &(*p);
+ let l1 : u32 = arrayLength(p2);
+}
diff --git a/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.hlsl b/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.hlsl
new file mode 100644
index 0000000..d5cdca5
--- /dev/null
+++ b/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.hlsl
@@ -0,0 +1,10 @@
+ByteAddressBuffer G : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ uint tint_symbol_1 = 0u;
+ G.GetDimensions(tint_symbol_1);
+ const uint tint_symbol_2 = (tint_symbol_1 / 4u);
+ const uint l1 = tint_symbol_2;
+ return;
+}
diff --git a/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.msl b/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.msl
new file mode 100644
index 0000000..0ab65d1
--- /dev/null
+++ b/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.msl
@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+ /* 0x0000 */ uint4 buffer_size[1];
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+ uint const l1 = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+ return;
+}
+
diff --git a/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.spvasm b/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.spvasm
new file mode 100644
index 0000000..23ceb5f
--- /dev/null
+++ b/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.spvasm
@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %G_block "G_block"
+ OpMemberName %G_block 0 "inner"
+ OpName %G "G"
+ OpName %main "main"
+ OpDecorate %G_block Block
+ OpMemberDecorate %G_block 0 Offset 0
+ OpDecorate %_runtimearr_int ArrayStride 4
+ OpDecorate %G NonWritable
+ OpDecorate %G DescriptorSet 0
+ OpDecorate %G Binding 0
+ %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+ %G_block = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_G_block = OpTypePointer StorageBuffer %G_block
+ %G = OpVariable %_ptr_StorageBuffer_G_block StorageBuffer
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %main = OpFunction %void None %6
+ %9 = OpLabel
+ %10 = OpArrayLength %uint %G 0
+ OpReturn
+ OpFunctionEnd
diff --git a/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.wgsl b/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.wgsl
new file mode 100644
index 0000000..56b0e0a
--- /dev/null
+++ b/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> G : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let p = &(G);
+ let p2 = &(*(p));
+ let l1 : u32 = arrayLength(p2);
+}
diff --git a/test/intrinsics/arrayLength/simple_no_struct.wgsl b/test/intrinsics/arrayLength/simple_no_struct.wgsl
new file mode 100644
index 0000000..0fcbfbe
--- /dev/null
+++ b/test/intrinsics/arrayLength/simple_no_struct.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<storage, read> G : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let l1 : u32 = arrayLength(&G);
+}
diff --git a/test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.hlsl b/test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.hlsl
new file mode 100644
index 0000000..d5cdca5
--- /dev/null
+++ b/test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.hlsl
@@ -0,0 +1,10 @@
+ByteAddressBuffer G : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ uint tint_symbol_1 = 0u;
+ G.GetDimensions(tint_symbol_1);
+ const uint tint_symbol_2 = (tint_symbol_1 / 4u);
+ const uint l1 = tint_symbol_2;
+ return;
+}
diff --git a/test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.msl b/test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.msl
new file mode 100644
index 0000000..0ab65d1
--- /dev/null
+++ b/test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.msl
@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+ /* 0x0000 */ uint4 buffer_size[1];
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+ uint const l1 = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+ return;
+}
+
diff --git a/test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.spvasm b/test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.spvasm
new file mode 100644
index 0000000..23ceb5f
--- /dev/null
+++ b/test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.spvasm
@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %G_block "G_block"
+ OpMemberName %G_block 0 "inner"
+ OpName %G "G"
+ OpName %main "main"
+ OpDecorate %G_block Block
+ OpMemberDecorate %G_block 0 Offset 0
+ OpDecorate %_runtimearr_int ArrayStride 4
+ OpDecorate %G NonWritable
+ OpDecorate %G DescriptorSet 0
+ OpDecorate %G Binding 0
+ %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+ %G_block = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_G_block = OpTypePointer StorageBuffer %G_block
+ %G = OpVariable %_ptr_StorageBuffer_G_block StorageBuffer
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %main = OpFunction %void None %6
+ %9 = OpLabel
+ %10 = OpArrayLength %uint %G 0
+ OpReturn
+ OpFunctionEnd
diff --git a/test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.wgsl b/test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.wgsl
new file mode 100644
index 0000000..cbe8d72
--- /dev/null
+++ b/test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<storage, read> G : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let l1 : u32 = arrayLength(&(G));
+}
diff --git a/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl b/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl
new file mode 100644
index 0000000..4e8a6c5
--- /dev/null
+++ b/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl
@@ -0,0 +1,9 @@
+[[group(0), binding(0)]] var<storage, read> G : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let p = &*&G;
+ let p2 = &*p;
+ let p3 = &(*p);
+ let l1 : u32 = arrayLength(&*p3);
+}
diff --git a/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.hlsl b/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.hlsl
new file mode 100644
index 0000000..d5cdca5
--- /dev/null
+++ b/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.hlsl
@@ -0,0 +1,10 @@
+ByteAddressBuffer G : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ uint tint_symbol_1 = 0u;
+ G.GetDimensions(tint_symbol_1);
+ const uint tint_symbol_2 = (tint_symbol_1 / 4u);
+ const uint l1 = tint_symbol_2;
+ return;
+}
diff --git a/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.msl b/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.msl
new file mode 100644
index 0000000..0ab65d1
--- /dev/null
+++ b/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.msl
@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+ /* 0x0000 */ uint4 buffer_size[1];
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+ uint const l1 = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+ return;
+}
+
diff --git a/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.spvasm b/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.spvasm
new file mode 100644
index 0000000..23ceb5f
--- /dev/null
+++ b/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.spvasm
@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %G_block "G_block"
+ OpMemberName %G_block 0 "inner"
+ OpName %G "G"
+ OpName %main "main"
+ OpDecorate %G_block Block
+ OpMemberDecorate %G_block 0 Offset 0
+ OpDecorate %_runtimearr_int ArrayStride 4
+ OpDecorate %G NonWritable
+ OpDecorate %G DescriptorSet 0
+ OpDecorate %G Binding 0
+ %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+ %G_block = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_G_block = OpTypePointer StorageBuffer %G_block
+ %G = OpVariable %_ptr_StorageBuffer_G_block StorageBuffer
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %main = OpFunction %void None %6
+ %9 = OpLabel
+ %10 = OpArrayLength %uint %G 0
+ OpReturn
+ OpFunctionEnd
diff --git a/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.wgsl b/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.wgsl
new file mode 100644
index 0000000..97964a4
--- /dev/null
+++ b/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.wgsl
@@ -0,0 +1,9 @@
+[[group(0), binding(0)]] var<storage, read> G : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let p = &(*(&(G)));
+ let p2 = &(*(p));
+ let p3 = &(*(p));
+ let l1 : u32 = arrayLength(&(*(p3)));
+}
diff --git a/test/intrinsics/arrayLength/via_let_no_struct.wgsl b/test/intrinsics/arrayLength/via_let_no_struct.wgsl
new file mode 100644
index 0000000..83ab810
--- /dev/null
+++ b/test/intrinsics/arrayLength/via_let_no_struct.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> G : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let p = &G;
+ let p2 = p;
+ let l1 : u32 = arrayLength(p2);
+}
diff --git a/test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.hlsl b/test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.hlsl
new file mode 100644
index 0000000..d5cdca5
--- /dev/null
+++ b/test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.hlsl
@@ -0,0 +1,10 @@
+ByteAddressBuffer G : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ uint tint_symbol_1 = 0u;
+ G.GetDimensions(tint_symbol_1);
+ const uint tint_symbol_2 = (tint_symbol_1 / 4u);
+ const uint l1 = tint_symbol_2;
+ return;
+}
diff --git a/test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.msl b/test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.msl
new file mode 100644
index 0000000..0ab65d1
--- /dev/null
+++ b/test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.msl
@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+ /* 0x0000 */ uint4 buffer_size[1];
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+ uint const l1 = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+ return;
+}
+
diff --git a/test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.spvasm b/test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.spvasm
new file mode 100644
index 0000000..23ceb5f
--- /dev/null
+++ b/test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.spvasm
@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %G_block "G_block"
+ OpMemberName %G_block 0 "inner"
+ OpName %G "G"
+ OpName %main "main"
+ OpDecorate %G_block Block
+ OpMemberDecorate %G_block 0 Offset 0
+ OpDecorate %_runtimearr_int ArrayStride 4
+ OpDecorate %G NonWritable
+ OpDecorate %G DescriptorSet 0
+ OpDecorate %G Binding 0
+ %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+ %G_block = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_G_block = OpTypePointer StorageBuffer %G_block
+ %G = OpVariable %_ptr_StorageBuffer_G_block StorageBuffer
+ %void = OpTypeVoid
+ %6 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %main = OpFunction %void None %6
+ %9 = OpLabel
+ %10 = OpArrayLength %uint %G 0
+ OpReturn
+ OpFunctionEnd
diff --git a/test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.wgsl b/test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.wgsl
new file mode 100644
index 0000000..731ba0c
--- /dev/null
+++ b/test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> G : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+ let p = &(G);
+ let p2 = p;
+ let l1 : u32 = arrayLength(p2);
+}
