blob: 54c00b70370c44fa3f0547f8e57dbbf9e9a0e004 [file] [log] [blame]
// 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/tint/transform/decompose_strided_array.h"
#include <memory>
#include <utility>
#include <vector>
#include "src/tint/program_builder.h"
#include "src/tint/transform/simplify_pointers.h"
#include "src/tint/transform/test_helper.h"
#include "src/tint/transform/unshadow.h"
using namespace tint::number_suffixes; // NOLINT
namespace tint::transform {
namespace {
using DecomposeStridedArrayTest = TransformTest;
TEST_F(DecomposeStridedArrayTest, ShouldRunEmptyModule) {
ProgramBuilder b;
EXPECT_FALSE(ShouldRun<DecomposeStridedArray>(Program(std::move(b))));
}
TEST_F(DecomposeStridedArrayTest, ShouldRunNonStridedArray) {
// var<private> arr : array<f32, 4u>
ProgramBuilder b;
b.GlobalVar("arr", b.ty.array<f32, 4u>(), builtin::AddressSpace::kPrivate);
EXPECT_FALSE(ShouldRun<DecomposeStridedArray>(Program(std::move(b))));
}
TEST_F(DecomposeStridedArrayTest, ShouldRunDefaultStridedArray) {
// var<private> arr : @stride(4) array<f32, 4u>
ProgramBuilder b;
b.GlobalVar("arr",
b.ty.array<f32, 4u>(utils::Vector{
b.Stride(4),
}),
builtin::AddressSpace::kPrivate);
EXPECT_TRUE(ShouldRun<DecomposeStridedArray>(Program(std::move(b))));
}
TEST_F(DecomposeStridedArrayTest, ShouldRunExplicitStridedArray) {
// var<private> arr : @stride(16) array<f32, 4u>
ProgramBuilder b;
b.GlobalVar("arr",
b.ty.array<f32, 4u>(utils::Vector{
b.Stride(16),
}),
builtin::AddressSpace::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, 4u>
//
// @compute @workgroup_size(1)
// fn f() {
// let a : @stride(4) array<f32, 4u> = a;
// let b : f32 = arr[1];
// }
ProgramBuilder b;
b.GlobalVar("arr",
b.ty.array<f32, 4u>(utils::Vector{
b.Stride(4),
}),
builtin::AddressSpace::kPrivate);
b.Func("f", utils::Empty, b.ty.void_(),
utils::Vector{
b.Decl(b.Let("a",
b.ty.array<f32, 4u>(utils::Vector{
b.Stride(4),
}),
b.Expr("arr"))),
b.Decl(b.Let("b", b.ty.f32(), b.IndexAccessor("arr", 1_i))),
},
utils::Vector{
b.Stage(ast::PipelineStage::kCompute),
b.WorkgroupSize(1_i),
});
auto* expect = R"(
var<private> arr : array<f32, 4u>;
@compute @workgroup_size(1i)
fn f() {
let a : array<f32, 4u> = arr;
let b : f32 = arr[1i];
}
)";
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, 4u>
//
// @compute @workgroup_size(1)
// fn f() {
// let a : @stride(32) array<f32, 4u> = a;
// let b : f32 = arr[1];
// }
ProgramBuilder b;
b.GlobalVar("arr",
b.ty.array<f32, 4u>(utils::Vector{
b.Stride(32),
}),
builtin::AddressSpace::kPrivate);
b.Func("f", utils::Empty, b.ty.void_(),
utils::Vector{
b.Decl(b.Let("a",
b.ty.array<f32, 4u>(utils::Vector{
b.Stride(32),
}),
b.Expr("arr"))),
b.Decl(b.Let("b", b.ty.f32(), b.IndexAccessor("arr", 1_i))),
},
utils::Vector{
b.Stage(ast::PipelineStage::kCompute),
b.WorkgroupSize(1_i),
});
auto* expect = R"(
struct strided_arr {
@size(32)
el : f32,
}
var<private> arr : array<strided_arr, 4u>;
@compute @workgroup_size(1i)
fn f() {
let a : array<strided_arr, 4u> = arr;
let b : f32 = arr[1i].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, 4u>,
// };
// @group(0) @binding(0) var<uniform> s : S;
//
// @compute @workgroup_size(1)
// fn f() {
// let a : @stride(32) array<f32, 4u> = s.a;
// let b : f32 = s.a[1];
// }
ProgramBuilder b;
auto* S = b.Structure("S", utils::Vector{b.Member("a", b.ty.array<f32, 4u>(utils::Vector{
b.Stride(32),
}))});
b.GlobalVar("s", b.ty.Of(S), builtin::AddressSpace::kUniform, b.Group(0_a), b.Binding(0_a));
b.Func("f", utils::Empty, b.ty.void_(),
utils::Vector{
b.Decl(b.Let("a",
b.ty.array<f32, 4u>(utils::Vector{
b.Stride(32),
}),
b.MemberAccessor("s", "a"))),
b.Decl(b.Let("b", b.ty.f32(), b.IndexAccessor(b.MemberAccessor("s", "a"), 1_i))),
},
utils::Vector{
b.Stage(ast::PipelineStage::kCompute),
b.WorkgroupSize(1_i),
});
auto* expect = R"(
struct strided_arr {
@size(32)
el : f32,
}
struct S {
a : array<strided_arr, 4u>,
}
@group(0) @binding(0) var<uniform> s : S;
@compute @workgroup_size(1i)
fn f() {
let a : array<strided_arr, 4u> = s.a;
let b : f32 = s.a[1i].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>, 4u>,
// };
// @group(0) @binding(0) var<uniform> s : S;
//
// @compute @workgroup_size(1)
// fn f() {
// let a : @stride(16) array<vec4<f32>, 4u> = s.a;
// let b : f32 = s.a[1][2];
// }
ProgramBuilder b;
auto* S = b.Structure("S", utils::Vector{b.Member("a", b.ty.array(b.ty.vec4<f32>(), 4_u,
utils::Vector{
b.Stride(16),
}))});
b.GlobalVar("s", b.ty.Of(S), builtin::AddressSpace::kUniform, b.Group(0_a), b.Binding(0_a));
b.Func(
"f", utils::Empty, b.ty.void_(),
utils::Vector{
b.Decl(b.Let("a",
b.ty.array(b.ty.vec4<f32>(), 4_u,
utils::Vector{
b.Stride(16),
}),
b.MemberAccessor("s", "a"))),
b.Decl(b.Let("b", b.ty.f32(),
b.IndexAccessor(b.IndexAccessor(b.MemberAccessor("s", "a"), 1_i), 2_i))),
},
utils::Vector{
b.Stage(ast::PipelineStage::kCompute),
b.WorkgroupSize(1_i),
});
auto* expect =
R"(
struct S {
a : array<vec4<f32>, 4u>,
}
@group(0) @binding(0) var<uniform> s : S;
@compute @workgroup_size(1i)
fn f() {
let a : array<vec4<f32>, 4u> = s.a;
let b : f32 = s.a[1i][2i];
}
)";
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, 4u>,
// };
// @group(0) @binding(0) var<storage> s : S;
//
// @compute @workgroup_size(1)
// fn f() {
// let a : @stride(32) array<f32, 4u> = s.a;
// let b : f32 = s.a[1];
// }
ProgramBuilder b;
auto* S = b.Structure("S", utils::Vector{b.Member("a", b.ty.array<f32, 4u>(utils::Vector{
b.Stride(32),
}))});
b.GlobalVar("s", b.ty.Of(S), builtin::AddressSpace::kStorage, b.Group(0_a), b.Binding(0_a));
b.Func("f", utils::Empty, b.ty.void_(),
utils::Vector{
b.Decl(b.Let("a",
b.ty.array<f32, 4u>(utils::Vector{
b.Stride(32),
}),
b.MemberAccessor("s", "a"))),
b.Decl(b.Let("b", b.ty.f32(), b.IndexAccessor(b.MemberAccessor("s", "a"), 1_i))),
},
utils::Vector{
b.Stage(ast::PipelineStage::kCompute),
b.WorkgroupSize(1_i),
});
auto* expect = R"(
struct strided_arr {
@size(32)
el : f32,
}
struct S {
a : array<strided_arr, 4u>,
}
@group(0) @binding(0) var<storage> s : S;
@compute @workgroup_size(1i)
fn f() {
let a : array<strided_arr, 4u> = s.a;
let b : f32 = s.a[1i].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, 4u>,
// };
// @group(0) @binding(0) var<storage> s : S;
//
// @compute @workgroup_size(1)
// fn f() {
// let a : @stride(4) array<f32, 4u> = s.a;
// let b : f32 = s.a[1];
// }
ProgramBuilder b;
auto* S = b.Structure("S", utils::Vector{b.Member("a", b.ty.array<f32, 4u>(utils::Vector{
b.Stride(4),
}))});
b.GlobalVar("s", b.ty.Of(S), builtin::AddressSpace::kStorage, b.Group(0_a), b.Binding(0_a));
b.Func("f", utils::Empty, b.ty.void_(),
utils::Vector{
b.Decl(b.Let("a",
b.ty.array<f32, 4u>(utils::Vector{
b.Stride(4),
}),
b.MemberAccessor("s", "a"))),
b.Decl(b.Let("b", b.ty.f32(), b.IndexAccessor(b.MemberAccessor("s", "a"), 1_i))),
},
utils::Vector{
b.Stage(ast::PipelineStage::kCompute),
b.WorkgroupSize(1_i),
});
auto* expect = R"(
struct S {
a : array<f32, 4u>,
}
@group(0) @binding(0) var<storage> s : S;
@compute @workgroup_size(1i)
fn f() {
let a : array<f32, 4u> = s.a;
let b : f32 = s.a[1i];
}
)";
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, 4u>,
// };
// @group(0) @binding(0) var<storage, read_write> s : S;
//
// @compute @workgroup_size(1)
// fn f() {
// s.a = @stride(32) array<f32, 4u>();
// s.a = @stride(32) array<f32, 4u>(1.0, 2.0, 3.0, 4.0);
// s.a[1i] = 5.0;
// }
ProgramBuilder b;
auto* S = b.Structure("S", utils::Vector{b.Member("a", b.ty.array<f32, 4u>(utils::Vector{
b.Stride(32),
}))});
b.GlobalVar("s", b.ty.Of(S), builtin::AddressSpace::kStorage, builtin::Access::kReadWrite,
b.Group(0_a), b.Binding(0_a));
b.Func("f", utils::Empty, b.ty.void_(),
utils::Vector{
b.Assign(b.MemberAccessor("s", "a"), b.Call(b.ty.array<f32, 4u>(utils::Vector{
b.Stride(32),
}))),
b.Assign(b.MemberAccessor("s", "a"), b.Call(b.ty.array<f32, 4u>(utils::Vector{
b.Stride(32),
}),
1_f, 2_f, 3_f, 4_f)),
b.Assign(b.IndexAccessor(b.MemberAccessor("s", "a"), 1_i), 5_f),
},
utils::Vector{
b.Stage(ast::PipelineStage::kCompute),
b.WorkgroupSize(1_i),
});
auto* expect =
R"(
struct strided_arr {
@size(32)
el : f32,
}
struct S {
a : array<strided_arr, 4u>,
}
@group(0) @binding(0) var<storage, read_write> s : S;
@compute @workgroup_size(1i)
fn f() {
s.a = array<strided_arr, 4u>();
s.a = array<strided_arr, 4u>(strided_arr(1.0f), strided_arr(2.0f), strided_arr(3.0f), strided_arr(4.0f));
s.a[1i].el = 5.0f;
}
)";
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, 4u>,
// };
// @group(0) @binding(0) var<storage, read_write> s : S;
//
// @compute @workgroup_size(1)
// fn f() {
// s.a = @stride(4) array<f32, 4u>();
// s.a = @stride(4) array<f32, 4u>(1.0, 2.0, 3.0, 4.0);
// s.a[1] = 5.0;
// }
ProgramBuilder b;
auto* S = b.Structure("S", utils::Vector{
b.Member("a", b.ty.array<f32, 4u>(utils::Vector{
b.Stride(4),
})),
});
b.GlobalVar("s", b.ty.Of(S), builtin::AddressSpace::kStorage, builtin::Access::kReadWrite,
b.Group(0_a), b.Binding(0_a));
b.Func("f", utils::Empty, b.ty.void_(),
utils::Vector{
b.Assign(b.MemberAccessor("s", "a"), b.Call(b.ty.array<f32, 4u>(utils::Vector{
b.Stride(4),
}))),
b.Assign(b.MemberAccessor("s", "a"), b.Call(b.ty.array<f32, 4u>(utils::Vector{
b.Stride(4),
}),
1_f, 2_f, 3_f, 4_f)),
b.Assign(b.IndexAccessor(b.MemberAccessor("s", "a"), 1_i), 5_f),
},
utils::Vector{
b.Stage(ast::PipelineStage::kCompute),
b.WorkgroupSize(1_i),
});
auto* expect =
R"(
struct S {
a : array<f32, 4u>,
}
@group(0) @binding(0) var<storage, read_write> s : S;
@compute @workgroup_size(1i)
fn f() {
s.a = array<f32, 4u>();
s.a = array<f32, 4u>(1.0f, 2.0f, 3.0f, 4.0f);
s.a[1i] = 5.0f;
}
)";
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, 4u>,
// };
// @group(0) @binding(0) var<storage, read_write> s : S;
//
// @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, 4u>(1.0, 2.0, 3.0, 4.0);
// (*b)[1] = 5.0;
// }
ProgramBuilder b;
auto* S = b.Structure("S", utils::Vector{b.Member("a", b.ty.array<f32, 4u>(utils::Vector{
b.Stride(32),
}))});
b.GlobalVar("s", b.ty.Of(S), builtin::AddressSpace::kStorage, builtin::Access::kReadWrite,
b.Group(0_a), b.Binding(0_a));
b.Func("f", utils::Empty, b.ty.void_(),
utils::Vector{
b.Decl(b.Let("a", b.AddressOf(b.MemberAccessor("s", "a")))),
b.Decl(b.Let("b", b.AddressOf(b.Deref(b.AddressOf(b.Deref("a")))))),
b.Decl(b.Let("c", b.Deref("b"))),
b.Decl(b.Let("d", b.IndexAccessor(b.Deref("b"), 1_i))),
b.Assign(b.Deref("b"), b.Call(b.ty.array<f32, 4u>(utils::Vector{
b.Stride(32),
}),
1_f, 2_f, 3_f, 4_f)),
b.Assign(b.IndexAccessor(b.Deref("b"), 1_i), 5_f),
},
utils::Vector{
b.Stage(ast::PipelineStage::kCompute),
b.WorkgroupSize(1_i),
});
auto* expect =
R"(
struct strided_arr {
@size(32)
el : f32,
}
struct S {
a : array<strided_arr, 4u>,
}
@group(0) @binding(0) var<storage, read_write> s : S;
@compute @workgroup_size(1i)
fn f() {
let c = s.a;
let d = s.a[1i].el;
s.a = array<strided_arr, 4u>(strided_arr(1.0f), strided_arr(2.0f), strided_arr(3.0f), strided_arr(4.0f));
s.a[1i].el = 5.0f;
}
)";
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, 4u>;
// struct S {
// a : ARR,
// };
// @group(0) @binding(0) var<storage, read_write> s : S;
//
// @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, 4u>(utils::Vector{
b.Stride(32),
}));
auto* S = b.Structure("S", utils::Vector{b.Member("a", b.ty("ARR"))});
b.GlobalVar("s", b.ty.Of(S), builtin::AddressSpace::kStorage, builtin::Access::kReadWrite,
b.Group(0_a), b.Binding(0_a));
b.Func("f", utils::Empty, b.ty.void_(),
utils::Vector{
b.Decl(b.Let("a", b.ty("ARR"), b.MemberAccessor("s", "a"))),
b.Decl(b.Let("b", b.ty.f32(), b.IndexAccessor(b.MemberAccessor("s", "a"), 1_i))),
b.Assign(b.MemberAccessor("s", "a"), b.Call("ARR")),
b.Assign(b.MemberAccessor("s", "a"), b.Call("ARR", 1_f, 2_f, 3_f, 4_f)),
b.Assign(b.IndexAccessor(b.MemberAccessor("s", "a"), 1_i), 5_f),
},
utils::Vector{
b.Stage(ast::PipelineStage::kCompute),
b.WorkgroupSize(1_i),
});
auto* expect = R"(
struct strided_arr {
@size(32)
el : f32,
}
alias ARR = array<strided_arr, 4u>;
struct S {
a : ARR,
}
@group(0) @binding(0) var<storage, read_write> s : S;
@compute @workgroup_size(1i)
fn f() {
let a : ARR = s.a;
let b : f32 = s.a[1i].el;
s.a = ARR();
s.a = ARR(strided_arr(1.0f), strided_arr(2.0f), strided_arr(3.0f), strided_arr(4.0f));
s.a[1i].el = 5.0f;
}
)";
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, 2u>;
// type ARR_B = @stride(128) array<@stride(16) array<ARR_A, 3u>, 4u>;
// struct S {
// a : ARR_B,
// };
// @group(0) @binding(0) var<storage, read_write> s : S;
//
// @compute @workgroup_size(1)
// fn f() {
// let a : ARR_B = s.a;
// let b : array<@stride(8) array<f32, 2u>, 3u> = 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>(utils::Vector{
b.Stride(8),
}));
b.Alias("ARR_B", b.ty.array( //
b.ty.array(b.ty("ARR_A"), 3_u,
utils::Vector{
b.Stride(16),
}),
4_u,
utils::Vector{
b.Stride(128),
}));
auto* S = b.Structure("S", utils::Vector{b.Member("a", b.ty("ARR_B"))});
b.GlobalVar("s", b.ty.Of(S), builtin::AddressSpace::kStorage, builtin::Access::kReadWrite,
b.Group(0_a), b.Binding(0_a));
b.Func("f", utils::Empty, b.ty.void_(),
utils::Vector{
b.Decl(b.Let("a", b.ty("ARR_B"), b.MemberAccessor("s", "a"))),
b.Decl(b.Let("b",
b.ty.array(b.ty("ARR_A"), 3_u,
utils::Vector{
b.Stride(16),
}),
b.IndexAccessor( //
b.MemberAccessor("s", "a"), //
3_i))),
b.Decl(b.Let("c", b.ty("ARR_A"),
b.IndexAccessor( //
b.IndexAccessor( //
b.MemberAccessor("s", "a"), //
3_i),
2_i))),
b.Decl(b.Let("d", b.ty.f32(),
b.IndexAccessor( //
b.IndexAccessor( //
b.IndexAccessor( //
b.MemberAccessor("s", "a"), //
3_i),
2_i),
1_i))),
b.Assign(b.MemberAccessor("s", "a"), b.Call("ARR_B")),
b.Assign(b.IndexAccessor( //
b.IndexAccessor( //
b.IndexAccessor( //
b.MemberAccessor("s", "a"), //
3_i),
2_i),
1_i),
5_f),
},
utils::Vector{
b.Stage(ast::PipelineStage::kCompute),
b.WorkgroupSize(1_i),
});
auto* expect =
R"(
struct strided_arr {
@size(8)
el : f32,
}
alias ARR_A = array<strided_arr, 2u>;
struct strided_arr_1 {
@size(128)
el : array<ARR_A, 3u>,
}
alias ARR_B = array<strided_arr_1, 4u>;
struct S {
a : ARR_B,
}
@group(0) @binding(0) var<storage, read_write> s : S;
@compute @workgroup_size(1i)
fn f() {
let a : ARR_B = s.a;
let b : array<ARR_A, 3u> = s.a[3i].el;
let c : ARR_A = s.a[3i].el[2i];
let d : f32 = s.a[3i].el[2i][1i].el;
s.a = ARR_B();
s.a[3i].el[2i][1i].el = 5.0f;
}
)";
auto got = Run<Unshadow, SimplifyPointers, DecomposeStridedArray>(Program(std::move(b)));
EXPECT_EQ(expect, str(got));
}
} // namespace
} // namespace tint::transform