// Copyright 2021 The Dawn & Tint Authors
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
//
// 1. Redistributions of source code must retain the above copyright notice, this
//    list of conditions and the following disclaimer.
//
// 2. Redistributions in binary form must reproduce the above copyright notice,
//    this list of conditions and the following disclaimer in the documentation
//    and/or other materials provided with the distribution.
//
// 3. Neither the name of the copyright holder nor the names of its
//    contributors may be used to endorse or promote products derived from
//    this software without specific prior written permission.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.

#include "src/tint/lang/wgsl/ast/transform/zero_init_workgroup_memory.h"

#include <utility>

#include "src/tint/lang/wgsl/ast/transform/helper_test.h"

namespace tint::ast::transform {
namespace {

using ZeroInitWorkgroupMemoryTest = TransformTest;

TEST_F(ZeroInitWorkgroupMemoryTest, ShouldRunEmptyModule) {
    auto* src = R"()";

    EXPECT_FALSE(ShouldRun<ZeroInitWorkgroupMemory>(src));
}

TEST_F(ZeroInitWorkgroupMemoryTest, ShouldRunHasNoWorkgroupVars) {
    auto* src = R"(
var<private> v : i32;
)";

    EXPECT_FALSE(ShouldRun<ZeroInitWorkgroupMemory>(src));
}

TEST_F(ZeroInitWorkgroupMemoryTest, ShouldRunHasWorkgroupVars) {
    auto* src = R"(
var<workgroup> a : i32;
)";

    EXPECT_TRUE(ShouldRun<ZeroInitWorkgroupMemory>(src));
}

TEST_F(ZeroInitWorkgroupMemoryTest, EmptyModule) {
    auto* src = "";
    auto* expect = src;

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, NoWorkgroupVars) {
    auto* src = R"(
var<private> v : i32;

fn f() {
  v = 1;
}
)";
    auto* expect = src;

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, UnreferencedWorkgroupVars) {
    auto* src = R"(
var<workgroup> a : i32;

var<workgroup> b : i32;

var<workgroup> c : i32;

fn unreferenced() {
  b = c;
}

@compute @workgroup_size(1)
fn f() {
}
)";
    auto* expect = src;

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, UnreferencedWorkgroupVars_OutOfOrder) {
    auto* src = R"(
@compute @workgroup_size(1)
fn f() {
}

fn unreferenced() {
  b = c;
}

var<workgroup> a : i32;

var<workgroup> b : i32;

var<workgroup> c : i32;
)";
    auto* expect = src;

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, SingleWorkgroupVar_ExistingLocalIndex) {
    auto* src = R"(
var<workgroup> v : i32;

@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  _ = v; // Initialization should be inserted above this statement
}
)";
    auto* expect = R"(
var<workgroup> v : i32;

@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  {
    v = i32();
  }
  workgroupBarrier();
  _ = v;
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, SingleWorkgroupVar_ExistingLocalIndex_OutOfOrder) {
    auto* src = R"(
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  _ = v; // Initialization should be inserted above this statement
}

var<workgroup> v : i32;
)";
    auto* expect = R"(
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  {
    v = i32();
  }
  workgroupBarrier();
  _ = v;
}

var<workgroup> v : i32;
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, SingleWorkgroupVar_ExistingLocalIndexInStruct) {
    auto* src = R"(
var<workgroup> v : i32;

struct Params {
  @builtin(local_invocation_index) local_idx : u32,
};

@compute @workgroup_size(1)
fn f(params : Params) {
  _ = v; // Initialization should be inserted above this statement
}
)";
    auto* expect = R"(
var<workgroup> v : i32;

struct Params {
  @builtin(local_invocation_index)
  local_idx : u32,
}

@compute @workgroup_size(1)
fn f(params : Params) {
  {
    v = i32();
  }
  workgroupBarrier();
  _ = v;
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, SingleWorkgroupVar_ExistingLocalIndexInStruct_OutOfOrder) {
    auto* src = R"(
@compute @workgroup_size(1)
fn f(params : Params) {
  _ = v; // Initialization should be inserted above this statement
}

struct Params {
  @builtin(local_invocation_index) local_idx : u32,
};

var<workgroup> v : i32;
)";
    auto* expect = R"(
