reader/spirv: Decompose arrays with strides
Transform any SPIR-V that has an array with a custom stride:
@stride(S) array<T, N>
into:
struct strided_arr {
@size(S) er : T;
};
array<strided_arr, N>
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 <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/src/BUILD.gn b/src/BUILD.gn
index 58b2f65..fce79f03 100644
--- a/src/BUILD.gn
+++ b/src/BUILD.gn
@@ -441,6 +441,8 @@
"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 600e91f..b0e0961 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -309,6 +309,8 @@
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 @@
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 2711733..e48357e 100644
--- a/src/reader/spirv/parser.cc
+++ b/src/reader/spirv/parser.cc
@@ -17,6 +17,7 @@
#include <utility>
#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 @@
manager.Add<transform::Unshadow>();
manager.Add<transform::SimplifyPointers>();
manager.Add<transform::DecomposeStridedMatrix>();
+ manager.Add<transform::DecomposeStridedArray>();
manager.Add<transform::RemoveUnreachableStatements>();
return manager.Run(&program).program;
}
diff --git a/src/sem/type_mappings.h b/src/sem/type_mappings.h
index 8d23a15..5dbc059 100644
--- a/src/sem/type_mappings.h
+++ b/src/sem/type_mappings.h
@@ -21,6 +21,7 @@
// Forward declarations
namespace ast {
+class Array;
class CallExpression;
class Expression;
class ElseStatement;
@@ -60,6 +61,7 @@
/// 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 0000000..106fa56
--- /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 <unordered_map>
+#include <utility>
+#include <vector>
+
+#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<const sem::Array*, Symbol>;
+
+} // 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<ast::Array>()) {
+ if (ast::GetAttribute<ast::StrideAttribute>(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<const sem::Array*, Symbol> 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::StrideAttribute>(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<sem::Array>()) {
+ 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<i32, 3>(1, 2, 3)`
+ // ->
+ // `array<strided_arr, 3>(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<sem::TypeConstructor>()) {
+ if (auto* arr = ctor->ReturnType()->As<sem::Array>()) {
+ // 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 0000000..27d4de0
--- /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<DecomposeStridedArray, Transform> {
+ 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 0000000..e982c9a
--- /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 <memory>
+#include <utility>
+#include <vector>
+
+#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<DecomposeStridedArray>(Program(std::move(b))));
+}
+
+TEST_F(DecomposeStridedArrayTest, ShouldRunNonStridedArray) {
+ // var<private> arr : array<f32, 4>
+
+ ProgramBuilder b;
+ b.Global("arr", b.ty.array<f32, 4>(), ast::StorageClass::kPrivate);
+ EXPECT_FALSE(ShouldRun<DecomposeStridedArray>(Program(std::move(b))));
+}
+
+TEST_F(DecomposeStridedArrayTest, ShouldRunDefaultStridedArray) {
+ // var<private> arr : @stride(4) array<f32, 4>
+
+ ProgramBuilder b;
+ b.Global("arr", b.ty.array<f32, 4>(4), ast::StorageClass::kPrivate);
+ EXPECT_TRUE(ShouldRun<DecomposeStridedArray>(Program(std::move(b))));
+}
+
+TEST_F(DecomposeStridedArrayTest, ShouldRunExplicitStridedArray) {
+ // var<private> arr : @stride(16) array<f32, 4>
+
+ ProgramBuilder b;
+ b.Global("arr", b.ty.array<f32, 4>(16), ast::StorageClass::kPrivate);
+ EXPECT_TRUE(ShouldRun<DecomposeStridedArray>(Program(std::move(b))));
+}
+
+TEST_F(DecomposeStridedArrayTest, Empty) {
+ auto* src = R"()";
+ auto* expect = src;
+
+ auto got = Run<DecomposeStridedArray>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(DecomposeStridedArrayTest, PrivateDefaultStridedArray) {
+ // var<private> arr : @stride(4) array<f32, 4>
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // let a : @stride(4) array<f32, 4> = a;
+ // let b : f32 = arr[1];
+ // }
+
+ ProgramBuilder b;
+ b.Global("arr", b.ty.array<f32, 4>(4), ast::StorageClass::kPrivate);
+ b.Func("f", {}, b.ty.void_(),
+ {
+ b.Decl(b.Const("a", b.ty.array<f32, 4>(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<private> arr : array<f32, 4>;
+
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a : array<f32, 4> = arr;
+ let b : f32 = arr[1];
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(
+ Program(std::move(b)));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(DecomposeStridedArrayTest, PrivateStridedArray) {
+ // var<private> arr : @stride(32) array<f32, 4>
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // let a : @stride(32) array<f32, 4> = a;
+ // let b : f32 = arr[1];
+ // }
+
+ ProgramBuilder b;
+ b.Global("arr", b.ty.array<f32, 4>(32), ast::StorageClass::kPrivate);
+ b.Func("f", {}, b.ty.void_(),
+ {
+ b.Decl(b.Const("a", b.ty.array<f32, 4>(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<private> arr : array<strided_arr, 4>;
+
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a : array<strided_arr, 4> = arr;
+ let b : f32 = arr[1].el;
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(
+ Program(std::move(b)));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(DecomposeStridedArrayTest, ReadUniformStridedArray) {
+ // struct S {
+ // a : @stride(32) array<f32, 4>;
+ // };
+ // @group(0) @binding(0) var<uniform> s : S;
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // let a : @stride(32) array<f32, 4> = s.a;
+ // let b : f32 = s.a[1];
+ // }
+ ProgramBuilder b;
+ auto* S = b.Structure("S", {b.Member("a", b.ty.array<f32, 4>(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<f32, 4>(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<strided_arr, 4>;
+}
+
+@group(0) @binding(0) var<uniform> s : S;
+
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a : array<strided_arr, 4> = s.a;
+ let b : f32 = s.a[1].el;
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(
+ Program(std::move(b)));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(DecomposeStridedArrayTest, ReadUniformDefaultStridedArray) {
+ // struct S {
+ // a : @stride(16) array<vec4<f32>, 4>;
+ // };
+ // @group(0) @binding(0) var<uniform> s : S;
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // let a : @stride(16) array<vec4<f32>, 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<f32>(), 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<f32>(), 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<vec4<f32>, 4>;
+}
+
+@group(0) @binding(0) var<uniform> s : S;
+
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a : array<vec4<f32>, 4> = s.a;
+ let b : f32 = s.a[1][2];
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(
+ Program(std::move(b)));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(DecomposeStridedArrayTest, ReadStorageStridedArray) {
+ // struct S {
+ // a : @stride(32) array<f32, 4>;
+ // };
+ // @group(0) @binding(0) var<storage> s : S;
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // let a : @stride(32) array<f32, 4> = s.a;
+ // let b : f32 = s.a[1];
+ // }
+ ProgramBuilder b;
+ auto* S = b.Structure("S", {b.Member("a", b.ty.array<f32, 4>(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<f32, 4>(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<strided_arr, 4>;
+}
+
+@group(0) @binding(0) var<storage> s : S;
+
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a : array<strided_arr, 4> = s.a;
+ let b : f32 = s.a[1].el;
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(
+ Program(std::move(b)));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(DecomposeStridedArrayTest, ReadStorageDefaultStridedArray) {
+ // struct S {
+ // a : @stride(4) array<f32, 4>;
+ // };
+ // @group(0) @binding(0) var<storage> s : S;
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // let a : @stride(4) array<f32, 4> = s.a;
+ // let b : f32 = s.a[1];
+ // }
+ ProgramBuilder b;
+ auto* S = b.Structure("S", {b.Member("a", b.ty.array<f32, 4>(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<f32, 4>(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<f32, 4>;
+}
+
+@group(0) @binding(0) var<storage> s : S;
+
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a : array<f32, 4> = s.a;
+ let b : f32 = s.a[1];
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(
+ Program(std::move(b)));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(DecomposeStridedArrayTest, WriteStorageStridedArray) {
+ // struct S {
+ // a : @stride(32) array<f32, 4>;
+ // };
+ // @group(0) @binding(0) var<storage, read_write> s : S;
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // s.a = @stride(32) array<f32, 4>();
+ // s.a = @stride(32) array<f32, 4>(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<f32, 4>(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<f32, 4>(32))),
+ b.Assign(b.MemberAccessor("s", "a"),
+ b.Construct(b.ty.array<f32, 4>(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<strided_arr, 4>;
+}
+
+@group(0) @binding(0) var<storage, read_write> s : S;
+
+@stage(compute) @workgroup_size(1)
+fn f() {
+ s.a = array<strided_arr, 4>();
+ s.a = array<strided_arr, 4>(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<Unshadow, SimplifyPointers, DecomposeStridedArray>(
+ Program(std::move(b)));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(DecomposeStridedArrayTest, WriteStorageDefaultStridedArray) {
+ // struct S {
+ // a : @stride(4) array<f32, 4>;
+ // };
+ // @group(0) @binding(0) var<storage, read_write> s : S;
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // s.a = @stride(4) array<f32, 4>();
+ // s.a = @stride(4) array<f32, 4>(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<f32, 4>(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<f32, 4>(4))),
+ b.Assign(b.MemberAccessor("s", "a"),
+ b.Construct(b.ty.array<f32, 4>(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<f32, 4>;
+}
+
+@group(0) @binding(0) var<storage, read_write> s : S;
+
+@stage(compute) @workgroup_size(1)
+fn f() {
+ s.a = array<f32, 4>();
+ s.a = array<f32, 4>(1.0, 2.0, 3.0, 4.0);
+ s.a[1] = 5.0;
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(
+ Program(std::move(b)));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(DecomposeStridedArrayTest, ReadWriteViaPointerLets) {
+ // struct S {
+ // a : @stride(32) array<f32, 4>;
+ // };
+ // @group(0) @binding(0) var<storage, read_write> 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<f32, 4>(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<f32, 4>(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<f32, 4>(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<strided_arr, 4>;
+}
+
+@group(0) @binding(0) var<storage, read_write> s : S;
+
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let c = s.a;
+ let d = s.a[1].el;
+ s.a = array<strided_arr, 4>(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<Unshadow, SimplifyPointers, DecomposeStridedArray>(
+ Program(std::move(b)));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(DecomposeStridedArrayTest, PrivateAliasedStridedArray) {
+ // type ARR = @stride(32) array<f32, 4>;
+ // struct S {
+ // a : ARR;
+ // };
+ // @group(0) @binding(0) var<storage, read_write> 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<f32, 4>(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<strided_arr, 4>;
+
+struct S {
+ a : ARR;
+}
+
+@group(0) @binding(0) var<storage, read_write> 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<Unshadow, SimplifyPointers, DecomposeStridedArray>(
+ Program(std::move(b)));
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(DecomposeStridedArrayTest, PrivateNestedStridedArray) {
+ // type ARR_A = @stride(8) array<f32, 2>;
+ // type ARR_B = @stride(128) array<@stride(16) array<ARR_A, 3>, 4>;
+ // struct S {
+ // a : ARR_B;
+ // };
+ // @group(0) @binding(0) var<storage, read_write> s : S;
+ //
+ // @stage(compute) @workgroup_size(1)
+ // fn f() {
+ // let a : ARR_B = s.a;
+ // let b : array<@stride(8) array<f32, 2>, 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<f32, 2>(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<strided_arr, 2>;
+
+struct strided_arr_1 {
+ @size(128)
+ el : array<ARR_A, 3>;
+}
+
+type ARR_B = array<strided_arr_1, 4>;
+
+struct S {
+ a : ARR_B;
+}
+
+@group(0) @binding(0) var<storage, read_write> s : S;
+
+@stage(compute) @workgroup_size(1)
+fn f() {
+ let a : ARR_B = s.a;
+ let b : array<ARR_A, 3> = 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<Unshadow, SimplifyPointers, DecomposeStridedArray>(
+ 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 a09ccba..3b3004a 100644
--- a/src/transform/test_helper.h
+++ b/src/transform/test_helper.h
@@ -81,6 +81,15 @@
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 <typename TRANSFORM>
+ 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 @@
bool ShouldRun(std::string in, const DataMap& data = {}) {
auto file = std::make_unique<Source::File>("test", in);
auto program = reader::wgsl::Parse(file.get());
- EXPECT_TRUE(program.IsValid()) << program.Diagnostics().str();
- return TRANSFORM().ShouldRun(&program, data);
+ return ShouldRun<TRANSFORM>(std::move(program), data);
}
/// @param output the output of the transform
diff --git a/test/BUILD.gn b/test/BUILD.gn
index 538c4a9..2eb530a 100644
--- a/test/BUILD.gn
+++ b/test/BUILD.gn
@@ -310,6 +310,7 @@
"../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 0000000..68fc329
--- /dev/null
+++ b/test/array/strides.spvasm
@@ -0,0 +1,71 @@
+; type ARR_A = @stride(8) array<f32, 2>;
+; type ARR_B = @stride(128) array<@stride(16) array<ARR_A, 4>, 3>;
+; struct S {
+; a : ARR_B;
+; };
+; @group(0) @binding(0) var<storage, read_write> s : S;
+;
+; @stage(compute) @workgroup_size(1)
+; fn f() {
+; let a : ARR_B = s.a;
+; let b : array<@stride(8) array<f32, 2>, 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 0000000..21fc3e8
--- /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 0000000..910a49a
--- /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 0000000..62f817c
--- /dev/null
+++ b/test/array/strides.spvasm.expected.msl
@@ -0,0 +1,40 @@
+#include <metal_stdlib>
+
+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 0000000..3108b19
--- /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 0000000..799ce0d
--- /dev/null
+++ b/test/array/strides.spvasm.expected.wgsl
@@ -0,0 +1,36 @@
+struct strided_arr {
+ @size(8)
+ el : f32;
+}
+
+type Arr = array<strided_arr, 2u>;
+
+type Arr_1 = array<Arr, 3u>;
+
+struct strided_arr_1 {
+ @size(128)
+ el : Arr_1;
+}
+
+type Arr_2 = array<strided_arr_1, 4u>;
+
+struct S {
+ a : Arr_2;
+}
+
+@group(0) @binding(0) var<storage, read_write> 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<strided_arr_1, 4u>();
+ 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 a94d391..dea6c16 100644
--- a/test/bug/tint/1088.spvasm.expected.glsl
+++ b/test/bug/tint/1088.spvasm.expected.glsl
@@ -5,7 +5,7 @@
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 @@
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 @@
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 c0dd650..1f074d5 100644
--- a/test/bug/tint/1088.spvasm.expected.msl
+++ b/test/bug/tint/1088.spvasm.expected.msl
@@ -4,12 +4,12 @@
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 721a011..f157cd1 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 c67da67..77f1063 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<mat4x4<f32>, 2u>;
+type Arr = array<mat4x4<f32>, 2u>;
-type Arr_1 = @stride(16) array<f32, 4u>;
+struct strided_arr {
+ @size(16)
+ el : f32;
+}
+
+type Arr_1 = array<strided_arr, 4u>;
struct LeftOver {
worldViewProjection : mat4x4<f32>;
@@ -31,7 +36,7 @@
let x_21 : vec4<f32> = q;
p = vec3<f32>(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 4af74ac..a04e265 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<i32, 6u>;
+type Arr = array<i32, 6u>;
struct sspp962805860buildInformationS {
footprint : vec4<f32>;
diff --git a/test/bug/tint/943.spvasm.expected.wgsl b/test/bug/tint/943.spvasm.expected.wgsl
index ea2d6ad..e587782 100644
--- a/test/bug/tint/943.spvasm.expected.wgsl
+++ b/test/bug/tint/943.spvasm.expected.wgsl
@@ -14,15 +14,15 @@
outShapeStrides : vec2<i32>;
}
-type RTArr = @stride(4) array<f32>;
+type RTArr = array<f32>;
-type RTArr_1 = @stride(4) array<f32>;
+type RTArr_1 = array<f32>;
struct ssbOut {
result : RTArr_1;
}
-type RTArr_2 = @stride(4) array<f32>;
+type RTArr_2 = array<f32>;
struct ssbA {
A : RTArr_1;
diff --git a/test/bug/tint/951.spvasm.expected.wgsl b/test/bug/tint/951.spvasm.expected.wgsl
index 63fd79e..865bac8 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<f32>;
+type RTArr = array<f32>;
-type RTArr_1 = @stride(4) array<f32>;
+type RTArr_1 = array<f32>;
struct ssbOut {
result : RTArr_1;
diff --git a/test/bug/tint/977.spvasm.expected.wgsl b/test/bug/tint/977.spvasm.expected.wgsl
index 8b5ad53..ac33a71 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<f32>;
+type RTArr = array<f32>;
-type RTArr_1 = @stride(4) array<f32>;
+type RTArr_1 = array<f32>;
struct ResultMatrix {
numbers : RTArr_1;
}
-type RTArr_2 = @stride(4) array<f32>;
+type RTArr_2 = array<f32>;
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 aca2256..e82307f 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 18bd54d..7e88aa7 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 ec498f4..6a0ca3e 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 <metal_stdlib>
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 @@
}
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 94c280b..f5ee801 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 a368ee1..669f8cc 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<f32>;
+}
+
struct SSBO {
- m : @stride(16) array<vec2<f32>, 2u>;
+ m : array<strided_arr, 2u>;
}
@group(0) @binding(0) var<storage, read_write> ssbo : SSBO;
-fn arr_to_mat2x2_stride_16(arr : @stride(16) array<vec2<f32>, 2u>) -> mat2x2<f32> {
- return mat2x2<f32>(arr[0u], arr[1u]);
+fn arr_to_mat2x2_stride_16(arr : array<strided_arr, 2u>) -> mat2x2<f32> {
+ return mat2x2<f32>(arr[0u].el, arr[1u].el);
}
-fn mat2x2_stride_16_to_arr(mat : mat2x2<f32>) -> @stride(16) array<vec2<f32>, 2u> {
- return @stride(16) array<vec2<f32>, 2u>(mat[0u], mat[1u]);
+fn mat2x2_stride_16_to_arr(mat : mat2x2<f32>) -> array<strided_arr, 2u> {
+ return array<strided_arr, 2u>(strided_arr(mat[0u]), strided_arr(mat[1u]));
}
fn f_1() {
