blob: 197aa02dcdd544f9c1d74b796c4ce7361e60e296 [file] [log] [blame]
// 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.
#include "src/transform/localize_struct_array_assignment.h"
#include "src/transform/simplify_pointers.h"
#include "src/transform/unshadow.h"
#include "src/transform/test_helper.h"
namespace tint {
namespace transform {
namespace {
using LocalizeStructArrayAssignmentTest = TransformTest;
TEST_F(LocalizeStructArrayAssignmentTest, MissingSimplifyPointers) {
auto* src = R"()";
auto* expect =
"error: tint::transform::LocalizeStructArrayAssignment depends on "
"tint::transform::SimplifyPointers but the dependency was not run";
auto got = Run<LocalizeStructArrayAssignment>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(LocalizeStructArrayAssignmentTest, EmptyModule) {
auto* src = R"()";
auto* expect = src;
auto got =
Run<Unshadow, SimplifyPointers, LocalizeStructArrayAssignment>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(LocalizeStructArrayAssignmentTest, StructArray) {
auto* src = R"(
[[block]] struct Uniforms {
i : u32;
};
struct InnerS {
v : i32;
};
struct OuterS {
a1 : array<InnerS, 8>;
};
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
[[stage(compute), workgroup_size(1)]]
fn main() {
var v : InnerS;
var s1 : OuterS;
s1.a1[uniforms.i] = v;
}
)";
auto* expect = R"(
[[block]]
struct Uniforms {
i : u32;
};
struct InnerS {
v : i32;
};
struct OuterS {
a1 : array<InnerS, 8>;
};
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
[[stage(compute), workgroup_size(1)]]
fn main() {
var v : InnerS;
var s1 : OuterS;
{
let tint_symbol = &(s1.a1);
var tint_symbol_1 = *(tint_symbol);
tint_symbol_1[uniforms.i] = v;
*(tint_symbol) = tint_symbol_1;
}
}
)";
auto got =
Run<Unshadow, SimplifyPointers, LocalizeStructArrayAssignment>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(LocalizeStructArrayAssignmentTest, StructStructArray) {
auto* src = R"(
[[block]] struct Uniforms {
i : u32;
};
struct InnerS {
v : i32;
};
struct S1 {
a : array<InnerS, 8>;
};
struct OuterS {
s2 : S1;
};
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
[[stage(compute), workgroup_size(1)]]
fn main() {
var v : InnerS;
var s1 : OuterS;
s1.s2.a[uniforms.i] = v;
}
)";
auto* expect = R"(
[[block]]
struct Uniforms {
i : u32;
};
struct InnerS {
v : i32;
};
struct S1 {
a : array<InnerS, 8>;
};
struct OuterS {
s2 : S1;
};
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
[[stage(compute), workgroup_size(1)]]
fn main() {
var v : InnerS;
var s1 : OuterS;
{
let tint_symbol = &(s1.s2.a);
var tint_symbol_1 = *(tint_symbol);
tint_symbol_1[uniforms.i] = v;
*(tint_symbol) = tint_symbol_1;
}
}
)";
auto got =
Run<Unshadow, SimplifyPointers, LocalizeStructArrayAssignment>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(LocalizeStructArrayAssignmentTest, StructArrayArray) {
auto* src = R"(
[[block]] struct Uniforms {
i : u32;
j : u32;
};
struct InnerS {
v : i32;
};
struct OuterS {
a1 : array<array<InnerS, 8>, 8>;
};
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
[[stage(compute), workgroup_size(1)]]
fn main() {
var v : InnerS;
var s1 : OuterS;
s1.a1[uniforms.i][uniforms.j] = v;
}
)";
auto* expect = R"(
[[block]]
struct Uniforms {
i : u32;
j : u32;
};
struct InnerS {
v : i32;
};
struct OuterS {
a1 : array<array<InnerS, 8>, 8>;
};
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
[[stage(compute), workgroup_size(1)]]
fn main() {
var v : InnerS;
var s1 : OuterS;
{
let tint_symbol = &(s1.a1);
var tint_symbol_1 = *(tint_symbol);
tint_symbol_1[uniforms.i][uniforms.j] = v;
*(tint_symbol) = tint_symbol_1;
}
}
)";
auto got =
Run<Unshadow, SimplifyPointers, LocalizeStructArrayAssignment>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(LocalizeStructArrayAssignmentTest, StructArrayStruct) {
auto* src = R"(
[[block]] struct Uniforms {
i : u32;
};
struct InnerS {
v : i32;
};
struct S1 {
s2 : InnerS;
};
struct OuterS {
a1 : array<S1, 8>;
};
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
[[stage(compute), workgroup_size(1)]]
fn main() {
var v : InnerS;
var s1 : OuterS;
s1.a1[uniforms.i].s2 = v;
}
)";
auto* expect = R"(
[[block]]
struct Uniforms {
i : u32;
};
struct InnerS {
v : i32;
};
struct S1 {
s2 : InnerS;
};
struct OuterS {
a1 : array<S1, 8>;
};
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
[[stage(compute), workgroup_size(1)]]
fn main() {
var v : InnerS;
var s1 : OuterS;
{
let tint_symbol = &(s1.a1);
var tint_symbol_1 = *(tint_symbol);
tint_symbol_1[uniforms.i].s2 = v;
*(tint_symbol) = tint_symbol_1;
}
}
)";
auto got =
Run<Unshadow, SimplifyPointers, LocalizeStructArrayAssignment>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(LocalizeStructArrayAssignmentTest, StructArrayStructArray) {
auto* src = R"(
[[block]] struct Uniforms {
i : u32;
j : u32;
};
struct InnerS {
v : i32;
};
struct S1 {
a2 : array<InnerS, 8>;
};
struct OuterS {
a1 : array<S1, 8>;
};
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
[[stage(compute), workgroup_size(1)]]
fn main() {
var v : InnerS;
var s : OuterS;
s.a1[uniforms.i].a2[uniforms.j] = v;
}
)";
auto* expect = R"(
[[block]]
struct Uniforms {
i : u32;
j : u32;
};
struct InnerS {
v : i32;
};
struct S1 {
a2 : array<InnerS, 8>;
};
struct OuterS {
a1 : array<S1, 8>;
};
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
[[stage(compute), workgroup_size(1)]]
fn main() {
var v : InnerS;
var s : OuterS;
{
let tint_symbol = &(s.a1);
var tint_symbol_1 = *(tint_symbol);
let tint_symbol_2 = &(tint_symbol_1[uniforms.i].a2);
var tint_symbol_3 = *(tint_symbol_2);
tint_symbol_3[uniforms.j] = v;
*(tint_symbol_2) = tint_symbol_3;
*(tint_symbol) = tint_symbol_1;
}
}
)";
auto got =
Run<Unshadow, SimplifyPointers, LocalizeStructArrayAssignment>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(LocalizeStructArrayAssignmentTest, IndexingWithSideEffectFunc) {
auto* src = R"(
[[block]] struct Uniforms {
i : u32;
j : u32;
};
struct InnerS {
v : i32;
};
struct S1 {
a2 : array<InnerS, 8>;
};
struct OuterS {
a1 : array<S1, 8>;
};
var<private> nextIndex : u32;
fn getNextIndex() -> u32 {
nextIndex = nextIndex + 1u;
return nextIndex;
}
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
[[stage(compute), workgroup_size(1)]]
fn main() {
var v : InnerS;
var s : OuterS;
s.a1[getNextIndex()].a2[uniforms.j] = v;
}
)";
auto* expect = R"(
[[block]]
struct Uniforms {
i : u32;
j : u32;
};
struct InnerS {
v : i32;
};
struct S1 {
a2 : array<InnerS, 8>;
};
struct OuterS {
a1 : array<S1, 8>;
};
var<private> nextIndex : u32;
fn getNextIndex() -> u32 {
nextIndex = (nextIndex + 1u);
return nextIndex;
}
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
[[stage(compute), workgroup_size(1)]]
fn main() {
var v : InnerS;
var s : OuterS;
{
let tint_symbol = &(s.a1);
var tint_symbol_1 = *(tint_symbol);
let tint_symbol_2 = &(tint_symbol_1[getNextIndex()].a2);
var tint_symbol_3 = *(tint_symbol_2);
tint_symbol_3[uniforms.j] = v;
*(tint_symbol_2) = tint_symbol_3;
*(tint_symbol) = tint_symbol_1;
}
}
)";
auto got =
Run<Unshadow, SimplifyPointers, LocalizeStructArrayAssignment>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(LocalizeStructArrayAssignmentTest, ViaPointerArg) {
auto* src = R"(
[[block]] struct Uniforms {
i : u32;
};
struct InnerS {
v : i32;
};
struct OuterS {
a1 : array<InnerS, 8>;
};
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
fn f(p : ptr<function, OuterS>) {
var v : InnerS;
(*p).a1[uniforms.i] = v;
}
[[stage(compute), workgroup_size(1)]]
fn main() {
var s1 : OuterS;
f(&s1);
}
)";
auto* expect = R"(
[[block]]
struct Uniforms {
i : u32;
};
struct InnerS {
v : i32;
};
struct OuterS {
a1 : array<InnerS, 8>;
};
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
fn f(p : ptr<function, OuterS>) {
var v : InnerS;
{
let tint_symbol = &((*(p)).a1);
var tint_symbol_1 = *(tint_symbol);
tint_symbol_1[uniforms.i] = v;
*(tint_symbol) = tint_symbol_1;
}
}
[[stage(compute), workgroup_size(1)]]
fn main() {
var s1 : OuterS;
f(&(s1));
}
)";
auto got =
Run<Unshadow, SimplifyPointers, LocalizeStructArrayAssignment>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(LocalizeStructArrayAssignmentTest, ViaPointerVar) {
auto* src = R"(
[[block]]
struct Uniforms {
i : u32;
};
struct InnerS {
v : i32;
};
struct OuterS {
a1 : array<InnerS, 8>;
};
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
fn f(p : ptr<function, InnerS>, v : InnerS) {
*(p) = v;
}
[[stage(compute), workgroup_size(1)]]
fn main() {
var v : InnerS;
var s1 : OuterS;
let p = &(s1.a1[uniforms.i]);
*(p) = v;
}
)";
auto* expect = R"(
[[block]]
struct Uniforms {
i : u32;
};
struct InnerS {
v : i32;
};
struct OuterS {
a1 : array<InnerS, 8>;
};
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
fn f(p : ptr<function, InnerS>, v : InnerS) {
*(p) = v;
}
[[stage(compute), workgroup_size(1)]]
fn main() {
var v : InnerS;
var s1 : OuterS;
let p_save = uniforms.i;
{
let tint_symbol = &(s1.a1);
var tint_symbol_1 = *(tint_symbol);
tint_symbol_1[p_save] = v;
*(tint_symbol) = tint_symbol_1;
}
}
)";
auto got =
Run<Unshadow, SimplifyPointers, LocalizeStructArrayAssignment>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(LocalizeStructArrayAssignmentTest, VectorAssignment) {
auto* src = R"(
[[block]]
struct Uniforms {
i : u32;
};
[[block]]
struct OuterS {
a1 : array<u32, 8>;
};
[[group(1), binding(4)]] var<uniform> uniforms : Uniforms;
fn f(i : u32) -> u32 {
return (i + 1u);
}
[[stage(compute), workgroup_size(1)]]
fn main() {
var s1 : OuterS;
var v : vec3<f32>;
v[s1.a1[uniforms.i]] = 1.0;
v[f(s1.a1[uniforms.i])] = 1.0;
}
)";
// Transform does nothing here as we're not actually assigning to the array in
// the struct.
auto* expect = src;
auto got =
Run<Unshadow, SimplifyPointers, LocalizeStructArrayAssignment>(src);
EXPECT_EQ(expect, str(got));
}
} // namespace
} // namespace transform
} // namespace tint