@compute @workgroup_size(1)
fn f(params : Params) {
  {
    v = i32();
  }
  workgroupBarrier();
  _ = v;
}

struct Params {
  @builtin(local_invocation_index)
  local_idx : u32,
}

var<workgroup> v : i32;
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, SingleWorkgroupVar_InjectedLocalIndex) {
    auto* src = R"(
var<workgroup> v : i32;

@compute @workgroup_size(1)
fn f() {
  _ = v; // Initialization should be inserted above this statement
}
)";
    auto* expect = R"(
var<workgroup> v : i32;

@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
  {
    v = i32();
  }
  workgroupBarrier();
  _ = v;
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, SingleWorkgroupVar_InjectedLocalIndex_OutOfOrder) {
    auto* src = R"(
@compute @workgroup_size(1)
fn f() {
  _ = v; // Initialization should be inserted above this statement
}

var<workgroup> v : i32;
)";
    auto* expect = R"(
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
  {
    v = i32();
  }
  workgroupBarrier();
  _ = v;
}

var<workgroup> v : i32;
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, MultipleWorkgroupVar_ExistingLocalIndex_Size1) {
    auto* src = R"(
struct S {
  x : i32,
  y : array<i32, 8>,
};

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  _ = a; // Initialization should be inserted above this statement
  _ = b;
  _ = c;
}
)";
    auto* expect = R"(
struct S {
  x : i32,
  y : array<i32, 8>,
}

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  {
    a = i32();
    b.x = i32();
  }
  for(var idx : u32 = local_idx; (idx < 8u); idx = (idx + 1u)) {
    let i : u32 = idx;
    b.y[i] = i32();
  }
  for(var idx_1 : u32 = local_idx; (idx_1 < 32u); idx_1 = (idx_1 + 1u)) {
    let i_1 : u32 = idx_1;
    c[i_1].x = i32();
  }
  for(var idx_2 : u32 = local_idx; (idx_2 < 256u); idx_2 = (idx_2 + 1u)) {
    let i_2 : u32 = (idx_2 / 8u);
    let i : u32 = (idx_2 % 8u);
    c[i_2].y[i] = i32();
  }
  workgroupBarrier();
  _ = a;
  _ = b;
  _ = c;
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, MultipleWorkgroupVar_ExistingLocalIndex_Size1_OutOfOrder) {
    auto* src = R"(
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  _ = a; // Initialization should be inserted above this statement
  _ = b;
  _ = c;
}

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

struct S {
  x : i32,
  y : array<i32, 8>,
};
)";
    auto* expect = R"(
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  {
    a = i32();
    b.x = i32();
  }
  for(var idx : u32 = local_idx; (idx < 8u); idx = (idx + 1u)) {
    let i : u32 = idx;
    b.y[i] = i32();
  }
  for(var idx_1 : u32 = local_idx; (idx_1 < 32u); idx_1 = (idx_1 + 1u)) {
    let i_1 : u32 = idx_1;
    c[i_1].x = i32();
  }
  for(var idx_2 : u32 = local_idx; (idx_2 < 256u); idx_2 = (idx_2 + 1u)) {
    let i_2 : u32 = (idx_2 / 8u);
    let i : u32 = (idx_2 % 8u);
    c[i_2].y[i] = i32();
  }
  workgroupBarrier();
  _ = a;
  _ = b;
  _ = c;
}

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

struct S {
  x : i32,
  y : array<i32, 8>,
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, MultipleWorkgroupVar_ExistingLocalIndex_Size_2_3) {
    auto* src = R"(
struct S {
  x : i32,
  y : array<i32, 8>,
};

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

@compute @workgroup_size(2, 3)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  _ = a; // Initialization should be inserted above this statement
  _ = b;
  _ = c;
}
)";
    auto* expect = R"(
struct S {
  x : i32,
  y : array<i32, 8>,
}

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

