diff --git a/src/transform/array_length_from_uniform.cc b/src/transform/array_length_from_uniform.cc index 5084a548f3..32289fe839 100644 --- a/src/transform/array_length_from_uniform.cc +++ b/src/transform/array_length_from_uniform.cc @@ -59,30 +59,26 @@ static void IterateArrayLengthOnStorageVar(CloneContext& ctx, F&& functor) { } // 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(); 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(); - 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()) { + storage_buffer_expr = accessor->structure; } - auto* storage_buffer_expr = accessor->structure; - auto* storage_buffer_sem = - sem.Get(storage_buffer_expr)->As(); + auto* storage_buffer_sem = sem.Get(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 @@ void ArrayLengthFromUniform::Run(CloneContext& ctx, // total_storage_buffer_size - array_offset // array_length = ---------------------------------------- // array_stride - auto* storage_buffer_type = - storage_buffer_sem->Type()->UnwrapRef()->As(); - 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()) { + // 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(); + total_size = ctx.dst->Sub(total_storage_buffer_size, + array_member_sem->Offset()); + } else if (auto* arr = storage_buffer_type->As()) { + 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 59f5a4d383..77fb0c3d75 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, Error_MissingSimplifyPointers) { TEST_F(ArrayLengthFromUniformTest, Basic) { auto* src = R"( +[[group(0), binding(0)]] var sb : array; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var len : u32 = arrayLength(&sb); +} +)"; + + auto* expect = R"( +struct tint_symbol { + buffer_size : array, 1u>; +}; + +[[group(0), binding(30)]] var tint_symbol_1 : tint_symbol; + +[[group(0), binding(0)]] var sb : array; + +[[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(std::move(cfg)); + + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); + EXPECT_EQ(std::unordered_set({0}), + got.data.Get()->used_size_indices); +} + +TEST_F(ArrayLengthFromUniformTest, BasicInStruct) { + auto* src = R"( struct SB { x : i32; arr : array; @@ -100,6 +138,44 @@ fn main() { TEST_F(ArrayLengthFromUniformTest, WithStride) { auto* src = R"( +[[group(0), binding(0)]] var sb : [[stride(64)]] array; + +[[stage(compute), workgroup_size(1)]] +fn main() { + var len : u32 = arrayLength(&sb); +} +)"; + + auto* expect = R"( +struct tint_symbol { + buffer_size : array, 1u>; +}; + +[[group(0), binding(30)]] var tint_symbol_1 : tint_symbol; + +[[group(0), binding(0)]] var sb : [[stride(64)]] array; + +[[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(std::move(cfg)); + + auto got = Run(src, data); + + EXPECT_EQ(expect, str(got)); + EXPECT_EQ(std::unordered_set({0}), + got.data.Get()->used_size_indices); +} + +TEST_F(ArrayLengthFromUniformTest, WithStride_InStruct) { + auto* src = R"( struct SB { x : i32; y : f32; @@ -158,32 +234,24 @@ struct SB2 { x : i32; arr2 : array>; }; -struct SB3 { - x : i32; - arr3 : array>; -}; struct SB4 { x : i32; arr4 : array>; }; -struct SB5 { - x : i32; - arr5 : array>; -}; [[group(0), binding(2)]] var sb1 : SB1; [[group(1), binding(2)]] var sb2 : SB2; -[[group(2), binding(2)]] var sb3 : SB3; +[[group(2), binding(2)]] var sb3 : array>; [[group(3), binding(2)]] var sb4 : SB4; -[[group(4), binding(2)]] var sb5 : SB5; +[[group(4), binding(2)]] var sb5 : array>; [[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 @@ struct SB2 { arr2 : array>; }; -struct SB3 { - x : i32; - arr3 : array>; -}; - struct SB4 { x : i32; arr4 : array>; }; -struct SB5 { - x : i32; - arr5 : array>; -}; - [[group(0), binding(2)]] var sb1 : SB1; [[group(1), binding(2)]] var sb2 : SB2; -[[group(2), binding(2)]] var sb3 : SB3; +[[group(2), binding(2)]] var sb3 : array>; [[group(3), binding(2)]] var sb4 : SB4; -[[group(4), binding(2)]] var sb5 : SB5; +[[group(4), binding(2)]] var sb5 : array>; [[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 @@ struct SB2 { x : i32; arr2 : array>; }; -struct SB3 { - x : i32; - arr3 : array>; -}; struct SB4 { x : i32; arr4 : array>; }; -struct SB5 { - x : i32; - arr5 : array>; -}; [[group(0), binding(2)]] var sb1 : SB1; [[group(1), binding(2)]] var sb2 : SB2; -[[group(2), binding(2)]] var sb3 : SB3; +[[group(2), binding(2)]] var sb3 : array>; [[group(3), binding(2)]] var sb4 : SB4; -[[group(4), binding(2)]] var sb5 : SB5; +[[group(4), binding(2)]] var sb5 : array>; [[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 @@ struct SB2 { arr2 : array>; }; -struct SB3 { - x : i32; - arr3 : array>; -}; - struct SB4 { x : i32; arr4 : array>; }; -struct SB5 { - x : i32; - arr5 : array>; -}; - [[group(0), binding(2)]] var sb1 : SB1; [[group(1), binding(2)]] var sb2 : SB2; -[[group(2), binding(2)]] var sb3 : SB3; +[[group(2), binding(2)]] var sb3 : array>; [[group(3), binding(2)]] var sb4 : SB4; -[[group(4), binding(2)]] var sb5 : SB5; +[[group(4), binding(2)]] var sb5 : array>; [[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 e06b95914e..6c6c6ffa5c 100644 --- a/src/transform/calculate_array_length.cc +++ b/src/transform/calculate_array_length.cc @@ -42,7 +42,7 @@ namespace { /// 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 @@ void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) { // 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 buffer_size_intrinsics; - auto get_buffer_size_intrinsic = [&](const sem::Struct* buffer_type) { + std::unordered_map 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( @@ -95,7 +94,7 @@ void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) { // in order for HLSL to emit this as a ByteAddressBuffer. ctx.dst->create( 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 @@ void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) { ctx.dst->ASTNodes().Create(ctx.dst->ID()), }, ast::DecorationList{}); - ctx.InsertAfter(ctx.src->AST().GlobalDeclarations(), - buffer_type->Declaration(), func); + if (auto* str = buffer_type->As()) { + ctx.InsertAfter(ctx.src->AST().GlobalDeclarations(), str->Declaration(), + func); + } else { + ctx.InsertFront(ctx.src->AST().GlobalDeclarations(), func); + } return name; }); }; @@ -123,71 +126,47 @@ void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) { 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(); 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(); - if (!accessor) { + auto* storage_buffer_expr = address_of->expr; + if (auto* accessor = + storage_buffer_expr->As()) { + storage_buffer_expr = accessor->structure; + } + auto* storage_buffer_sem = + sem.Get(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(); + 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()) { - 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 @@ void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) { // 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()) { + // 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(); + total_size = + ctx.dst->Sub(total_size, array_member_sem->Offset()); + } else if (auto* arr = storage_buffer_type->As()) { + 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 4522768795..4b42bf9084 100644 --- a/src/transform/calculate_array_length_test.cc +++ b/src/transform/calculate_array_length_test.cc @@ -38,6 +38,36 @@ TEST_F(CalculateArrayLengthTest, Error_MissingCalculateArrayLength) { TEST_F(CalculateArrayLengthTest, Basic) { auto* src = R"( +[[group(0), binding(0)]] var sb : array; + +[[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, result : ptr) + +[[group(0), binding(0)]] var sb : array; + +[[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(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(CalculateArrayLengthTest, BasicInStruct) { + auto* src = R"( struct SB { x : i32; arr : array; @@ -78,6 +108,40 @@ fn main() { TEST_F(CalculateArrayLengthTest, InSameBlock) { auto* src = R"( +[[group(0), binding(0)]] var sb : array;; + +[[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, result : ptr) + +[[group(0), binding(0)]] var sb : array; + +[[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(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(CalculateArrayLengthTest, InSameBlock_Struct) { + auto* src = R"( struct SB { x : i32; arr : array; @@ -122,6 +186,36 @@ fn main() { TEST_F(CalculateArrayLengthTest, WithStride) { auto* src = R"( +[[group(0), binding(0)]] var sb : [[stride(64)]] array; + +[[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, result : ptr) + +[[group(0), binding(0)]] var sb : [[stride(64)]] array; + +[[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(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(CalculateArrayLengthTest, WithStride_InStruct) { + auto* src = R"( struct SB { x : i32; y : f32; @@ -233,15 +327,21 @@ struct SB2 { [[group(0), binding(1)]] var sb2 : SB2; +[[group(0), binding(2)]] var sb3 : array; + [[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, result : ptr) + struct SB1 { x : i32; arr1 : array; @@ -262,6 +362,8 @@ fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_pa [[group(0), binding(1)]] var sb2 : SB2; +[[group(0), binding(2)]] var sb3 : array; + [[stage(compute), workgroup_size(1)]] fn main() { var tint_symbol_1 : u32 = 0u; @@ -270,9 +372,13 @@ fn main() { 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 0000000000..2e2966d627 --- /dev/null +++ b/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl @@ -0,0 +1,8 @@ +[[group(0), binding(0)]] var G : array; + +[[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 0000000000..d5cdca5b40 --- /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 0000000000..0ab65d100a --- /dev/null +++ b/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.msl @@ -0,0 +1,12 @@ +#include + +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 0000000000..23ceb5f344 --- /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 0000000000..56b0e0ab8b --- /dev/null +++ b/test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.wgsl @@ -0,0 +1,8 @@ +[[group(0), binding(0)]] var G : array; + +[[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 0000000000..0fcbfbecb4 --- /dev/null +++ b/test/intrinsics/arrayLength/simple_no_struct.wgsl @@ -0,0 +1,6 @@ +[[group(0), binding(0)]] var G : array; + +[[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 0000000000..d5cdca5b40 --- /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 0000000000..0ab65d100a --- /dev/null +++ b/test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.msl @@ -0,0 +1,12 @@ +#include + +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 0000000000..23ceb5f344 --- /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 0000000000..cbe8d72f41 --- /dev/null +++ b/test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +[[group(0), binding(0)]] var G : array; + +[[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 0000000000..4e8a6c5a37 --- /dev/null +++ b/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl @@ -0,0 +1,9 @@ +[[group(0), binding(0)]] var G : array; + +[[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 0000000000..d5cdca5b40 --- /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 0000000000..0ab65d100a --- /dev/null +++ b/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.msl @@ -0,0 +1,12 @@ +#include + +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 0000000000..23ceb5f344 --- /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 0000000000..97964a41e4 --- /dev/null +++ b/test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.wgsl @@ -0,0 +1,9 @@ +[[group(0), binding(0)]] var G : array; + +[[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 0000000000..83ab810423 --- /dev/null +++ b/test/intrinsics/arrayLength/via_let_no_struct.wgsl @@ -0,0 +1,8 @@ +[[group(0), binding(0)]] var G : array; + +[[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 0000000000..d5cdca5b40 --- /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 0000000000..0ab65d100a --- /dev/null +++ b/test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.msl @@ -0,0 +1,12 @@ +#include + +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 0000000000..23ceb5f344 --- /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 0000000000..731ba0c459 --- /dev/null +++ b/test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.wgsl @@ -0,0 +1,8 @@ +[[group(0), binding(0)]] var G : array; + +[[stage(compute), workgroup_size(1)]] +fn main() { + let p = &(G); + let p2 = p; + let l1 : u32 = arrayLength(p2); +}