dawn-cmake/src/transform/localize_struct_array_assignment_test.cc

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