@compute @workgroup_size(2, 3)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  if ((local_idx < 1u)) {
    a = i32();
    b.x = i32();
  }
  for(var idx : u32 = local_idx; (idx < 8u); idx = (idx + 6u)) {
    let i : u32 = idx;
    b.y[i] = i32();
  }
  for(var idx_1 : u32 = local_idx; (idx_1 < 32u); idx_1 = (idx_1 + 6u)) {
    let i_1 : u32 = idx_1;
    c[i_1].x = i32();
  }
  for(var idx_2 : u32 = local_idx; (idx_2 < 256u); idx_2 = (idx_2 + 6u)) {
    let i_2 : u32 = (idx_2 / 8u);
    let i : u32 = (idx_2 % 8u);
    c[i_2].y[i] = i32();
  }
  workgroupBarrier();
  _ = a;
  _ = b;
  _ = c;
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, MultipleWorkgroupVar_ExistingLocalIndex_Size_2_3_X) {
    auto* src = R"(
struct S {
  x : i32,
  y : array<i32, 8>,
};

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

@id(1) override X : i32;

@compute @workgroup_size(2, 3, X)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  _ = a; // Initialization should be inserted above this statement
  _ = b;
  _ = c;
}
)";
    auto* expect =
        R"(
struct S {
  x : i32,
  y : array<i32, 8>,
}

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

@id(1) override X : i32;

@compute @workgroup_size(2, 3, X)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  for(var idx : u32 = local_idx; (idx < 1u); idx = (idx + (u32(X) * 6u))) {
    a = i32();
    b.x = i32();
  }
  for(var idx_1 : u32 = local_idx; (idx_1 < 8u); idx_1 = (idx_1 + (u32(X) * 6u))) {
    let i : u32 = idx_1;
    b.y[i] = i32();
  }
  for(var idx_2 : u32 = local_idx; (idx_2 < 32u); idx_2 = (idx_2 + (u32(X) * 6u))) {
    let i_1 : u32 = idx_2;
    c[i_1].x = i32();
  }
  for(var idx_3 : u32 = local_idx; (idx_3 < 256u); idx_3 = (idx_3 + (u32(X) * 6u))) {
    let i_2 : u32 = (idx_3 / 8u);
    let i : u32 = (idx_3 % 8u);
    c[i_2].y[i] = i32();
  }
  workgroupBarrier();
  _ = a;
  _ = b;
  _ = c;
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, MultipleWorkgroupVar_ExistingLocalIndex_Size_5u_X_10u) {
    auto* src = R"(
struct S {
  x : array<array<i32, 8>, 10>,
  y : array<i32, 8>,
  z : array<array<array<i32, 8>, 10>, 20>,
};

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

@id(1) override X : u32;

@compute @workgroup_size(5u, X, 10u)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  _ = a; // Initialization should be inserted above this statement
  _ = b;
  _ = c;
}
)";
    auto* expect =
        R"(
struct S {
  x : array<array<i32, 8>, 10>,
  y : array<i32, 8>,
  z : array<array<array<i32, 8>, 10>, 20>,
}

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

@id(1) override X : u32;

@compute @workgroup_size(5u, X, 10u)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  for(var idx : u32 = local_idx; (idx < 1u); idx = (idx + (X * 50u))) {
    a = i32();
  }
  for(var idx_1 : u32 = local_idx; (idx_1 < 8u); idx_1 = (idx_1 + (X * 50u))) {
    let i_1 : u32 = idx_1;
    b.y[i_1] = i32();
  }
  for(var idx_2 : u32 = local_idx; (idx_2 < 80u); idx_2 = (idx_2 + (X * 50u))) {
    let i : u32 = (idx_2 / 8u);
    let i_1 : u32 = (idx_2 % 8u);
    b.x[i][i_1] = i32();
  }
  for(var idx_3 : u32 = local_idx; (idx_3 < 256u); idx_3 = (idx_3 + (X * 50u))) {
    let i_4 : u32 = (idx_3 / 8u);
    let i_1 : u32 = (idx_3 % 8u);
    c[i_4].y[i_1] = i32();
  }
  for(var idx_4 : u32 = local_idx; (idx_4 < 1600u); idx_4 = (idx_4 + (X * 50u))) {
    let i_2 : u32 = (idx_4 / 80u);
    let i : u32 = ((idx_4 % 80u) / 8u);
    let i_1 : u32 = (idx_4 % 8u);
    b.z[i_2][i][i_1] = i32();
  }
  for(var idx_5 : u32 = local_idx; (idx_5 < 2560u); idx_5 = (idx_5 + (X * 50u))) {
    let i_3 : u32 = (idx_5 / 80u);
    let i : u32 = ((idx_5 % 80u) / 8u);
    let i_1 : u32 = (idx_5 % 8u);
    c[i_3].x[i][i_1] = i32();
  }
  for(var idx_6 : u32 = local_idx; (idx_6 < 51200u); idx_6 = (idx_6 + (X * 50u))) {
    let i_5 : u32 = (idx_6 / 1600u);
    let i_2 : u32 = ((idx_6 % 1600u) / 80u);
    let i : u32 = ((idx_6 % 80u) / 8u);
    let i_1 : u32 = (idx_6 % 8u);
    c[i_5].z[i_2][i][i_1] = i32();
  }
  workgroupBarrier();
  _ = a;
  _ = b;
  _ = c;
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, MultipleWorkgroupVar_InjectedLocalIndex) {
    auto* src = R"(
struct S {
  x : i32,
  y : array<i32, 8>,
};

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

@compute @workgroup_size(1)
fn f(@builtin(local_invocation_id) local_invocation_id : vec3<u32>) {
  _ = a; // Initialization should be inserted above this statement
  _ = b;
  _ = c;
}
)";
    auto* expect = R"(
