blob: 4d4ede32b52e9c144df8ccc66068867adebef02c [file] [log] [blame]
// 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/single_entry_point.h"
#include <utility>
#include "src/tint/lang/wgsl/ast/transform/helper_test.h"
namespace tint::ast::transform {
namespace {
using SingleEntryPointTest = TransformTest;
TEST_F(SingleEntryPointTest, Error_MissingTransformData) {
auto* src = "";
auto* expect = "error: missing transform data for tint::ast::transform::SingleEntryPoint";
auto got = Run<SingleEntryPoint>(src);
EXPECT_EQ(expect, str(got));
}
TEST_F(SingleEntryPointTest, Error_NoEntryPoints) {
auto* src = "";
auto* expect = "error: entry point 'main' not found";
DataMap data;
data.Add<SingleEntryPoint::Config>("main");
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
TEST_F(SingleEntryPointTest, Error_InvalidEntryPoint) {
auto* src = R"(
@vertex
fn main() -> @builtin(position) vec4<f32> {
return vec4<f32>();
}
)";
auto* expect = "error: entry point '_' not found";
SingleEntryPoint::Config cfg("_");
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
TEST_F(SingleEntryPointTest, Error_NotAnEntryPoint) {
auto* src = R"(
fn foo() {}
@fragment
fn main() {}
)";
auto* expect = "error: entry point 'foo' not found";
SingleEntryPoint::Config cfg("foo");
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
TEST_F(SingleEntryPointTest, SingleEntryPoint) {
auto* src = R"(
@compute @workgroup_size(1)
fn main() {
}
)";
SingleEntryPoint::Config cfg("main");
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(src, str(got));
}
TEST_F(SingleEntryPointTest, MultipleEntryPoints) {
auto* src = R"(
@vertex
fn vert_main() -> @builtin(position) vec4<f32> {
return vec4<f32>();
}
@fragment
fn frag_main() {
}
@compute @workgroup_size(1)
fn comp_main1() {
}
@compute @workgroup_size(1)
fn comp_main2() {
}
)";
auto* expect = R"(
@compute @workgroup_size(1)
fn comp_main1() {
}
)";
SingleEntryPoint::Config cfg("comp_main1");
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
TEST_F(SingleEntryPointTest, GlobalVariables) {
auto* src = R"(
var<private> a : f32;
var<private> b : f32;
var<private> c : f32;
var<private> d : f32;
@vertex
fn vert_main() -> @builtin(position) vec4<f32> {
a = 0.0;
return vec4<f32>();
}
@fragment
fn frag_main() {
b = 0.0;
}
@compute @workgroup_size(1)
fn comp_main1() {
c = 0.0;
}
@compute @workgroup_size(1)
fn comp_main2() {
d = 0.0;
}
)";
auto* expect = R"(
var<private> c : f32;
@compute @workgroup_size(1)
fn comp_main1() {
c = 0.0;
}
)";
SingleEntryPoint::Config cfg("comp_main1");
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
TEST_F(SingleEntryPointTest, GlobalConstants) {
auto* src = R"(
const a : f32 = 1.0;
const b : f32 = 1.0;
const c : f32 = 1.0;
const d : f32 = 1.0;
@vertex
fn vert_main() -> @builtin(position) vec4<f32> {
let local_a : f32 = a;
return vec4<f32>();
}
@fragment
fn frag_main() {
let local_b : f32 = b;
}
@compute @workgroup_size(1)
fn comp_main1() {
let local_c : f32 = c;
}
@compute @workgroup_size(1)
fn comp_main2() {
let local_d : f32 = d;
}
)";
auto* expect = R"(
const a : f32 = 1.0;
const b : f32 = 1.0;
const c : f32 = 1.0;
const d : f32 = 1.0;
@compute @workgroup_size(1)
fn comp_main1() {
let local_c : f32 = c;
}
)";
SingleEntryPoint::Config cfg("comp_main1");
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
TEST_F(SingleEntryPointTest, WorkgroupSizeConstPreserved) {
auto* src = R"(
const size : i32 = 1;
@compute @workgroup_size(size)
fn main() {
}
)";
auto* expect = src;
SingleEntryPoint::Config cfg("main");
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
TEST_F(SingleEntryPointTest, OverridableConstants) {
auto* src = R"(
@id(1001) override c1 : u32 = 1u;
override c2 : u32 = 1u;
@id(0) override c3 : u32 = 1u;
@id(9999) override c4 : u32 = 1u;
@compute @workgroup_size(1)
fn comp_main1() {
let local_d = c1;
}
@compute @workgroup_size(1)
fn comp_main2() {
let local_d = c2;
}
@compute @workgroup_size(1)
fn comp_main3() {
let local_d = c3;
}
@compute @workgroup_size(1)
fn comp_main4() {
let local_d = c4;
}
@compute @workgroup_size(1)
fn comp_main5() {
let local_d = 1u;
}
)";
{
SingleEntryPoint::Config cfg("comp_main1");
auto* expect = R"(
@id(1001) override c1 : u32 = 1u;
@compute @workgroup_size(1)
fn comp_main1() {
let local_d = c1;
}
)";
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
{
SingleEntryPoint::Config cfg("comp_main2");
// The decorator is replaced with the one with explicit id
// And should not be affected by other constants stripped away
auto* expect = R"(
@id(1) override c2 : u32 = 1u;
@compute @workgroup_size(1)
fn comp_main2() {
let local_d = c2;
}
)";
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
{
SingleEntryPoint::Config cfg("comp_main3");
auto* expect = R"(
@id(0) override c3 : u32 = 1u;
@compute @workgroup_size(1)
fn comp_main3() {
let local_d = c3;
}
)";
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
{
SingleEntryPoint::Config cfg("comp_main4");
auto* expect = R"(
@id(9999) override c4 : u32 = 1u;
@compute @workgroup_size(1)
fn comp_main4() {
let local_d = c4;
}
)";
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
{
SingleEntryPoint::Config cfg("comp_main5");
auto* expect = R"(
@compute @workgroup_size(1)
fn comp_main5() {
let local_d = 1u;
}
)";
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
}
TEST_F(SingleEntryPointTest, OverridableConstants_TransitiveUses) {
// Make sure we do not strip away transitive uses of overridable constants.
auto* src = R"(
@id(0) override c0 : u32;
@id(1) override c1 : u32 = (2 * c0);
@id(2) override c2 : u32;
@id(3) override c3 : u32 = (2 * c2);
@id(4) override c4 : u32;
@id(5) override c5 : u32 = (2 * c4);
alias arr_ty = array<i32, (2 * c5)>;
var<workgroup> arr : arr_ty;
@compute @workgroup_size(1, 1, (2 * c3))
fn main() {
let local_d = c1;
arr[0] = 42;
}
)";
auto* expect = src;
SingleEntryPoint::Config cfg("main");
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
TEST_F(SingleEntryPointTest, OverridableConstants_UnusedAliasForOverrideSizedArray) {
// Make sure we strip away aliases that reference unused overridable constants.
auto* src = R"(
@id(0) override c0 : u32;
// This is all unused by the target entry point.
@id(1) override c1 : u32;
alias arr_ty = array<i32, c1>;
var<workgroup> arr : arr_ty;
@compute @workgroup_size(64)
fn unused() {
arr[0] = 42;
}
@compute @workgroup_size(64)
fn main() {
let local_d = c0;
}
)";
auto* expect = R"(
@id(0) override c0 : u32;
@compute @workgroup_size(64)
fn main() {
let local_d = c0;
}
)";
SingleEntryPoint::Config cfg("main");
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
TEST_F(SingleEntryPointTest, CalledFunctions) {
auto* src = R"(
fn inner1() {
}
fn inner2() {
}
fn inner_shared() {
}
fn outer1() {
inner1();
inner_shared();
}
fn outer2() {
inner2();
inner_shared();
}
@compute @workgroup_size(1)
fn comp_main1() {
outer1();
}
@compute @workgroup_size(1)
fn comp_main2() {
outer2();
}
)";
auto* expect = R"(
fn inner1() {
}
fn inner_shared() {
}
fn outer1() {
inner1();
inner_shared();
}
@compute @workgroup_size(1)
fn comp_main1() {
outer1();
}
)";
SingleEntryPoint::Config cfg("comp_main1");
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
TEST_F(SingleEntryPointTest, GlobalsReferencedByCalledFunctions) {
auto* src = R"(
var<private> inner1_var : f32;
var<private> inner2_var : f32;
var<private> inner_shared_var : f32;
var<private> outer1_var : f32;
var<private> outer2_var : f32;
fn inner1() {
inner1_var = 0.0;
}
fn inner2() {
inner2_var = 0.0;
}
fn inner_shared() {
inner_shared_var = 0.0;
}
fn outer1() {
inner1();
inner_shared();
outer1_var = 0.0;
}
fn outer2() {
inner2();
inner_shared();
outer2_var = 0.0;
}
@compute @workgroup_size(1)
fn comp_main1() {
outer1();
}
@compute @workgroup_size(1)
fn comp_main2() {
outer2();
}
)";
auto* expect = R"(
var<private> inner1_var : f32;
var<private> inner_shared_var : f32;
var<private> outer1_var : f32;
fn inner1() {
inner1_var = 0.0;
}
fn inner_shared() {
inner_shared_var = 0.0;
}
fn outer1() {
inner1();
inner_shared();
outer1_var = 0.0;
}
@compute @workgroup_size(1)
fn comp_main1() {
outer1();
}
)";
SingleEntryPoint::Config cfg("comp_main1");
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
TEST_F(SingleEntryPointTest, GlobalConstUsedAsArraySize) {
// See crbug.com/tint/1598
auto* src = R"(
const MY_SIZE = 5u;
alias Arr = array<i32, MY_SIZE>;
@fragment
fn main() {
}
)";
auto* expect = src;
SingleEntryPoint::Config cfg("main");
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
TEST_F(SingleEntryPointTest, Directives) {
// Make sure that directives are preserved.
auto* src = R"(
enable f16;
diagnostic(off, derivative_uniformity);
@compute @workgroup_size(1)
fn main() {
}
)";
SingleEntryPoint::Config cfg("main");
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(src, str(got));
}
TEST_F(SingleEntryPointTest, Requires) {
// Make sure that requires are handled (and dropped).
auto* src = R"(
requires readonly_and_readwrite_storage_textures;
@compute @workgroup_size(1)
fn main() {
}
)";
auto* expect = R"(
@compute @workgroup_size(1)
fn main() {
}
)";
SingleEntryPoint::Config cfg("main");
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
TEST_F(SingleEntryPointTest, ConstAssert_ModuleScope) {
// module scope const_assert is preserved
auto* src = R"(
const C = 42;
const_assert (C == 42);
@compute @workgroup_size(1)
fn main() {
}
)";
auto* expect = src;
SingleEntryPoint::Config cfg("main");
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
TEST_F(SingleEntryPointTest, ConstAssert_FnScope) {
// function scope const_assert is preserved
auto* src = R"(
const C = 42;
@compute @workgroup_size(1)
fn main() {
const_assert (C == 42);
}
)";
auto* expect = src;
SingleEntryPoint::Config cfg("main");
DataMap data;
data.Add<SingleEntryPoint::Config>(cfg);
auto got = Run<SingleEntryPoint>(src, data);
EXPECT_EQ(expect, str(got));
}
} // namespace
} // namespace tint::ast::transform