From 009d129103eebe3eb60222971a3734c7274ee3c5 Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Fri, 4 Feb 2022 15:39:34 +0000 Subject: [PATCH] reader/spirv: Decompose arrays with strides Transform any SPIR-V that has an array with a custom stride: @stride(S) array into: struct strided_arr { @size(S) er : T; }; array Also remove any @stride decorations that match the default array stride. Bug: tint:1394 Bug: tint:1381 Change-Id: I8be8f3a76c5335fdb2bc5183388366091dbc7642 Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/78781 Reviewed-by: David Neto Kokoro: Kokoro Commit-Queue: Ben Clayton --- src/BUILD.gn | 2 + src/CMakeLists.txt | 3 + src/reader/spirv/parser.cc | 2 + src/sem/type_mappings.h | 2 + src/transform/decompose_strided_array.cc | 162 ++++ src/transform/decompose_strided_array.h | 61 ++ src/transform/decompose_strided_array_test.cc | 698 ++++++++++++++++++ src/transform/test_helper.h | 12 +- test/BUILD.gn | 1 + test/array/strides.spvasm | 71 ++ test/array/strides.spvasm.expected.glsl | 38 + test/array/strides.spvasm.expected.hlsl | 103 +++ test/array/strides.spvasm.expected.msl | 40 + test/array/strides.spvasm.expected.spvasm | 74 ++ test/array/strides.spvasm.expected.wgsl | 36 + test/bug/tint/1088.spvasm.expected.glsl | 6 +- test/bug/tint/1088.spvasm.expected.msl | 4 +- test/bug/tint/1088.spvasm.expected.spvasm | 158 ++-- test/bug/tint/1088.spvasm.expected.wgsl | 11 +- test/bug/tint/870.spvasm.expected.wgsl | 2 +- test/bug/tint/943.spvasm.expected.wgsl | 6 +- test/bug/tint/951.spvasm.expected.wgsl | 4 +- test/bug/tint/977.spvasm.expected.wgsl | 6 +- .../mat2x2/stride/16.spvasm.expected.glsl | 16 +- .../mat2x2/stride/16.spvasm.expected.hlsl | 35 +- .../mat2x2/stride/16.spvasm.expected.msl | 20 +- .../mat2x2/stride/16.spvasm.expected.spvasm | 72 +- .../mat2x2/stride/16.spvasm.expected.wgsl | 15 +- 28 files changed, 1499 insertions(+), 161 deletions(-) create mode 100644 src/transform/decompose_strided_array.cc create mode 100644 src/transform/decompose_strided_array.h create mode 100644 src/transform/decompose_strided_array_test.cc create mode 100644 test/array/strides.spvasm create mode 100644 test/array/strides.spvasm.expected.glsl create mode 100644 test/array/strides.spvasm.expected.hlsl create mode 100644 test/array/strides.spvasm.expected.msl create mode 100644 test/array/strides.spvasm.expected.spvasm create mode 100644 test/array/strides.spvasm.expected.wgsl diff --git a/src/BUILD.gn b/src/BUILD.gn index 58b2f65ee2..fce79f0308 100644 --- a/src/BUILD.gn +++ b/src/BUILD.gn @@ -441,6 +441,8 @@ libtint_source_set("libtint_core_all_src") { "transform/combine_samplers.h", "transform/decompose_memory_access.cc", "transform/decompose_memory_access.h", + "transform/decompose_strided_array.cc", + "transform/decompose_strided_array.h", "transform/decompose_strided_matrix.cc", "transform/decompose_strided_matrix.h", "transform/external_texture_transform.cc", diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 600e91f7e2..b0e09618b1 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -309,6 +309,8 @@ set(TINT_LIB_SRCS transform/canonicalize_entry_point_io.h transform/decompose_memory_access.cc transform/decompose_memory_access.h + transform/decompose_strided_array.cc + transform/decompose_strided_array.h transform/decompose_strided_matrix.cc transform/decompose_strided_matrix.h transform/external_texture_transform.cc @@ -984,6 +986,7 @@ if(TINT_BUILD_TESTS) transform/canonicalize_entry_point_io_test.cc transform/combine_samplers_test.cc transform/decompose_memory_access_test.cc + transform/decompose_strided_array_test.cc transform/decompose_strided_matrix_test.cc transform/external_texture_transform_test.cc transform/first_index_offset_test.cc diff --git a/src/reader/spirv/parser.cc b/src/reader/spirv/parser.cc index 2711733f9b..e48357e1a2 100644 --- a/src/reader/spirv/parser.cc +++ b/src/reader/spirv/parser.cc @@ -17,6 +17,7 @@ #include #include "src/reader/spirv/parser_impl.h" +#include "src/transform/decompose_strided_array.h" #include "src/transform/decompose_strided_matrix.h" #include "src/transform/manager.h" #include "src/transform/remove_unreachable_statements.h" @@ -54,6 +55,7 @@ Program Parse(const std::vector& input) { manager.Add(); manager.Add(); manager.Add(); + manager.Add(); manager.Add(); return manager.Run(&program).program; } diff --git a/src/sem/type_mappings.h b/src/sem/type_mappings.h index 8d23a1538a..5dbc059995 100644 --- a/src/sem/type_mappings.h +++ b/src/sem/type_mappings.h @@ -21,6 +21,7 @@ namespace tint { // Forward declarations namespace ast { +class Array; class CallExpression; class Expression; class ElseStatement; @@ -60,6 +61,7 @@ class Variable; /// rules will be used to infer the return type based on the argument type. struct TypeMappings { //! @cond Doxygen_Suppress + Array* operator()(ast::Array*); Call* operator()(ast::CallExpression*); Expression* operator()(ast::Expression*); ElseStatement* operator()(ast::ElseStatement*); diff --git a/src/transform/decompose_strided_array.cc b/src/transform/decompose_strided_array.cc new file mode 100644 index 0000000000..106fa56d09 --- /dev/null +++ b/src/transform/decompose_strided_array.cc @@ -0,0 +1,162 @@ +// Copyright 2022 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/transform/decompose_strided_array.h" + +#include +#include +#include + +#include "src/program_builder.h" +#include "src/sem/call.h" +#include "src/sem/expression.h" +#include "src/sem/member_accessor_expression.h" +#include "src/sem/type_constructor.h" +#include "src/transform/simplify_pointers.h" +#include "src/utils/hash.h" +#include "src/utils/map.h" + +TINT_INSTANTIATE_TYPEINFO(tint::transform::DecomposeStridedArray); + +namespace tint { +namespace transform { +namespace { + +using DecomposedArrays = std::unordered_map; + +} // namespace + +DecomposeStridedArray::DecomposeStridedArray() = default; + +DecomposeStridedArray::~DecomposeStridedArray() = default; + +bool DecomposeStridedArray::ShouldRun(const Program* program, + const DataMap&) const { + for (auto* node : program->ASTNodes().Objects()) { + if (auto* ast = node->As()) { + if (ast::GetAttribute(ast->attributes)) { + return true; + } + } + } + return false; +} + +void DecomposeStridedArray::Run(CloneContext& ctx, + const DataMap&, + DataMap&) const { + const auto& sem = ctx.src->Sem(); + + static constexpr const char* kMemberName = "el"; + + // Maps an array type in the source program to the name of the struct wrapper + // type in the target program. + std::unordered_map decomposed; + + // Find and replace all arrays with a @stride attribute with a array that has + // the @stride removed. If the source array stride does not match the natural + // stride for the array element type, then replace the array element type with + // a structure, holding a single field with a @size attribute equal to the + // array stride. + ctx.ReplaceAll([&](const ast::Array* ast) -> const ast::Array* { + if (auto* arr = sem.Get(ast)) { + if (!arr->IsStrideImplicit()) { + auto el_ty = utils::GetOrCreate(decomposed, arr, [&] { + auto name = ctx.dst->Symbols().New("strided_arr"); + auto* member_ty = ctx.Clone(ast->type); + auto* member = ctx.dst->Member(kMemberName, member_ty, + {ctx.dst->MemberSize(arr->Stride())}); + ctx.dst->Structure(name, {member}); + return name; + }); + auto* count = ctx.Clone(ast->count); + return ctx.dst->ty.array(ctx.dst->ty.type_name(el_ty), count); + } + if (ast::GetAttribute(ast->attributes)) { + // Strip the @stride attribute + auto* ty = ctx.Clone(ast->type); + auto* count = ctx.Clone(ast->count); + return ctx.dst->ty.array(ty, count); + } + } + return nullptr; + }); + + // Find all array index-accessors expressions for arrays that have had their + // element changed to a single field structure. These expressions are adjusted + // to insert an additional member accessor for the single structure field. + // Example: `arr[i]` -> `arr[i].el` + ctx.ReplaceAll( + [&](const ast::IndexAccessorExpression* idx) -> const ast::Expression* { + if (auto* ty = ctx.src->TypeOf(idx->object)) { + if (auto* arr = ty->UnwrapRef()->As()) { + if (!arr->IsStrideImplicit()) { + auto* expr = ctx.CloneWithoutTransform(idx); + return ctx.dst->MemberAccessor(expr, kMemberName); + } + } + } + return nullptr; + }); + + // Find all array type constructor expressions for array types that have had + // their element changed to a single field structure. These constructors are + // adjusted to wrap each of the arguments with an additional constructor for + // the new element structure type. + // Example: + // `@stride(32) array(1, 2, 3)` + // -> + // `array(strided_arr(1), strided_arr(2), strided_arr(3))` + ctx.ReplaceAll( + [&](const ast::CallExpression* expr) -> const ast::Expression* { + if (!expr->args.empty()) { + if (auto* call = sem.Get(expr)) { + if (auto* ctor = call->Target()->As()) { + if (auto* arr = ctor->ReturnType()->As()) { + // Begin by cloning the array constructor type or name + // If this is an unaliased array, this may add a new entry to + // decomposed. + // If this is an aliased array, decomposed should already be + // populated with any strided aliases. + ast::CallExpression::Target target; + if (expr->target.type) { + target.type = ctx.Clone(expr->target.type); + } else { + target.name = ctx.Clone(expr->target.name); + } + + ast::ExpressionList args; + if (auto it = decomposed.find(arr); it != decomposed.end()) { + args.reserve(expr->args.size()); + for (auto* arg : expr->args) { + args.emplace_back( + ctx.dst->Call(it->second, ctx.Clone(arg))); + } + } else { + args = ctx.Clone(expr->args); + } + + return target.type ? ctx.dst->Construct(target.type, args) + : ctx.dst->Call(target.name, args); + } + } + } + } + return nullptr; + }); + ctx.Clone(); +} + +} // namespace transform +} // namespace tint diff --git a/src/transform/decompose_strided_array.h b/src/transform/decompose_strided_array.h new file mode 100644 index 0000000000..27d4de0b5a --- /dev/null +++ b/src/transform/decompose_strided_array.h @@ -0,0 +1,61 @@ +// 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. + +#ifndef SRC_TRANSFORM_DECOMPOSE_STRIDED_ARRAY_H_ +#define SRC_TRANSFORM_DECOMPOSE_STRIDED_ARRAY_H_ + +#include "src/transform/transform.h" + +namespace tint { +namespace transform { + +/// DecomposeStridedArray transforms replaces arrays with a non-default +/// `@stride` attribute with an array of structure elements, where the +/// structure contains a single field with an equivalent `@size` attribute. +/// `@stride` attributes on arrays that match the default stride are also +/// removed. +/// +/// @note Depends on the following transforms to have been run first: +/// * SimplifyPointers +class DecomposeStridedArray + : public Castable { + public: + /// Constructor + DecomposeStridedArray(); + + /// Destructor + ~DecomposeStridedArray() override; + + /// @param program the program to inspect + /// @param data optional extra transform-specific input data + /// @returns true if this transform should be run for the given program + bool ShouldRun(const Program* program, + const DataMap& data = {}) const override; + + protected: + /// Runs the transform using the CloneContext built for transforming a + /// program. Run() is responsible for calling Clone() on the CloneContext. + /// @param ctx the CloneContext primed with the input program and + /// ProgramBuilder + /// @param inputs optional extra transform-specific input data + /// @param outputs optional extra transform-specific output data + void Run(CloneContext& ctx, + const DataMap& inputs, + DataMap& outputs) const override; +}; + +} // namespace transform +} // namespace tint + +#endif // SRC_TRANSFORM_DECOMPOSE_STRIDED_ARRAY_H_ diff --git a/src/transform/decompose_strided_array_test.cc b/src/transform/decompose_strided_array_test.cc new file mode 100644 index 0000000000..e982c9a91a --- /dev/null +++ b/src/transform/decompose_strided_array_test.cc @@ -0,0 +1,698 @@ +// Copyright 2022 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/transform/decompose_strided_array.h" + +#include +#include +#include + +#include "src/program_builder.h" +#include "src/transform/simplify_pointers.h" +#include "src/transform/test_helper.h" +#include "src/transform/unshadow.h" + +namespace tint { +namespace transform { +namespace { + +using DecomposeStridedArrayTest = TransformTest; +using f32 = ProgramBuilder::f32; + +TEST_F(DecomposeStridedArrayTest, ShouldRunEmptyModule) { + ProgramBuilder b; + EXPECT_FALSE(ShouldRun(Program(std::move(b)))); +} + +TEST_F(DecomposeStridedArrayTest, ShouldRunNonStridedArray) { + // var arr : array + + ProgramBuilder b; + b.Global("arr", b.ty.array(), ast::StorageClass::kPrivate); + EXPECT_FALSE(ShouldRun(Program(std::move(b)))); +} + +TEST_F(DecomposeStridedArrayTest, ShouldRunDefaultStridedArray) { + // var arr : @stride(4) array + + ProgramBuilder b; + b.Global("arr", b.ty.array(4), ast::StorageClass::kPrivate); + EXPECT_TRUE(ShouldRun(Program(std::move(b)))); +} + +TEST_F(DecomposeStridedArrayTest, ShouldRunExplicitStridedArray) { + // var arr : @stride(16) array + + ProgramBuilder b; + b.Global("arr", b.ty.array(16), ast::StorageClass::kPrivate); + EXPECT_TRUE(ShouldRun(Program(std::move(b)))); +} + +TEST_F(DecomposeStridedArrayTest, Empty) { + auto* src = R"()"; + auto* expect = src; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(DecomposeStridedArrayTest, PrivateDefaultStridedArray) { + // var arr : @stride(4) array + // + // @stage(compute) @workgroup_size(1) + // fn f() { + // let a : @stride(4) array = a; + // let b : f32 = arr[1]; + // } + + ProgramBuilder b; + b.Global("arr", b.ty.array(4), ast::StorageClass::kPrivate); + b.Func("f", {}, b.ty.void_(), + { + b.Decl(b.Const("a", b.ty.array(4), b.Expr("arr"))), + b.Decl(b.Const("b", b.ty.f32(), b.IndexAccessor("arr", 1))), + }, + { + b.Stage(ast::PipelineStage::kCompute), + b.WorkgroupSize(1), + }); + + auto* expect = R"( +var arr : array; + +@stage(compute) @workgroup_size(1) +fn f() { + let a : array = arr; + let b : f32 = arr[1]; +} +)"; + + auto got = Run( + Program(std::move(b))); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(DecomposeStridedArrayTest, PrivateStridedArray) { + // var arr : @stride(32) array + // + // @stage(compute) @workgroup_size(1) + // fn f() { + // let a : @stride(32) array = a; + // let b : f32 = arr[1]; + // } + + ProgramBuilder b; + b.Global("arr", b.ty.array(32), ast::StorageClass::kPrivate); + b.Func("f", {}, b.ty.void_(), + { + b.Decl(b.Const("a", b.ty.array(32), b.Expr("arr"))), + b.Decl(b.Const("b", b.ty.f32(), b.IndexAccessor("arr", 1))), + }, + { + b.Stage(ast::PipelineStage::kCompute), + b.WorkgroupSize(1), + }); + + auto* expect = R"( +struct strided_arr { + @size(32) + el : f32; +} + +var arr : array; + +@stage(compute) @workgroup_size(1) +fn f() { + let a : array = arr; + let b : f32 = arr[1].el; +} +)"; + + auto got = Run( + Program(std::move(b))); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(DecomposeStridedArrayTest, ReadUniformStridedArray) { + // struct S { + // a : @stride(32) array; + // }; + // @group(0) @binding(0) var s : S; + // + // @stage(compute) @workgroup_size(1) + // fn f() { + // let a : @stride(32) array = s.a; + // let b : f32 = s.a[1]; + // } + ProgramBuilder b; + auto* S = b.Structure("S", {b.Member("a", b.ty.array(32))}); + b.Global("s", b.ty.Of(S), ast::StorageClass::kUniform, + b.GroupAndBinding(0, 0)); + b.Func("f", {}, b.ty.void_(), + { + b.Decl(b.Const("a", b.ty.array(32), + b.MemberAccessor("s", "a"))), + b.Decl(b.Const("b", b.ty.f32(), + b.IndexAccessor(b.MemberAccessor("s", "a"), 1))), + }, + { + b.Stage(ast::PipelineStage::kCompute), + b.WorkgroupSize(1), + }); + + auto* expect = R"( +struct strided_arr { + @size(32) + el : f32; +} + +struct S { + a : array; +} + +@group(0) @binding(0) var s : S; + +@stage(compute) @workgroup_size(1) +fn f() { + let a : array = s.a; + let b : f32 = s.a[1].el; +} +)"; + + auto got = Run( + Program(std::move(b))); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(DecomposeStridedArrayTest, ReadUniformDefaultStridedArray) { + // struct S { + // a : @stride(16) array, 4>; + // }; + // @group(0) @binding(0) var s : S; + // + // @stage(compute) @workgroup_size(1) + // fn f() { + // let a : @stride(16) array, 4> = s.a; + // let b : f32 = s.a[1][2]; + // } + ProgramBuilder b; + auto* S = + b.Structure("S", {b.Member("a", b.ty.array(b.ty.vec4(), 4, 16))}); + b.Global("s", b.ty.Of(S), ast::StorageClass::kUniform, + b.GroupAndBinding(0, 0)); + b.Func("f", {}, b.ty.void_(), + { + b.Decl(b.Const("a", b.ty.array(b.ty.vec4(), 4, 16), + b.MemberAccessor("s", "a"))), + b.Decl(b.Const( + "b", b.ty.f32(), + b.IndexAccessor(b.IndexAccessor(b.MemberAccessor("s", "a"), 1), + 2))), + }, + { + b.Stage(ast::PipelineStage::kCompute), + b.WorkgroupSize(1), + }); + + auto* expect = + R"( +struct S { + a : array, 4>; +} + +@group(0) @binding(0) var s : S; + +@stage(compute) @workgroup_size(1) +fn f() { + let a : array, 4> = s.a; + let b : f32 = s.a[1][2]; +} +)"; + + auto got = Run( + Program(std::move(b))); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(DecomposeStridedArrayTest, ReadStorageStridedArray) { + // struct S { + // a : @stride(32) array; + // }; + // @group(0) @binding(0) var s : S; + // + // @stage(compute) @workgroup_size(1) + // fn f() { + // let a : @stride(32) array = s.a; + // let b : f32 = s.a[1]; + // } + ProgramBuilder b; + auto* S = b.Structure("S", {b.Member("a", b.ty.array(32))}); + b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage, + b.GroupAndBinding(0, 0)); + b.Func("f", {}, b.ty.void_(), + { + b.Decl(b.Const("a", b.ty.array(32), + b.MemberAccessor("s", "a"))), + b.Decl(b.Const("b", b.ty.f32(), + b.IndexAccessor(b.MemberAccessor("s", "a"), 1))), + }, + { + b.Stage(ast::PipelineStage::kCompute), + b.WorkgroupSize(1), + }); + + auto* expect = R"( +struct strided_arr { + @size(32) + el : f32; +} + +struct S { + a : array; +} + +@group(0) @binding(0) var s : S; + +@stage(compute) @workgroup_size(1) +fn f() { + let a : array = s.a; + let b : f32 = s.a[1].el; +} +)"; + + auto got = Run( + Program(std::move(b))); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(DecomposeStridedArrayTest, ReadStorageDefaultStridedArray) { + // struct S { + // a : @stride(4) array; + // }; + // @group(0) @binding(0) var s : S; + // + // @stage(compute) @workgroup_size(1) + // fn f() { + // let a : @stride(4) array = s.a; + // let b : f32 = s.a[1]; + // } + ProgramBuilder b; + auto* S = b.Structure("S", {b.Member("a", b.ty.array(4))}); + b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage, + b.GroupAndBinding(0, 0)); + b.Func("f", {}, b.ty.void_(), + { + b.Decl(b.Const("a", b.ty.array(4), + b.MemberAccessor("s", "a"))), + b.Decl(b.Const("b", b.ty.f32(), + b.IndexAccessor(b.MemberAccessor("s", "a"), 1))), + }, + { + b.Stage(ast::PipelineStage::kCompute), + b.WorkgroupSize(1), + }); + + auto* expect = R"( +struct S { + a : array; +} + +@group(0) @binding(0) var s : S; + +@stage(compute) @workgroup_size(1) +fn f() { + let a : array = s.a; + let b : f32 = s.a[1]; +} +)"; + + auto got = Run( + Program(std::move(b))); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(DecomposeStridedArrayTest, WriteStorageStridedArray) { + // struct S { + // a : @stride(32) array; + // }; + // @group(0) @binding(0) var s : S; + // + // @stage(compute) @workgroup_size(1) + // fn f() { + // s.a = @stride(32) array(); + // s.a = @stride(32) array(1.0, 2.0, 3.0, 4.0); + // s.a[1] = 5.0; + // } + ProgramBuilder b; + auto* S = b.Structure("S", {b.Member("a", b.ty.array(32))}); + b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage, + ast::Access::kReadWrite, b.GroupAndBinding(0, 0)); + b.Func( + "f", {}, b.ty.void_(), + { + b.Assign(b.MemberAccessor("s", "a"), + b.Construct(b.ty.array(32))), + b.Assign(b.MemberAccessor("s", "a"), + b.Construct(b.ty.array(32), 1.0f, 2.0f, 3.0f, 4.0f)), + b.Assign(b.IndexAccessor(b.MemberAccessor("s", "a"), 1), 5.0f), + }, + { + b.Stage(ast::PipelineStage::kCompute), + b.WorkgroupSize(1), + }); + + auto* expect = + R"( +struct strided_arr { + @size(32) + el : f32; +} + +struct S { + a : array; +} + +@group(0) @binding(0) var s : S; + +@stage(compute) @workgroup_size(1) +fn f() { + s.a = array(); + s.a = array(strided_arr(1.0), strided_arr(2.0), strided_arr(3.0), strided_arr(4.0)); + s.a[1].el = 5.0; +} +)"; + + auto got = Run( + Program(std::move(b))); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(DecomposeStridedArrayTest, WriteStorageDefaultStridedArray) { + // struct S { + // a : @stride(4) array; + // }; + // @group(0) @binding(0) var s : S; + // + // @stage(compute) @workgroup_size(1) + // fn f() { + // s.a = @stride(4) array(); + // s.a = @stride(4) array(1.0, 2.0, 3.0, 4.0); + // s.a[1] = 5.0; + // } + ProgramBuilder b; + auto* S = b.Structure("S", {b.Member("a", b.ty.array(4))}); + b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage, + ast::Access::kReadWrite, b.GroupAndBinding(0, 0)); + b.Func( + "f", {}, b.ty.void_(), + { + b.Assign(b.MemberAccessor("s", "a"), + b.Construct(b.ty.array(4))), + b.Assign(b.MemberAccessor("s", "a"), + b.Construct(b.ty.array(4), 1.0f, 2.0f, 3.0f, 4.0f)), + b.Assign(b.IndexAccessor(b.MemberAccessor("s", "a"), 1), 5.0f), + }, + { + b.Stage(ast::PipelineStage::kCompute), + b.WorkgroupSize(1), + }); + + auto* expect = + R"( +struct S { + a : array; +} + +@group(0) @binding(0) var s : S; + +@stage(compute) @workgroup_size(1) +fn f() { + s.a = array(); + s.a = array(1.0, 2.0, 3.0, 4.0); + s.a[1] = 5.0; +} +)"; + + auto got = Run( + Program(std::move(b))); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(DecomposeStridedArrayTest, ReadWriteViaPointerLets) { + // struct S { + // a : @stride(32) array; + // }; + // @group(0) @binding(0) var s : S; + // + // @stage(compute) @workgroup_size(1) + // fn f() { + // let a = &s.a; + // let b = &*&*(a); + // let c = *b; + // let d = (*b)[1]; + // (*b) = @stride(32) array(1.0, 2.0, 3.0, 4.0); + // (*b)[1] = 5.0; + // } + ProgramBuilder b; + auto* S = b.Structure("S", {b.Member("a", b.ty.array(32))}); + b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage, + ast::Access::kReadWrite, b.GroupAndBinding(0, 0)); + b.Func("f", {}, b.ty.void_(), + { + b.Decl(b.Const("a", nullptr, + b.AddressOf(b.MemberAccessor("s", "a")))), + b.Decl(b.Const("b", nullptr, + b.AddressOf(b.Deref(b.AddressOf(b.Deref("a")))))), + b.Decl(b.Const("c", nullptr, b.Deref("b"))), + b.Decl(b.Const("d", nullptr, b.IndexAccessor(b.Deref("b"), 1))), + b.Assign(b.Deref("b"), b.Construct(b.ty.array(32), 1.0f, + 2.0f, 3.0f, 4.0f)), + b.Assign(b.IndexAccessor(b.Deref("b"), 1), 5.0f), + }, + { + b.Stage(ast::PipelineStage::kCompute), + b.WorkgroupSize(1), + }); + + auto* expect = + R"( +struct strided_arr { + @size(32) + el : f32; +} + +struct S { + a : array; +} + +@group(0) @binding(0) var s : S; + +@stage(compute) @workgroup_size(1) +fn f() { + let c = s.a; + let d = s.a[1].el; + s.a = array(strided_arr(1.0), strided_arr(2.0), strided_arr(3.0), strided_arr(4.0)); + s.a[1].el = 5.0; +} +)"; + + auto got = Run( + Program(std::move(b))); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(DecomposeStridedArrayTest, PrivateAliasedStridedArray) { + // type ARR = @stride(32) array; + // struct S { + // a : ARR; + // }; + // @group(0) @binding(0) var s : S; + // + // @stage(compute) @workgroup_size(1) + // fn f() { + // let a : ARR = s.a; + // let b : f32 = s.a[1]; + // s.a = ARR(); + // s.a = ARR(1.0, 2.0, 3.0, 4.0); + // s.a[1] = 5.0; + // } + ProgramBuilder b; + b.Alias("ARR", b.ty.array(32)); + auto* S = b.Structure("S", {b.Member("a", b.ty.type_name("ARR"))}); + b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage, + ast::Access::kReadWrite, b.GroupAndBinding(0, 0)); + b.Func( + "f", {}, b.ty.void_(), + { + b.Decl( + b.Const("a", b.ty.type_name("ARR"), b.MemberAccessor("s", "a"))), + b.Decl(b.Const("b", b.ty.f32(), + b.IndexAccessor(b.MemberAccessor("s", "a"), 1))), + b.Assign(b.MemberAccessor("s", "a"), + b.Construct(b.ty.type_name("ARR"))), + b.Assign(b.MemberAccessor("s", "a"), + b.Construct(b.ty.type_name("ARR"), 1.0f, 2.0f, 3.0f, 4.0f)), + b.Assign(b.IndexAccessor(b.MemberAccessor("s", "a"), 1), 5.0f), + }, + { + b.Stage(ast::PipelineStage::kCompute), + b.WorkgroupSize(1), + }); + + auto* expect = R"( +struct strided_arr { + @size(32) + el : f32; +} + +type ARR = array; + +struct S { + a : ARR; +} + +@group(0) @binding(0) var s : S; + +@stage(compute) @workgroup_size(1) +fn f() { + let a : ARR = s.a; + let b : f32 = s.a[1].el; + s.a = ARR(); + s.a = ARR(strided_arr(1.0), strided_arr(2.0), strided_arr(3.0), strided_arr(4.0)); + s.a[1].el = 5.0; +} +)"; + + auto got = Run( + Program(std::move(b))); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(DecomposeStridedArrayTest, PrivateNestedStridedArray) { + // type ARR_A = @stride(8) array; + // type ARR_B = @stride(128) array<@stride(16) array, 4>; + // struct S { + // a : ARR_B; + // }; + // @group(0) @binding(0) var s : S; + // + // @stage(compute) @workgroup_size(1) + // fn f() { + // let a : ARR_B = s.a; + // let b : array<@stride(8) array, 3> = s.a[3]; + // let c = s.a[3][2]; + // let d = s.a[3][2][1]; + // s.a = ARR_B(); + // s.a[3][2][1] = 5.0; + // } + + ProgramBuilder b; + b.Alias("ARR_A", b.ty.array(8)); + b.Alias("ARR_B", + b.ty.array( // + b.ty.array(b.ty.type_name("ARR_A"), 3, 16), // + 4, 128)); + auto* S = b.Structure("S", {b.Member("a", b.ty.type_name("ARR_B"))}); + b.Global("s", b.ty.Of(S), ast::StorageClass::kStorage, + ast::Access::kReadWrite, b.GroupAndBinding(0, 0)); + b.Func("f", {}, b.ty.void_(), + { + b.Decl(b.Const("a", b.ty.type_name("ARR_B"), + b.MemberAccessor("s", "a"))), + b.Decl(b.Const("b", b.ty.array(b.ty.type_name("ARR_A"), 3, 16), + b.IndexAccessor( // + b.MemberAccessor("s", "a"), // + 3))), + b.Decl(b.Const("c", b.ty.type_name("ARR_A"), + b.IndexAccessor( // + b.IndexAccessor( // + b.MemberAccessor("s", "a"), // + 3), + 2))), + b.Decl(b.Const("d", b.ty.f32(), + b.IndexAccessor( // + b.IndexAccessor( // + b.IndexAccessor( // + b.MemberAccessor("s", "a"), // + 3), + 2), + 1))), + b.Assign(b.MemberAccessor("s", "a"), + b.Construct(b.ty.type_name("ARR_B"))), + b.Assign(b.IndexAccessor( // + b.IndexAccessor( // + b.IndexAccessor( // + b.MemberAccessor("s", "a"), // + 3), + 2), + 1), + 5.0f), + }, + { + b.Stage(ast::PipelineStage::kCompute), + b.WorkgroupSize(1), + }); + + auto* expect = + R"( +struct strided_arr { + @size(8) + el : f32; +} + +type ARR_A = array; + +struct strided_arr_1 { + @size(128) + el : array; +} + +type ARR_B = array; + +struct S { + a : ARR_B; +} + +@group(0) @binding(0) var s : S; + +@stage(compute) @workgroup_size(1) +fn f() { + let a : ARR_B = s.a; + let b : array = s.a[3].el; + let c : ARR_A = s.a[3].el[2]; + let d : f32 = s.a[3].el[2][1].el; + s.a = ARR_B(); + s.a[3].el[2][1].el = 5.0; +} +)"; + + auto got = Run( + Program(std::move(b))); + + EXPECT_EQ(expect, str(got)); +} +} // namespace +} // namespace transform +} // namespace tint diff --git a/src/transform/test_helper.h b/src/transform/test_helper.h index a09ccba66c..3b3004a4d8 100644 --- a/src/transform/test_helper.h +++ b/src/transform/test_helper.h @@ -81,6 +81,15 @@ class TransformTestBase : public BASE { return manager.Run(&program, data); } + /// @param program the input program + /// @param data the optional DataMap to pass to Transform::Run() + /// @return true if the transform should be run for the given input. + template + bool ShouldRun(Program&& program, const DataMap& data = {}) { + EXPECT_TRUE(program.IsValid()) << program.Diagnostics().str(); + return TRANSFORM().ShouldRun(&program, data); + } + /// @param in the input WGSL source /// @param data the optional DataMap to pass to Transform::Run() /// @return true if the transform should be run for the given input. @@ -88,8 +97,7 @@ class TransformTestBase : public BASE { bool ShouldRun(std::string in, const DataMap& data = {}) { auto file = std::make_unique("test", in); auto program = reader::wgsl::Parse(file.get()); - EXPECT_TRUE(program.IsValid()) << program.Diagnostics().str(); - return TRANSFORM().ShouldRun(&program, data); + return ShouldRun(std::move(program), data); } /// @param output the output of the transform diff --git a/test/BUILD.gn b/test/BUILD.gn index 538c4a9639..2eb530a26d 100644 --- a/test/BUILD.gn +++ b/test/BUILD.gn @@ -310,6 +310,7 @@ tint_unittests_source_set("tint_unittests_transform_src") { "../src/transform/canonicalize_entry_point_io_test.cc", "../src/transform/combine_samplers_test.cc", "../src/transform/decompose_memory_access_test.cc", + "../src/transform/decompose_strided_array_test.cc", "../src/transform/decompose_strided_matrix_test.cc", "../src/transform/external_texture_transform_test.cc", "../src/transform/first_index_offset_test.cc", diff --git a/test/array/strides.spvasm b/test/array/strides.spvasm new file mode 100644 index 0000000000..68fc329f93 --- /dev/null +++ b/test/array/strides.spvasm @@ -0,0 +1,71 @@ +; type ARR_A = @stride(8) array; +; type ARR_B = @stride(128) array<@stride(16) array, 3>; +; struct S { +; a : ARR_B; +; }; +; @group(0) @binding(0) var s : S; +; +; @stage(compute) @workgroup_size(1) +; fn f() { +; let a : ARR_B = s.a; +; let b : array<@stride(8) array, 3> = s.a[3]; +; let c = s.a[3][2]; +; let d = s.a[3][2][1]; +; s.a = ARR_B(); +; s.a[3][2][1] = 5.0; +; } + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %S "S" + OpMemberName %S 0 "a" + OpName %s "s" + OpName %f "f" + OpDecorate %S Block + OpMemberDecorate %S 0 Offset 0 + OpDecorate %_arr_float_uint_2 ArrayStride 8 + OpDecorate %_arr__arr_float_uint_2_uint_3 ArrayStride 16 + OpDecorate %_arr__arr__arr_float_uint_2_uint_3_uint_4 ArrayStride 128 + OpDecorate %s DescriptorSet 0 + OpDecorate %s Binding 0 + %float = OpTypeFloat 32 + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 + %_arr_float_uint_2 = OpTypeArray %float %uint_2 + %uint_3 = OpConstant %uint 3 + %_arr__arr_float_uint_2_uint_3 = OpTypeArray %_arr_float_uint_2 %uint_3 + %uint_4 = OpConstant %uint 4 + %_arr__arr__arr_float_uint_2_uint_3_uint_4 = OpTypeArray %_arr__arr_float_uint_2_uint_3 %uint_4 + %S = OpTypeStruct %_arr__arr__arr_float_uint_2_uint_3_uint_4 + %_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S + %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer + %void = OpTypeVoid + %12 = OpTypeFunction %void + %uint_0 = OpConstant %uint 0 +%_ptr_StorageBuffer__arr__arr__arr_float_uint_2_uint_3_uint_4 = OpTypePointer StorageBuffer %_arr__arr__arr_float_uint_2_uint_3_uint_4 + %int = OpTypeInt 32 1 + %int_3 = OpConstant %int 3 + %_ptr_StorageBuffer__arr__arr_float_uint_2_uint_3 = OpTypePointer StorageBuffer %_arr__arr_float_uint_2_uint_3 + %int_2 = OpConstant %int 2 + %_ptr_StorageBuffer__arr_float_uint_2 = OpTypePointer StorageBuffer %_arr_float_uint_2 + %int_1 = OpConstant %int 1 + %_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float + %34 = OpConstantNull %_arr__arr__arr_float_uint_2_uint_3_uint_4 + %float_5 = OpConstant %float 5 + %f = OpFunction %void None %12 + %15 = OpLabel + %18 = OpAccessChain %_ptr_StorageBuffer__arr__arr__arr_float_uint_2_uint_3_uint_4 %s %uint_0 + %19 = OpLoad %_arr__arr__arr_float_uint_2_uint_3_uint_4 %18 + %23 = OpAccessChain %_ptr_StorageBuffer__arr__arr_float_uint_2_uint_3 %s %uint_0 %int_3 + %24 = OpLoad %_arr__arr_float_uint_2_uint_3 %23 + %27 = OpAccessChain %_ptr_StorageBuffer__arr_float_uint_2 %s %uint_0 %int_3 %int_2 + %28 = OpLoad %_arr_float_uint_2 %27 + %31 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %int_3 %int_2 %int_1 + %32 = OpLoad %float %31 + %33 = OpAccessChain %_ptr_StorageBuffer__arr__arr__arr_float_uint_2_uint_3_uint_4 %s %uint_0 + OpStore %33 %34 + %35 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %int_3 %int_2 %int_1 + OpStore %35 %float_5 + OpReturn + OpFunctionEnd diff --git a/test/array/strides.spvasm.expected.glsl b/test/array/strides.spvasm.expected.glsl new file mode 100644 index 0000000000..21fc3e81a4 --- /dev/null +++ b/test/array/strides.spvasm.expected.glsl @@ -0,0 +1,38 @@ +#version 310 es +precision mediump float; + +struct strided_arr { + float el; +}; + +struct strided_arr_1 { + strided_arr el[3][2]; +}; + +struct S { + strided_arr_1 a[4]; +}; + +layout(binding = 0) buffer S_1 { + strided_arr_1 a[4]; +} s; +void f_1() { + strided_arr_1 x_19[4] = s.a; + strided_arr x_24[3][2] = s.a[3].el; + strided_arr x_28[2] = s.a[3].el[2]; + float x_32 = s.a[3].el[2][1].el; + strided_arr_1 tint_symbol[4] = strided_arr_1[4](strided_arr_1(strided_arr[3][2](strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)), strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)), strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)))), strided_arr_1(strided_arr[3][2](strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)), strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)), strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)))), strided_arr_1(strided_arr[3][2](strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)), strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)), strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)))), strided_arr_1(strided_arr[3][2](strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)), strided_arr[2](strided_arr(0.0f), strided_arr(0.0f)), strided_arr[2](strided_arr(0.0f), strided_arr(0.0f))))); + s.a = tint_symbol; + s.a[3].el[2][1].el = 5.0f; + return; +} + +void f() { + f_1(); +} + +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; +void main() { + f(); + return; +} diff --git a/test/array/strides.spvasm.expected.hlsl b/test/array/strides.spvasm.expected.hlsl new file mode 100644 index 0000000000..910a49a935 --- /dev/null +++ b/test/array/strides.spvasm.expected.hlsl @@ -0,0 +1,103 @@ +struct strided_arr { + float el; +}; +struct strided_arr_1 { + strided_arr el[3][2]; +}; + +RWByteAddressBuffer s : register(u0, space0); + +strided_arr tint_symbol_4(RWByteAddressBuffer buffer, uint offset) { + const strided_arr tint_symbol_12 = {asfloat(buffer.Load((offset + 0u)))}; + return tint_symbol_12; +} + +typedef strided_arr tint_symbol_3_ret[2]; +tint_symbol_3_ret tint_symbol_3(RWByteAddressBuffer buffer, uint offset) { + strided_arr arr[2] = (strided_arr[2])0; + { + [loop] for(uint i = 0u; (i < 2u); i = (i + 1u)) { + arr[i] = tint_symbol_4(buffer, (offset + (i * 8u))); + } + } + return arr; +} + +typedef strided_arr tint_symbol_2_ret[3][2]; +tint_symbol_2_ret tint_symbol_2(RWByteAddressBuffer buffer, uint offset) { + strided_arr arr_1[3][2] = (strided_arr[3][2])0; + { + [loop] for(uint i_1 = 0u; (i_1 < 3u); i_1 = (i_1 + 1u)) { + arr_1[i_1] = tint_symbol_3(buffer, (offset + (i_1 * 16u))); + } + } + return arr_1; +} + +strided_arr_1 tint_symbol_1(RWByteAddressBuffer buffer, uint offset) { + const strided_arr_1 tint_symbol_13 = {tint_symbol_2(buffer, (offset + 0u))}; + return tint_symbol_13; +} + +typedef strided_arr_1 tint_symbol_ret[4]; +tint_symbol_ret tint_symbol(RWByteAddressBuffer buffer, uint offset) { + strided_arr_1 arr_2[4] = (strided_arr_1[4])0; + { + [loop] for(uint i_2 = 0u; (i_2 < 4u); i_2 = (i_2 + 1u)) { + arr_2[i_2] = tint_symbol_1(buffer, (offset + (i_2 * 128u))); + } + } + return arr_2; +} + +void tint_symbol_10(RWByteAddressBuffer buffer, uint offset, strided_arr value) { + buffer.Store((offset + 0u), asuint(value.el)); +} + +void tint_symbol_9(RWByteAddressBuffer buffer, uint offset, strided_arr value[2]) { + strided_arr array_2[2] = value; + { + [loop] for(uint i_3 = 0u; (i_3 < 2u); i_3 = (i_3 + 1u)) { + tint_symbol_10(buffer, (offset + (i_3 * 8u)), array_2[i_3]); + } + } +} + +void tint_symbol_8(RWByteAddressBuffer buffer, uint offset, strided_arr value[3][2]) { + strided_arr array_1[3][2] = value; + { + [loop] for(uint i_4 = 0u; (i_4 < 3u); i_4 = (i_4 + 1u)) { + tint_symbol_9(buffer, (offset + (i_4 * 16u)), array_1[i_4]); + } + } +} + +void tint_symbol_7(RWByteAddressBuffer buffer, uint offset, strided_arr_1 value) { + tint_symbol_8(buffer, (offset + 0u), value.el); +} + +void tint_symbol_6(RWByteAddressBuffer buffer, uint offset, strided_arr_1 value[4]) { + strided_arr_1 array[4] = value; + { + [loop] for(uint i_5 = 0u; (i_5 < 4u); i_5 = (i_5 + 1u)) { + tint_symbol_7(buffer, (offset + (i_5 * 128u)), array[i_5]); + } + } +} + +void f_1() { + const strided_arr_1 x_19[4] = tint_symbol(s, 0u); + const strided_arr x_24[3][2] = tint_symbol_2(s, 384u); + const strided_arr x_28[2] = tint_symbol_3(s, 416u); + const float x_32 = asfloat(s.Load(424u)); + const strided_arr_1 tint_symbol_14[4] = (strided_arr_1[4])0; + tint_symbol_6(s, 0u, tint_symbol_14); + s.Store(424u, asuint(5.0f)); + return; +} + +[numthreads(1, 1, 1)] +void f() { + f_1(); + return; +} diff --git a/test/array/strides.spvasm.expected.msl b/test/array/strides.spvasm.expected.msl new file mode 100644 index 0000000000..62f817c85b --- /dev/null +++ b/test/array/strides.spvasm.expected.msl @@ -0,0 +1,40 @@ +#include + +using namespace metal; +struct strided_arr { + /* 0x0000 */ float el; + /* 0x0004 */ int8_t tint_pad[4]; +}; +struct tint_array_wrapper { + /* 0x0000 */ strided_arr arr[2]; +}; +struct tint_array_wrapper_1 { + /* 0x0000 */ tint_array_wrapper arr[3]; +}; +struct strided_arr_1 { + /* 0x0000 */ tint_array_wrapper_1 el; + /* 0x0030 */ int8_t tint_pad_1[80]; +}; +struct tint_array_wrapper_2 { + /* 0x0000 */ strided_arr_1 arr[4]; +}; +struct S { + /* 0x0000 */ tint_array_wrapper_2 a; +}; + +void f_1(device S* const tint_symbol_1) { + tint_array_wrapper_2 const x_19 = (*(tint_symbol_1)).a; + tint_array_wrapper_1 const x_24 = (*(tint_symbol_1)).a.arr[3].el; + tint_array_wrapper const x_28 = (*(tint_symbol_1)).a.arr[3].el.arr[2]; + float const x_32 = (*(tint_symbol_1)).a.arr[3].el.arr[2].arr[1].el; + tint_array_wrapper_2 const tint_symbol = {.arr={}}; + (*(tint_symbol_1)).a = tint_symbol; + (*(tint_symbol_1)).a.arr[3].el.arr[2].arr[1].el = 5.0f; + return; +} + +kernel void f(device S* tint_symbol_2 [[buffer(0)]]) { + f_1(tint_symbol_2); + return; +} + diff --git a/test/array/strides.spvasm.expected.spvasm b/test/array/strides.spvasm.expected.spvasm new file mode 100644 index 0000000000..3108b19428 --- /dev/null +++ b/test/array/strides.spvasm.expected.spvasm @@ -0,0 +1,74 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 42 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %f "f" + OpExecutionMode %f LocalSize 1 1 1 + OpName %S "S" + OpMemberName %S 0 "a" + OpName %strided_arr_1 "strided_arr_1" + OpMemberName %strided_arr_1 0 "el" + OpName %strided_arr "strided_arr" + OpMemberName %strided_arr 0 "el" + OpName %s "s" + OpName %f_1 "f_1" + OpName %f "f" + OpDecorate %S Block + OpMemberDecorate %S 0 Offset 0 + OpMemberDecorate %strided_arr_1 0 Offset 0 + OpMemberDecorate %strided_arr 0 Offset 0 + OpDecorate %_arr_strided_arr_uint_2 ArrayStride 8 + OpDecorate %_arr__arr_strided_arr_uint_2_uint_3 ArrayStride 16 + OpDecorate %_arr_strided_arr_1_uint_4 ArrayStride 128 + OpDecorate %s DescriptorSet 0 + OpDecorate %s Binding 0 + %float = OpTypeFloat 32 +%strided_arr = OpTypeStruct %float + %uint = OpTypeInt 32 0 + %uint_2 = OpConstant %uint 2 +%_arr_strided_arr_uint_2 = OpTypeArray %strided_arr %uint_2 + %uint_3 = OpConstant %uint 3 +%_arr__arr_strided_arr_uint_2_uint_3 = OpTypeArray %_arr_strided_arr_uint_2 %uint_3 +%strided_arr_1 = OpTypeStruct %_arr__arr_strided_arr_uint_2_uint_3 + %uint_4 = OpConstant %uint 4 +%_arr_strided_arr_1_uint_4 = OpTypeArray %strided_arr_1 %uint_4 + %S = OpTypeStruct %_arr_strided_arr_1_uint_4 +%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S + %s = OpVariable %_ptr_StorageBuffer_S StorageBuffer + %void = OpTypeVoid + %14 = OpTypeFunction %void + %uint_0 = OpConstant %uint 0 +%_ptr_StorageBuffer__arr_strided_arr_1_uint_4 = OpTypePointer StorageBuffer %_arr_strided_arr_1_uint_4 + %int = OpTypeInt 32 1 + %int_3 = OpConstant %int 3 +%_ptr_StorageBuffer__arr__arr_strided_arr_uint_2_uint_3 = OpTypePointer StorageBuffer %_arr__arr_strided_arr_uint_2_uint_3 + %int_2 = OpConstant %int 2 +%_ptr_StorageBuffer__arr_strided_arr_uint_2 = OpTypePointer StorageBuffer %_arr_strided_arr_uint_2 + %int_1 = OpConstant %int 1 +%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float + %36 = OpConstantNull %_arr_strided_arr_1_uint_4 + %float_5 = OpConstant %float 5 + %f_1 = OpFunction %void None %14 + %17 = OpLabel + %20 = OpAccessChain %_ptr_StorageBuffer__arr_strided_arr_1_uint_4 %s %uint_0 + %21 = OpLoad %_arr_strided_arr_1_uint_4 %20 + %25 = OpAccessChain %_ptr_StorageBuffer__arr__arr_strided_arr_uint_2_uint_3 %s %uint_0 %int_3 %uint_0 + %26 = OpLoad %_arr__arr_strided_arr_uint_2_uint_3 %25 + %29 = OpAccessChain %_ptr_StorageBuffer__arr_strided_arr_uint_2 %s %uint_0 %int_3 %uint_0 %int_2 + %30 = OpLoad %_arr_strided_arr_uint_2 %29 + %33 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %int_3 %uint_0 %int_2 %int_1 %uint_0 + %34 = OpLoad %float %33 + %35 = OpAccessChain %_ptr_StorageBuffer__arr_strided_arr_1_uint_4 %s %uint_0 + OpStore %35 %36 + %37 = OpAccessChain %_ptr_StorageBuffer_float %s %uint_0 %int_3 %uint_0 %int_2 %int_1 %uint_0 + OpStore %37 %float_5 + OpReturn + OpFunctionEnd + %f = OpFunction %void None %14 + %40 = OpLabel + %41 = OpFunctionCall %void %f_1 + OpReturn + OpFunctionEnd diff --git a/test/array/strides.spvasm.expected.wgsl b/test/array/strides.spvasm.expected.wgsl new file mode 100644 index 0000000000..799ce0d870 --- /dev/null +++ b/test/array/strides.spvasm.expected.wgsl @@ -0,0 +1,36 @@ +struct strided_arr { + @size(8) + el : f32; +} + +type Arr = array; + +type Arr_1 = array; + +struct strided_arr_1 { + @size(128) + el : Arr_1; +} + +type Arr_2 = array; + +struct S { + a : Arr_2; +} + +@group(0) @binding(0) var s : S; + +fn f_1() { + let x_19 : Arr_2 = s.a; + let x_24 : Arr_1 = s.a[3].el; + let x_28 : Arr = s.a[3].el[2]; + let x_32 : f32 = s.a[3].el[2][1].el; + s.a = array(); + s.a[3].el[2][1].el = 5.0; + return; +} + +@stage(compute) @workgroup_size(1, 1, 1) +fn f() { + f_1(); +} diff --git a/test/bug/tint/1088.spvasm.expected.glsl b/test/bug/tint/1088.spvasm.expected.glsl index a94d3917d6..dea6c16a5a 100644 --- a/test/bug/tint/1088.spvasm.expected.glsl +++ b/test/bug/tint/1088.spvasm.expected.glsl @@ -5,7 +5,7 @@ layout(location = 0) in vec3 position_param_1; layout(location = 2) in vec2 uv_param_1; layout(location = 1) in vec3 normal_param_1; layout(location = 0) out vec2 vUV_1_1; -struct tint_padded_array_element { +struct strided_arr { float el; }; @@ -13,7 +13,7 @@ struct LeftOver { mat4 worldViewProjection; float time; mat4 test2[2]; - tint_padded_array_element test[4]; + strided_arr test[4]; }; vec3 position = vec3(0.0f, 0.0f, 0.0f); @@ -21,7 +21,7 @@ layout(binding = 2) uniform LeftOver_1 { mat4 worldViewProjection; float time; mat4 test2[2]; - tint_padded_array_element test[4]; + strided_arr test[4]; } x_14; vec2 vUV = vec2(0.0f, 0.0f); diff --git a/test/bug/tint/1088.spvasm.expected.msl b/test/bug/tint/1088.spvasm.expected.msl index c0dd6507a5..1f074d5da4 100644 --- a/test/bug/tint/1088.spvasm.expected.msl +++ b/test/bug/tint/1088.spvasm.expected.msl @@ -4,12 +4,12 @@ using namespace metal; struct tint_array_wrapper { /* 0x0000 */ float4x4 arr[2]; }; -struct tint_padded_array_element { +struct strided_arr { /* 0x0000 */ float el; /* 0x0004 */ int8_t tint_pad[12]; }; struct tint_array_wrapper_1 { - /* 0x0000 */ tint_padded_array_element arr[4]; + /* 0x0000 */ strided_arr arr[4]; }; struct LeftOver { /* 0x0000 */ float4x4 worldViewProjection; diff --git a/test/bug/tint/1088.spvasm.expected.spvasm b/test/bug/tint/1088.spvasm.expected.spvasm index 721a011a1d..f157cd1751 100644 --- a/test/bug/tint/1088.spvasm.expected.spvasm +++ b/test/bug/tint/1088.spvasm.expected.spvasm @@ -1,10 +1,10 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 121 +; Bound: 122 ; Schema: 0 OpCapability Shader - %74 = OpExtInstImport "GLSL.std.450" + %75 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 OpEntryPoint Vertex %main "main" %position_param_1 %uv_param_1 %normal_param_1 %gl_Position_1 %vUV_1_1 %vertex_point_size OpName %position_param_1 "position_param_1" @@ -19,6 +19,8 @@ OpMemberName %LeftOver 1 "time" OpMemberName %LeftOver 2 "test2" OpMemberName %LeftOver 3 "test" + OpName %strided_arr "strided_arr" + OpMemberName %strided_arr 0 "el" OpName %x_14 "x_14" OpName %vUV "vUV" OpName %uv "uv" @@ -51,7 +53,8 @@ OpMemberDecorate %LeftOver 2 MatrixStride 16 OpDecorate %_arr_mat4v4float_uint_2 ArrayStride 64 OpMemberDecorate %LeftOver 3 Offset 208 - OpDecorate %_arr_float_uint_4 ArrayStride 16 + OpMemberDecorate %strided_arr 0 Offset 0 + OpDecorate %_arr_strided_arr_uint_4 ArrayStride 16 OpDecorate %x_14 NonWritable OpDecorate %x_14 DescriptorSet 2 OpDecorate %x_14 Binding 2 @@ -82,9 +85,10 @@ %uint = OpTypeInt 32 0 %uint_2 = OpConstant %uint 2 %_arr_mat4v4float_uint_2 = OpTypeArray %mat4v4float %uint_2 +%strided_arr = OpTypeStruct %float %uint_4 = OpConstant %uint 4 -%_arr_float_uint_4 = OpTypeArray %float %uint_4 - %LeftOver = OpTypeStruct %mat4v4float %float %_arr_mat4v4float_uint_2 %_arr_float_uint_4 +%_arr_strided_arr_uint_4 = OpTypeArray %strided_arr %uint_4 + %LeftOver = OpTypeStruct %mat4v4float %float %_arr_mat4v4float_uint_2 %_arr_strided_arr_uint_4 %_ptr_Uniform_LeftOver = OpTypePointer Uniform %LeftOver %x_14 = OpVariable %_ptr_Uniform_LeftOver Uniform %_ptr_Private_v2float = OpTypePointer Private %v2float @@ -94,7 +98,7 @@ %_ptr_Private_v4float = OpTypePointer Private %v4float %gl_Position = OpVariable %_ptr_Private_v4float Private %12 %void = OpTypeVoid - %37 = OpTypeFunction %void + %38 = OpTypeFunction %void %_ptr_Function_v4float = OpTypePointer Function %v4float %_ptr_Function_v3float = OpTypePointer Function %v3float %float_1 = OpConstant %float 1 @@ -110,88 +114,88 @@ %_ptr_Uniform_mat4v4float = OpTypePointer Uniform %mat4v4float %float_n1 = OpConstant %float -1 %main_out = OpTypeStruct %v4float %v2float - %102 = OpTypeFunction %main_out %v3float %v2float %v3float - %main_1 = OpFunction %void None %37 - %40 = OpLabel + %103 = OpTypeFunction %main_out %v3float %v2float %v3float + %main_1 = OpFunction %void None %38 + %41 = OpLabel %q = OpVariable %_ptr_Function_v4float Function %12 %p = OpVariable %_ptr_Function_v3float Function %21 - %45 = OpLoad %v3float %position - %46 = OpCompositeExtract %float %45 0 - %47 = OpCompositeExtract %float %45 1 - %48 = OpCompositeExtract %float %45 2 - %50 = OpCompositeConstruct %v4float %46 %47 %48 %float_1 - OpStore %q %50 - %51 = OpLoad %v4float %q - %52 = OpCompositeExtract %float %51 0 - %53 = OpCompositeExtract %float %51 1 - %54 = OpCompositeExtract %float %51 2 - %55 = OpCompositeConstruct %v3float %52 %53 %54 - OpStore %p %55 - %58 = OpAccessChain %_ptr_Function_float %p %uint_0 - %59 = OpLoad %float %58 - %64 = OpAccessChain %_ptr_Uniform_float %x_14 %uint_3 %int_0 - %65 = OpLoad %float %64 - %68 = OpAccessChain %_ptr_Private_float %position %uint_1 - %69 = OpLoad %float %68 - %70 = OpAccessChain %_ptr_Uniform_float %x_14 %uint_1 - %71 = OpLoad %float %70 - %72 = OpAccessChain %_ptr_Function_float %p %uint_0 - %75 = OpFMul %float %65 %69 - %76 = OpFAdd %float %75 %71 - %73 = OpExtInst %float %74 Sin %76 - %77 = OpFAdd %float %59 %73 - OpStore %72 %77 - %78 = OpAccessChain %_ptr_Function_float %p %uint_1 - %79 = OpLoad %float %78 - %80 = OpAccessChain %_ptr_Uniform_float %x_14 %uint_1 - %81 = OpLoad %float %80 - %82 = OpAccessChain %_ptr_Function_float %p %uint_1 - %85 = OpFAdd %float %81 %float_4 - %83 = OpExtInst %float %74 Sin %85 - %86 = OpFAdd %float %79 %83 - OpStore %82 %86 - %88 = OpAccessChain %_ptr_Uniform_mat4v4float %x_14 %uint_0 - %89 = OpLoad %mat4v4float %88 - %90 = OpLoad %v3float %p - %91 = OpCompositeExtract %float %90 0 - %92 = OpCompositeExtract %float %90 1 - %93 = OpCompositeExtract %float %90 2 - %94 = OpCompositeConstruct %v4float %91 %92 %93 %float_1 - %95 = OpMatrixTimesVector %v4float %89 %94 - OpStore %gl_Position %95 - %96 = OpLoad %v2float %uv - OpStore %vUV %96 - %97 = OpAccessChain %_ptr_Private_float %gl_Position %uint_1 - %98 = OpLoad %float %97 - %99 = OpAccessChain %_ptr_Private_float %gl_Position %uint_1 - %101 = OpFMul %float %98 %float_n1 - OpStore %99 %101 + %46 = OpLoad %v3float %position + %47 = OpCompositeExtract %float %46 0 + %48 = OpCompositeExtract %float %46 1 + %49 = OpCompositeExtract %float %46 2 + %51 = OpCompositeConstruct %v4float %47 %48 %49 %float_1 + OpStore %q %51 + %52 = OpLoad %v4float %q + %53 = OpCompositeExtract %float %52 0 + %54 = OpCompositeExtract %float %52 1 + %55 = OpCompositeExtract %float %52 2 + %56 = OpCompositeConstruct %v3float %53 %54 %55 + OpStore %p %56 + %59 = OpAccessChain %_ptr_Function_float %p %uint_0 + %60 = OpLoad %float %59 + %65 = OpAccessChain %_ptr_Uniform_float %x_14 %uint_3 %int_0 %uint_0 + %66 = OpLoad %float %65 + %69 = OpAccessChain %_ptr_Private_float %position %uint_1 + %70 = OpLoad %float %69 + %71 = OpAccessChain %_ptr_Uniform_float %x_14 %uint_1 + %72 = OpLoad %float %71 + %73 = OpAccessChain %_ptr_Function_float %p %uint_0 + %76 = OpFMul %float %66 %70 + %77 = OpFAdd %float %76 %72 + %74 = OpExtInst %float %75 Sin %77 + %78 = OpFAdd %float %60 %74 + OpStore %73 %78 + %79 = OpAccessChain %_ptr_Function_float %p %uint_1 + %80 = OpLoad %float %79 + %81 = OpAccessChain %_ptr_Uniform_float %x_14 %uint_1 + %82 = OpLoad %float %81 + %83 = OpAccessChain %_ptr_Function_float %p %uint_1 + %86 = OpFAdd %float %82 %float_4 + %84 = OpExtInst %float %75 Sin %86 + %87 = OpFAdd %float %80 %84 + OpStore %83 %87 + %89 = OpAccessChain %_ptr_Uniform_mat4v4float %x_14 %uint_0 + %90 = OpLoad %mat4v4float %89 + %91 = OpLoad %v3float %p + %92 = OpCompositeExtract %float %91 0 + %93 = OpCompositeExtract %float %91 1 + %94 = OpCompositeExtract %float %91 2 + %95 = OpCompositeConstruct %v4float %92 %93 %94 %float_1 + %96 = OpMatrixTimesVector %v4float %90 %95 + OpStore %gl_Position %96 + %97 = OpLoad %v2float %uv + OpStore %vUV %97 + %98 = OpAccessChain %_ptr_Private_float %gl_Position %uint_1 + %99 = OpLoad %float %98 + %100 = OpAccessChain %_ptr_Private_float %gl_Position %uint_1 + %102 = OpFMul %float %99 %float_n1 + OpStore %100 %102 OpReturn OpFunctionEnd - %main_inner = OpFunction %main_out None %102 + %main_inner = OpFunction %main_out None %103 %position_param = OpFunctionParameter %v3float %uv_param = OpFunctionParameter %v2float %normal_param = OpFunctionParameter %v3float - %108 = OpLabel + %109 = OpLabel OpStore %position %position_param OpStore %uv %uv_param OpStore %normal %normal_param - %109 = OpFunctionCall %void %main_1 - %110 = OpLoad %v4float %gl_Position - %111 = OpLoad %v2float %vUV - %112 = OpCompositeConstruct %main_out %110 %111 - OpReturnValue %112 + %110 = OpFunctionCall %void %main_1 + %111 = OpLoad %v4float %gl_Position + %112 = OpLoad %v2float %vUV + %113 = OpCompositeConstruct %main_out %111 %112 + OpReturnValue %113 OpFunctionEnd - %main = OpFunction %void None %37 - %114 = OpLabel - %116 = OpLoad %v3float %position_param_1 - %117 = OpLoad %v2float %uv_param_1 - %118 = OpLoad %v3float %normal_param_1 - %115 = OpFunctionCall %main_out %main_inner %116 %117 %118 - %119 = OpCompositeExtract %v4float %115 0 - OpStore %gl_Position_1 %119 - %120 = OpCompositeExtract %v2float %115 1 - OpStore %vUV_1_1 %120 + %main = OpFunction %void None %38 + %115 = OpLabel + %117 = OpLoad %v3float %position_param_1 + %118 = OpLoad %v2float %uv_param_1 + %119 = OpLoad %v3float %normal_param_1 + %116 = OpFunctionCall %main_out %main_inner %117 %118 %119 + %120 = OpCompositeExtract %v4float %116 0 + OpStore %gl_Position_1 %120 + %121 = OpCompositeExtract %v2float %116 1 + OpStore %vUV_1_1 %121 OpStore %vertex_point_size %float_1 OpReturn OpFunctionEnd diff --git a/test/bug/tint/1088.spvasm.expected.wgsl b/test/bug/tint/1088.spvasm.expected.wgsl index c67da676d0..77f10638a9 100644 --- a/test/bug/tint/1088.spvasm.expected.wgsl +++ b/test/bug/tint/1088.spvasm.expected.wgsl @@ -1,6 +1,11 @@ -type Arr = @stride(64) array, 2u>; +type Arr = array, 2u>; -type Arr_1 = @stride(16) array; +struct strided_arr { + @size(16) + el : f32; +} + +type Arr_1 = array; struct LeftOver { worldViewProjection : mat4x4; @@ -31,7 +36,7 @@ fn main_1() { let x_21 : vec4 = q; p = vec3(x_21.x, x_21.y, x_21.z); let x_27 : f32 = p.x; - let x_41 : f32 = x_14.test[0]; + let x_41 : f32 = x_14.test[0].el; let x_45 : f32 = position.y; let x_49 : f32 = x_14.time; p.x = (x_27 + sin(((x_41 * x_45) + x_49))); diff --git a/test/bug/tint/870.spvasm.expected.wgsl b/test/bug/tint/870.spvasm.expected.wgsl index 4af74ac3d4..a04e265f31 100644 --- a/test/bug/tint/870.spvasm.expected.wgsl +++ b/test/bug/tint/870.spvasm.expected.wgsl @@ -1,4 +1,4 @@ -type Arr = @stride(4) array; +type Arr = array; struct sspp962805860buildInformationS { footprint : vec4; diff --git a/test/bug/tint/943.spvasm.expected.wgsl b/test/bug/tint/943.spvasm.expected.wgsl index ea2d6ade12..e587782691 100644 --- a/test/bug/tint/943.spvasm.expected.wgsl +++ b/test/bug/tint/943.spvasm.expected.wgsl @@ -14,15 +14,15 @@ struct Uniforms { outShapeStrides : vec2; } -type RTArr = @stride(4) array; +type RTArr = array; -type RTArr_1 = @stride(4) array; +type RTArr_1 = array; struct ssbOut { result : RTArr_1; } -type RTArr_2 = @stride(4) array; +type RTArr_2 = array; struct ssbA { A : RTArr_1; diff --git a/test/bug/tint/951.spvasm.expected.wgsl b/test/bug/tint/951.spvasm.expected.wgsl index 63fd79ea60..865bac8ef5 100644 --- a/test/bug/tint/951.spvasm.expected.wgsl +++ b/test/bug/tint/951.spvasm.expected.wgsl @@ -1,6 +1,6 @@ -type RTArr = @stride(4) array; +type RTArr = array; -type RTArr_1 = @stride(4) array; +type RTArr_1 = array; struct ssbOut { result : RTArr_1; diff --git a/test/bug/tint/977.spvasm.expected.wgsl b/test/bug/tint/977.spvasm.expected.wgsl index 8b5ad53aa0..ac33a715e1 100644 --- a/test/bug/tint/977.spvasm.expected.wgsl +++ b/test/bug/tint/977.spvasm.expected.wgsl @@ -1,12 +1,12 @@ -type RTArr = @stride(4) array; +type RTArr = array; -type RTArr_1 = @stride(4) array; +type RTArr_1 = array; struct ResultMatrix { numbers : RTArr_1; } -type RTArr_2 = @stride(4) array; +type RTArr_2 = array; struct FirstMatrix { numbers : RTArr_1; diff --git a/test/layout/storage/mat2x2/stride/16.spvasm.expected.glsl b/test/layout/storage/mat2x2/stride/16.spvasm.expected.glsl index aca2256026..e82307f88e 100644 --- a/test/layout/storage/mat2x2/stride/16.spvasm.expected.glsl +++ b/test/layout/storage/mat2x2/stride/16.spvasm.expected.glsl @@ -1,24 +1,26 @@ #version 310 es precision mediump float; -struct tint_padded_array_element { +struct strided_arr { vec2 el; }; struct SSBO { - tint_padded_array_element m[2]; + strided_arr m[2]; }; layout(binding = 0) buffer SSBO_1 { - tint_padded_array_element m[2]; + strided_arr m[2]; } ssbo; -mat2 arr_to_mat2x2_stride_16(tint_padded_array_element arr[2]) { +mat2 arr_to_mat2x2_stride_16(strided_arr arr[2]) { return mat2(arr[0u].el, arr[1u].el); } -tint_padded_array_element[2] mat2x2_stride_16_to_arr(mat2 mat) { - tint_padded_array_element tint_symbol[2] = tint_padded_array_element[2](tint_padded_array_element(mat[0u]), tint_padded_array_element(mat[1u])); - return tint_symbol; +strided_arr[2] mat2x2_stride_16_to_arr(mat2 mat) { + strided_arr tint_symbol = strided_arr(mat[0u]); + strided_arr tint_symbol_1 = strided_arr(mat[1u]); + strided_arr tint_symbol_2[2] = strided_arr[2](tint_symbol, tint_symbol_1); + return tint_symbol_2; } void f_1() { diff --git a/test/layout/storage/mat2x2/stride/16.spvasm.expected.hlsl b/test/layout/storage/mat2x2/stride/16.spvasm.expected.hlsl index 18bd54dcf5..7e88aa71be 100644 --- a/test/layout/storage/mat2x2/stride/16.spvasm.expected.hlsl +++ b/test/layout/storage/mat2x2/stride/16.spvasm.expected.hlsl @@ -1,42 +1,53 @@ -struct tint_padded_array_element { +struct strided_arr { float2 el; }; RWByteAddressBuffer ssbo : register(u0, space0); -float2x2 arr_to_mat2x2_stride_16(tint_padded_array_element arr[2]) { +float2x2 arr_to_mat2x2_stride_16(strided_arr arr[2]) { return float2x2(arr[0u].el, arr[1u].el); } -typedef tint_padded_array_element mat2x2_stride_16_to_arr_ret[2]; +typedef strided_arr mat2x2_stride_16_to_arr_ret[2]; mat2x2_stride_16_to_arr_ret mat2x2_stride_16_to_arr(float2x2 mat) { - const tint_padded_array_element tint_symbol_4[2] = {{mat[0u]}, {mat[1u]}}; - return tint_symbol_4; + const strided_arr tint_symbol_6 = {mat[0u]}; + const strided_arr tint_symbol_7 = {mat[1u]}; + const strided_arr tint_symbol_8[2] = {tint_symbol_6, tint_symbol_7}; + return tint_symbol_8; } -typedef tint_padded_array_element tint_symbol_ret[2]; +strided_arr tint_symbol_1(RWByteAddressBuffer buffer, uint offset) { + const strided_arr tint_symbol_9 = {asfloat(buffer.Load2((offset + 0u)))}; + return tint_symbol_9; +} + +typedef strided_arr tint_symbol_ret[2]; tint_symbol_ret tint_symbol(RWByteAddressBuffer buffer, uint offset) { - tint_padded_array_element arr_1[2] = (tint_padded_array_element[2])0; + strided_arr arr_1[2] = (strided_arr[2])0; { [loop] for(uint i = 0u; (i < 2u); i = (i + 1u)) { - arr_1[i].el = asfloat(buffer.Load2((offset + (i * 16u)))); + arr_1[i] = tint_symbol_1(buffer, (offset + (i * 16u))); } } return arr_1; } -void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, tint_padded_array_element value[2]) { - tint_padded_array_element array[2] = value; +void tint_symbol_4(RWByteAddressBuffer buffer, uint offset, strided_arr value) { + buffer.Store2((offset + 0u), asuint(value.el)); +} + +void tint_symbol_3(RWByteAddressBuffer buffer, uint offset, strided_arr value[2]) { + strided_arr array[2] = value; { [loop] for(uint i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) { - buffer.Store2((offset + (i_1 * 16u)), asuint(array[i_1].el)); + tint_symbol_4(buffer, (offset + (i_1 * 16u)), array[i_1]); } } } void f_1() { const float2x2 x_15 = arr_to_mat2x2_stride_16(tint_symbol(ssbo, 0u)); - tint_symbol_2(ssbo, 0u, mat2x2_stride_16_to_arr(x_15)); + tint_symbol_3(ssbo, 0u, mat2x2_stride_16_to_arr(x_15)); return; } diff --git a/test/layout/storage/mat2x2/stride/16.spvasm.expected.msl b/test/layout/storage/mat2x2/stride/16.spvasm.expected.msl index ec498f464e..6a0ca3e1cd 100644 --- a/test/layout/storage/mat2x2/stride/16.spvasm.expected.msl +++ b/test/layout/storage/mat2x2/stride/16.spvasm.expected.msl @@ -1,12 +1,12 @@ #include using namespace metal; -struct tint_padded_array_element { +struct strided_arr { /* 0x0000 */ float2 el; /* 0x0008 */ int8_t tint_pad[8]; }; struct tint_array_wrapper { - /* 0x0000 */ tint_padded_array_element arr[2]; + /* 0x0000 */ strided_arr arr[2]; }; struct SSBO { /* 0x0000 */ tint_array_wrapper m; @@ -17,18 +17,20 @@ float2x2 arr_to_mat2x2_stride_16(tint_array_wrapper arr) { } tint_array_wrapper mat2x2_stride_16_to_arr(float2x2 mat) { - tint_array_wrapper const tint_symbol = {.arr={{.el=mat[0u]}, {.el=mat[1u]}}}; - return tint_symbol; + strided_arr const tint_symbol = {.el=mat[0u]}; + strided_arr const tint_symbol_1 = {.el=mat[1u]}; + tint_array_wrapper const tint_symbol_2 = {.arr={tint_symbol, tint_symbol_1}}; + return tint_symbol_2; } -void f_1(device SSBO* const tint_symbol_1) { - float2x2 const x_15 = arr_to_mat2x2_stride_16((*(tint_symbol_1)).m); - (*(tint_symbol_1)).m = mat2x2_stride_16_to_arr(x_15); +void f_1(device SSBO* const tint_symbol_3) { + float2x2 const x_15 = arr_to_mat2x2_stride_16((*(tint_symbol_3)).m); + (*(tint_symbol_3)).m = mat2x2_stride_16_to_arr(x_15); return; } -kernel void f(device SSBO* tint_symbol_2 [[buffer(0)]]) { - f_1(tint_symbol_2); +kernel void f(device SSBO* tint_symbol_4 [[buffer(0)]]) { + f_1(tint_symbol_4); return; } diff --git a/test/layout/storage/mat2x2/stride/16.spvasm.expected.spvasm b/test/layout/storage/mat2x2/stride/16.spvasm.expected.spvasm index 94c280b7cd..f5ee801824 100644 --- a/test/layout/storage/mat2x2/stride/16.spvasm.expected.spvasm +++ b/test/layout/storage/mat2x2/stride/16.spvasm.expected.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.3 ; Generator: Google Tint Compiler; 0 -; Bound: 39 +; Bound: 44 ; Schema: 0 OpCapability Shader OpMemoryModel Logical GLSL450 @@ -9,6 +9,8 @@ OpExecutionMode %f LocalSize 1 1 1 OpName %SSBO "SSBO" OpMemberName %SSBO 0 "m" + OpName %strided_arr "strided_arr" + OpMemberName %strided_arr 0 "el" OpName %ssbo "ssbo" OpName %arr_to_mat2x2_stride_16 "arr_to_mat2x2_stride_16" OpName %arr "arr" @@ -18,53 +20,59 @@ OpName %f "f" OpDecorate %SSBO Block OpMemberDecorate %SSBO 0 Offset 0 - OpDecorate %_arr_v2float_uint_2 ArrayStride 16 + OpMemberDecorate %strided_arr 0 Offset 0 + OpDecorate %_arr_strided_arr_uint_2 ArrayStride 16 OpDecorate %ssbo DescriptorSet 0 OpDecorate %ssbo Binding 0 %float = OpTypeFloat 32 %v2float = OpTypeVector %float 2 +%strided_arr = OpTypeStruct %v2float %uint = OpTypeInt 32 0 %uint_2 = OpConstant %uint 2 -%_arr_v2float_uint_2 = OpTypeArray %v2float %uint_2 - %SSBO = OpTypeStruct %_arr_v2float_uint_2 +%_arr_strided_arr_uint_2 = OpTypeArray %strided_arr %uint_2 + %SSBO = OpTypeStruct %_arr_strided_arr_uint_2 %_ptr_StorageBuffer_SSBO = OpTypePointer StorageBuffer %SSBO %ssbo = OpVariable %_ptr_StorageBuffer_SSBO StorageBuffer %mat2v2float = OpTypeMatrix %v2float 2 - %9 = OpTypeFunction %mat2v2float %_arr_v2float_uint_2 + %10 = OpTypeFunction %mat2v2float %_arr_strided_arr_uint_2 %uint_0 = OpConstant %uint 0 %uint_1 = OpConstant %uint 1 - %19 = OpTypeFunction %_arr_v2float_uint_2 %mat2v2float + %22 = OpTypeFunction %_arr_strided_arr_uint_2 %mat2v2float %void = OpTypeVoid - %26 = OpTypeFunction %void -%_ptr_StorageBuffer__arr_v2float_uint_2 = OpTypePointer StorageBuffer %_arr_v2float_uint_2 -%arr_to_mat2x2_stride_16 = OpFunction %mat2v2float None %9 - %arr = OpFunctionParameter %_arr_v2float_uint_2 - %13 = OpLabel - %15 = OpCompositeExtract %v2float %arr 0 - %17 = OpCompositeExtract %v2float %arr 1 - %18 = OpCompositeConstruct %mat2v2float %15 %17 - OpReturnValue %18 + %31 = OpTypeFunction %void +%_ptr_StorageBuffer__arr_strided_arr_uint_2 = OpTypePointer StorageBuffer %_arr_strided_arr_uint_2 +%arr_to_mat2x2_stride_16 = OpFunction %mat2v2float None %10 + %arr = OpFunctionParameter %_arr_strided_arr_uint_2 + %14 = OpLabel + %16 = OpCompositeExtract %strided_arr %arr 0 + %17 = OpCompositeExtract %v2float %16 0 + %19 = OpCompositeExtract %strided_arr %arr 1 + %20 = OpCompositeExtract %v2float %19 0 + %21 = OpCompositeConstruct %mat2v2float %17 %20 + OpReturnValue %21 OpFunctionEnd -%mat2x2_stride_16_to_arr = OpFunction %_arr_v2float_uint_2 None %19 +%mat2x2_stride_16_to_arr = OpFunction %_arr_strided_arr_uint_2 None %22 %mat = OpFunctionParameter %mat2v2float - %22 = OpLabel - %23 = OpCompositeExtract %v2float %mat 0 - %24 = OpCompositeExtract %v2float %mat 1 - %25 = OpCompositeConstruct %_arr_v2float_uint_2 %23 %24 - OpReturnValue %25 + %25 = OpLabel + %26 = OpCompositeExtract %v2float %mat 0 + %27 = OpCompositeConstruct %strided_arr %26 + %28 = OpCompositeExtract %v2float %mat 1 + %29 = OpCompositeConstruct %strided_arr %28 + %30 = OpCompositeConstruct %_arr_strided_arr_uint_2 %27 %29 + OpReturnValue %30 OpFunctionEnd - %f_1 = OpFunction %void None %26 - %29 = OpLabel - %32 = OpAccessChain %_ptr_StorageBuffer__arr_v2float_uint_2 %ssbo %uint_0 - %33 = OpLoad %_arr_v2float_uint_2 %32 - %30 = OpFunctionCall %mat2v2float %arr_to_mat2x2_stride_16 %33 - %34 = OpAccessChain %_ptr_StorageBuffer__arr_v2float_uint_2 %ssbo %uint_0 - %35 = OpFunctionCall %_arr_v2float_uint_2 %mat2x2_stride_16_to_arr %30 - OpStore %34 %35 + %f_1 = OpFunction %void None %31 + %34 = OpLabel + %37 = OpAccessChain %_ptr_StorageBuffer__arr_strided_arr_uint_2 %ssbo %uint_0 + %38 = OpLoad %_arr_strided_arr_uint_2 %37 + %35 = OpFunctionCall %mat2v2float %arr_to_mat2x2_stride_16 %38 + %39 = OpAccessChain %_ptr_StorageBuffer__arr_strided_arr_uint_2 %ssbo %uint_0 + %40 = OpFunctionCall %_arr_strided_arr_uint_2 %mat2x2_stride_16_to_arr %35 + OpStore %39 %40 OpReturn OpFunctionEnd - %f = OpFunction %void None %26 - %37 = OpLabel - %38 = OpFunctionCall %void %f_1 + %f = OpFunction %void None %31 + %42 = OpLabel + %43 = OpFunctionCall %void %f_1 OpReturn OpFunctionEnd diff --git a/test/layout/storage/mat2x2/stride/16.spvasm.expected.wgsl b/test/layout/storage/mat2x2/stride/16.spvasm.expected.wgsl index a368ee1721..669f8cc489 100644 --- a/test/layout/storage/mat2x2/stride/16.spvasm.expected.wgsl +++ b/test/layout/storage/mat2x2/stride/16.spvasm.expected.wgsl @@ -1,15 +1,20 @@ +struct strided_arr { + @size(16) + el : vec2; +} + struct SSBO { - m : @stride(16) array, 2u>; + m : array; } @group(0) @binding(0) var ssbo : SSBO; -fn arr_to_mat2x2_stride_16(arr : @stride(16) array, 2u>) -> mat2x2 { - return mat2x2(arr[0u], arr[1u]); +fn arr_to_mat2x2_stride_16(arr : array) -> mat2x2 { + return mat2x2(arr[0u].el, arr[1u].el); } -fn mat2x2_stride_16_to_arr(mat : mat2x2) -> @stride(16) array, 2u> { - return @stride(16) array, 2u>(mat[0u], mat[1u]); +fn mat2x2_stride_16_to_arr(mat : mat2x2) -> array { + return array(strided_arr(mat[0u]), strided_arr(mat[1u])); } fn f_1() {