struct S {
  x : i32,
  y : array<i32, 8>,
}

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

@compute @workgroup_size(1)
fn f(@builtin(local_invocation_id) local_invocation_id : vec3<u32>, @builtin(local_invocation_index) local_invocation_index : u32) {
  {
    a = i32();
    b.x = i32();
  }
  for(var idx : u32 = local_invocation_index; (idx < 8u); idx = (idx + 1u)) {
    let i : u32 = idx;
    b.y[i] = i32();
  }
  for(var idx_1 : u32 = local_invocation_index; (idx_1 < 32u); idx_1 = (idx_1 + 1u)) {
    let i_1 : u32 = idx_1;
    c[i_1].x = i32();
  }
  for(var idx_2 : u32 = local_invocation_index; (idx_2 < 256u); idx_2 = (idx_2 + 1u)) {
    let i_2 : u32 = (idx_2 / 8u);
    let i : u32 = (idx_2 % 8u);
    c[i_2].y[i] = i32();
  }
  workgroupBarrier();
  _ = a;
  _ = b;
  _ = c;
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, MultipleWorkgroupVar_InjectedLocalIndex_OutOfOrder) {
    auto* src = R"(
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_id) local_invocation_id : vec3<u32>) {
  _ = a; // Initialization should be inserted above this statement
  _ = b;
  _ = c;
}

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

struct S {
  x : i32,
  y : array<i32, 8>,
};
)";
    auto* expect = R"(
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_id) local_invocation_id : vec3<u32>, @builtin(local_invocation_index) local_invocation_index : u32) {
  {
    a = i32();
    b.x = i32();
  }
  for(var idx : u32 = local_invocation_index; (idx < 8u); idx = (idx + 1u)) {
    let i : u32 = idx;
    b.y[i] = i32();
  }
  for(var idx_1 : u32 = local_invocation_index; (idx_1 < 32u); idx_1 = (idx_1 + 1u)) {
    let i_1 : u32 = idx_1;
    c[i_1].x = i32();
  }
  for(var idx_2 : u32 = local_invocation_index; (idx_2 < 256u); idx_2 = (idx_2 + 1u)) {
    let i_2 : u32 = (idx_2 / 8u);
    let i : u32 = (idx_2 % 8u);
    c[i_2].y[i] = i32();
  }
  workgroupBarrier();
  _ = a;
  _ = b;
  _ = c;
}

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

struct S {
  x : i32,
  y : array<i32, 8>,
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, MultipleWorkgroupVar_MultipleEntryPoints) {
    auto* src = R"(
struct S {
  x : i32,
  y : array<i32, 8>,
};

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

@compute @workgroup_size(1)
fn f1() {
  _ = a; // Initialization should be inserted above this statement
  _ = c;
}

@compute @workgroup_size(1, 2, 3)
fn f2(@builtin(local_invocation_id) local_invocation_id : vec3<u32>) {
  _ = b; // Initialization should be inserted above this statement
}

@compute @workgroup_size(4, 5, 6)
fn f3() {
  _ = c; // Initialization should be inserted above this statement
  _ = a;
}
)";
    auto* expect = R"(
