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>
This commit is contained in:
Ben Clayton 2022-02-04 15:39:34 +00:00 committed by Tint LUCI CQ
parent de857e1c58
commit 009d129103
28 changed files with 1499 additions and 161 deletions

View File

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

View File

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

View File

@ -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 @@ Program Parse(const std::vector<uint32_t>& input) {
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;
}

View File

@ -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*);

View File

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

View File

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

View File

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

View File

@ -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 <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 @@ class TransformTestBase : public BASE {
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

View File

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

71
test/array/strides.spvasm Normal file
View File

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

View File

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

View File

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

View File

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

View File

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

View File

@ -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();
}

View File

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

View File

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

View File

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

View File

@ -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 @@ fn main_1() {
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)));

View File

@ -1,4 +1,4 @@
type Arr = @stride(4) array<i32, 6u>;
type Arr = array<i32, 6u>;
struct sspp962805860buildInformationS {
footprint : vec4<f32>;

View File

@ -14,15 +14,15 @@ struct Uniforms {
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;

View File

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

View File

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

View File

@ -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() {

View File

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

View File

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

View File

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

View File

@ -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() {