struct S {
  x : i32,
  y : array<i32, 8>,
}

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

@compute @workgroup_size(1)
fn f1(@builtin(local_invocation_index) local_invocation_index : u32) {
  {
    a = i32();
  }
  for(var idx : u32 = local_invocation_index; (idx < 32u); idx = (idx + 1u)) {
    let i : u32 = idx;
    c[i].x = i32();
  }
  for(var idx_1 : u32 = local_invocation_index; (idx_1 < 256u); idx_1 = (idx_1 + 1u)) {
    let i_1 : u32 = (idx_1 / 8u);
    let i_2 : u32 = (idx_1 % 8u);
    c[i_1].y[i_2] = i32();
  }
  workgroupBarrier();
  _ = a;
  _ = c;
}

@compute @workgroup_size(1, 2, 3)
fn f2(@builtin(local_invocation_id) local_invocation_id : vec3<u32>, @builtin(local_invocation_index) local_invocation_index_1 : u32) {
  if ((local_invocation_index_1 < 1u)) {
    b.x = i32();
  }
  for(var idx_2 : u32 = local_invocation_index_1; (idx_2 < 8u); idx_2 = (idx_2 + 6u)) {
    let i_3 : u32 = idx_2;
    b.y[i_3] = i32();
  }
  workgroupBarrier();
  _ = b;
}

@compute @workgroup_size(4, 5, 6)
fn f3(@builtin(local_invocation_index) local_invocation_index_2 : u32) {
  if ((local_invocation_index_2 < 1u)) {
    a = i32();
  }
  if ((local_invocation_index_2 < 32u)) {
    let i_4 : u32 = local_invocation_index_2;
    c[i_4].x = i32();
  }
  for(var idx_3 : u32 = local_invocation_index_2; (idx_3 < 256u); idx_3 = (idx_3 + 120u)) {
    let i_5 : u32 = (idx_3 / 8u);
    let i_6 : u32 = (idx_3 % 8u);
    c[i_5].y[i_6] = i32();
  }
  workgroupBarrier();
  _ = c;
  _ = a;
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, MultipleWorkgroupVar_MultipleEntryPoints_OutOfOrder) {
    auto* src = R"(
@compute @workgroup_size(1)
fn f1() {
  _ = a; // Initialization should be inserted above this statement
  _ = c;
}

@compute @workgroup_size(1, 2, 3)
fn f2(@builtin(local_invocation_id) local_invocation_id : vec3<u32>) {
  _ = b; // Initialization should be inserted above this statement
}

@compute @workgroup_size(4, 5, 6)
fn f3() {
  _ = c; // Initialization should be inserted above this statement
  _ = a;
}

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

struct S {
  x : i32,
  y : array<i32, 8>,
};
)";
    auto* expect = R"(
@compute @workgroup_size(1)
fn f1(@builtin(local_invocation_index) local_invocation_index : u32) {
  {
    a = i32();
  }
  for(var idx : u32 = local_invocation_index; (idx < 32u); idx = (idx + 1u)) {
    let i : u32 = idx;
    c[i].x = i32();
  }
  for(var idx_1 : u32 = local_invocation_index; (idx_1 < 256u); idx_1 = (idx_1 + 1u)) {
    let i_1 : u32 = (idx_1 / 8u);
    let i_2 : u32 = (idx_1 % 8u);
    c[i_1].y[i_2] = i32();
  }
  workgroupBarrier();
  _ = a;
  _ = c;
}

@compute @workgroup_size(1, 2, 3)
fn f2(@builtin(local_invocation_id) local_invocation_id : vec3<u32>, @builtin(local_invocation_index) local_invocation_index_1 : u32) {
  if ((local_invocation_index_1 < 1u)) {
    b.x = i32();
  }
  for(var idx_2 : u32 = local_invocation_index_1; (idx_2 < 8u); idx_2 = (idx_2 + 6u)) {
    let i_3 : u32 = idx_2;
    b.y[i_3] = i32();
  }
  workgroupBarrier();
  _ = b;
}

@compute @workgroup_size(4, 5, 6)
fn f3(@builtin(local_invocation_index) local_invocation_index_2 : u32) {
  if ((local_invocation_index_2 < 1u)) {
    a = i32();
  }
  if ((local_invocation_index_2 < 32u)) {
    let i_4 : u32 = local_invocation_index_2;
    c[i_4].x = i32();
  }
  for(var idx_3 : u32 = local_invocation_index_2; (idx_3 < 256u); idx_3 = (idx_3 + 120u)) {
    let i_5 : u32 = (idx_3 / 8u);
    let i_6 : u32 = (idx_3 % 8u);
    c[i_5].y[i_6] = i32();
  }
  workgroupBarrier();
  _ = c;
  _ = a;
}

var<workgroup> a : i32;

var<workgroup> b : S;

var<workgroup> c : array<S, 32>;

struct S {
  x : i32,
  y : array<i32, 8>,
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, TransitiveUsage) {
    auto* src = R"(
var<workgroup> v : i32;

fn use_v() {
  _ = v;
}

fn call_use_v() {
  use_v();
}

@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  call_use_v(); // Initialization should be inserted above this statement
}
)";
    auto* expect = R"(
var<workgroup> v : i32;

fn use_v() {
  _ = v;
}

fn call_use_v() {
  use_v();
}

@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  {
    v = i32();
  }
  workgroupBarrier();
  call_use_v();
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, TransitiveUsage_OutOfOrder) {
    auto* src = R"(
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  call_use_v(); // Initialization should be inserted above this statement
}

fn call_use_v() {
  use_v();
}

fn use_v() {
  _ = v;
}

var<workgroup> v : i32;
)";
    auto* expect = R"(
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_idx : u32) {
  {
    v = i32();
  }
  workgroupBarrier();
  call_use_v();
}

fn call_use_v() {
  use_v();
}

fn use_v() {
  _ = v;
}

var<workgroup> v : i32;
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupAtomics) {
    auto* src = R"(
var<workgroup> i : atomic<i32>;
var<workgroup> u : atomic<u32>;

@compute @workgroup_size(1)
fn f() {
  atomicLoad(&(i)); // Initialization should be inserted above this statement
  atomicLoad(&(u));
}
)";
    auto* expect = R"(
var<workgroup> i : atomic<i32>;

var<workgroup> u : atomic<u32>;

@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
  {
    atomicStore(&(i), i32());
    atomicStore(&(u), u32());
  }
  workgroupBarrier();
  atomicLoad(&(i));
  atomicLoad(&(u));
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupAtomics_OutOfOrder) {
    auto* src = R"(
@compute @workgroup_size(1)
fn f() {
  atomicLoad(&(i)); // Initialization should be inserted above this statement
  atomicLoad(&(u));
}

var<workgroup> i : atomic<i32>;
var<workgroup> u : atomic<u32>;
)";
    auto* expect = R"(
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
  {
    atomicStore(&(i), i32());
    atomicStore(&(u), u32());
  }
  workgroupBarrier();
  atomicLoad(&(i));
  atomicLoad(&(u));
}

var<workgroup> i : atomic<i32>;

var<workgroup> u : atomic<u32>;
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupStructOfAtomics) {
    auto* src = R"(
struct S {
  a : i32,
  i : atomic<i32>,
  b : f32,
  u : atomic<u32>,
  c : u32,
};

var<workgroup> w : S;

@compute @workgroup_size(1)
fn f() {
  _ = w.a; // Initialization should be inserted above this statement
}
)";
    auto* expect = R"(
struct S {
  a : i32,
  i : atomic<i32>,
  b : f32,
  u : atomic<u32>,
  c : u32,
}

var<workgroup> w : S;

@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
  {
    w.a = i32();
    atomicStore(&(w.i), i32());
    w.b = f32();
    atomicStore(&(w.u), u32());
    w.c = u32();
  }
  workgroupBarrier();
  _ = w.a;
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupStructOfAtomics_OutOfOrder) {
    auto* src = R"(
@compute @workgroup_size(1)
fn f() {
  _ = w.a; // Initialization should be inserted above this statement
}

var<workgroup> w : S;

struct S {
  a : i32,
  i : atomic<i32>,
  b : f32,
  u : atomic<u32>,
  c : u32,
};
)";
    auto* expect = R"(
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
  {
    w.a = i32();
    atomicStore(&(w.i), i32());
    w.b = f32();
    atomicStore(&(w.u), u32());
    w.c = u32();
  }
  workgroupBarrier();
  _ = w.a;
}

var<workgroup> w : S;

struct S {
  a : i32,
  i : atomic<i32>,
  b : f32,
  u : atomic<u32>,
  c : u32,
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupArrayOfAtomics) {
    auto* src = R"(
var<workgroup> w : array<atomic<u32>, 4>;

@compute @workgroup_size(1)
fn f() {
  atomicLoad(&w[0]); // Initialization should be inserted above this statement
}
)";
    auto* expect = R"(
var<workgroup> w : array<atomic<u32>, 4>;

@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
  for(var idx : u32 = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
    let i : u32 = idx;
    atomicStore(&(w[i]), u32());
  }
  workgroupBarrier();
  atomicLoad(&(w[0]));
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupArrayOfAtomics_OutOfOrder) {
    auto* src = R"(
@compute @workgroup_size(1)
fn f() {
  atomicLoad(&w[0]); // Initialization should be inserted above this statement
}

var<workgroup> w : array<atomic<u32>, 4>;
)";
    auto* expect = R"(
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
  for(var idx : u32 = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
    let i : u32 = idx;
    atomicStore(&(w[i]), u32());
  }
  workgroupBarrier();
  atomicLoad(&(w[0]));
}

var<workgroup> w : array<atomic<u32>, 4>;
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupArrayOfStructOfAtomics) {
    auto* src = R"(
struct S {
  a : i32,
  i : atomic<i32>,
  b : f32,
  u : atomic<u32>,
  c : u32,
};

var<workgroup> w : array<S, 4>;

@compute @workgroup_size(1)
fn f() {
  _ = w[0].a; // Initialization should be inserted above this statement
}
)";
    auto* expect = R"(
struct S {
  a : i32,
  i : atomic<i32>,
  b : f32,
  u : atomic<u32>,
  c : u32,
}

var<workgroup> w : array<S, 4>;

@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
  for(var idx : u32 = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
    let i_1 : u32 = idx;
    w[i_1].a = i32();
    atomicStore(&(w[i_1].i), i32());
    w[i_1].b = f32();
    atomicStore(&(w[i_1].u), u32());
    w[i_1].c = u32();
  }
  workgroupBarrier();
  _ = w[0].a;
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, WorkgroupArrayOfStructOfAtomics_OutOfOrder) {
    auto* src = R"(
@compute @workgroup_size(1)
fn f() {
  _ = w[0].a; // Initialization should be inserted above this statement
}

var<workgroup> w : array<S, 4>;

struct S {
  a : i32,
  i : atomic<i32>,
  b : f32,
  u : atomic<u32>,
  c : u32,
};
)";
    auto* expect = R"(
@compute @workgroup_size(1)
fn f(@builtin(local_invocation_index) local_invocation_index : u32) {
  for(var idx : u32 = local_invocation_index; (idx < 4u); idx = (idx + 1u)) {
    let i_1 : u32 = idx;
    w[i_1].a = i32();
    atomicStore(&(w[i_1].i), i32());
    w[i_1].b = f32();
    atomicStore(&(w[i_1].u), u32());
    w[i_1].c = u32();
  }
  workgroupBarrier();
  _ = w[0].a;
}

var<workgroup> w : array<S, 4>;

struct S {
  a : i32,
  i : atomic<i32>,
  b : f32,
  u : atomic<u32>,
  c : u32,
}
)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

TEST_F(ZeroInitWorkgroupMemoryTest, ArrayWithOverrideCount) {
    auto* src =
        R"(override O = 123;
alias A = array<i32, O*2>;

var<workgroup> W : A;

@compute @workgroup_size(1)
fn main() {
    let p : ptr<workgroup, A> = &W;
    (*p)[0] = 42;
}
)";

    auto* expect =
        R"(error: array size is an override-expression, when expected a constant-expression.
Was the SubstituteOverride transform run?)";

    auto got = Run<ZeroInitWorkgroupMemory>(src);

    EXPECT_EQ(expect, str(got));
}

}  // namespace
}  // namespace tint::ast::transform