tint: Add MergeReturn transform
Add a transform that merges return statements into a single return statement at the end of the function. Bug: tint:1627 Change-Id: I971cc298fd9814634c82b49a07e15c5c0f3da404 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/93660 Kokoro: Kokoro <noreply+kokoro@google.com> Reviewed-by: Ben Clayton <bclayton@google.com> Commit-Queue: James Price <jrprice@google.com>
This commit is contained in:
James Price
Dawn LUCI CQ
parent
5ff6a6a496
commit
5de8b436fa
|
|
@ -511,6 +511,8 @@ libtint_source_set("libtint_core_all_src") {
|
|||
"transform/localize_struct_array_assignment.h",
|
||||
"transform/manager.cc",
|
||||
"transform/manager.h",
|
||||
"transform/merge_return.cc",
|
||||
"transform/merge_return.h",
|
||||
"transform/module_scope_var_to_entry_point_param.cc",
|
||||
"transform/module_scope_var_to_entry_point_param.h",
|
||||
"transform/multiplanar_external_texture.cc",
|
||||
|
|
@ -1224,6 +1226,7 @@ if (tint_build_unittests) {
|
|||
"transform/first_index_offset_test.cc",
|
||||
"transform/for_loop_to_loop_test.cc",
|
||||
"transform/localize_struct_array_assignment_test.cc",
|
||||
"transform/merge_return_test.cc",
|
||||
"transform/module_scope_var_to_entry_point_param_test.cc",
|
||||
"transform/multiplanar_external_texture_test.cc",
|
||||
"transform/num_workgroups_from_uniform_test.cc",
|
||||
|
|
|
|||
|
|
@ -444,6 +444,8 @@ set(TINT_LIB_SRCS
|
|||
transform/localize_struct_array_assignment.h
|
||||
transform/manager.cc
|
||||
transform/manager.h
|
||||
transform/merge_return.cc
|
||||
transform/merge_return.h
|
||||
transform/module_scope_var_to_entry_point_param.cc
|
||||
transform/module_scope_var_to_entry_point_param.h
|
||||
transform/multiplanar_external_texture.cc
|
||||
|
|
@ -1162,6 +1164,7 @@ if(TINT_BUILD_TESTS)
|
|||
transform/for_loop_to_loop_test.cc
|
||||
transform/expand_compound_assignment.cc
|
||||
transform/localize_struct_array_assignment_test.cc
|
||||
transform/merge_return.cc
|
||||
transform/module_scope_var_to_entry_point_param_test.cc
|
||||
transform/multiplanar_external_texture_test.cc
|
||||
transform/num_workgroups_from_uniform_test.cc
|
||||
|
|
|
|||
|
|
@ -0,0 +1,239 @@
|
|||
// 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/merge_return.h"
|
||||
|
||||
#include <utility>
|
||||
|
||||
#include "src/tint/program_builder.h"
|
||||
#include "src/tint/sem/statement.h"
|
||||
#include "src/tint/utils/scoped_assignment.h"
|
||||
|
||||
TINT_INSTANTIATE_TYPEINFO(tint::transform::MergeReturn);
|
||||
|
||||
using namespace tint::number_suffixes; // NOLINT
|
||||
|
||||
namespace tint::transform {
|
||||
|
||||
namespace {
|
||||
|
||||
/// Returns `true` if `stmt` has the behavior `behavior`.
|
||||
bool HasBehavior(const Program* program, const ast::Statement* stmt, sem::Behavior behavior) {
|
||||
return program->Sem().Get(stmt)->Behaviors().Contains(behavior);
|
||||
}
|
||||
|
||||
/// Returns `true` if `func` needs to be transformed.
|
||||
bool NeedsTransform(const Program* program, const ast::Function* func) {
|
||||
// Entry points and intrinsic declarations never need transforming.
|
||||
if (func->IsEntryPoint() || func->body == nullptr) {
|
||||
return false;
|
||||
}
|
||||
|
||||
// Avoid transforming functions that only have a single exit point.
|
||||
// TODO(jrprice): Alternatively, use the uniformity analysis to decide which
|
||||
// functions need to be transformed.
|
||||
for (auto* s : func->body->statements) {
|
||||
// Find the first statement that has contains the Return behavior.
|
||||
if (HasBehavior(program, s, sem::Behavior::kReturn)) {
|
||||
// If this statement is itself a return, it will be the only exit point,
|
||||
// so no need to apply the transform to the function.
|
||||
if (s->Is<ast::ReturnStatement>()) {
|
||||
return false;
|
||||
} else {
|
||||
// Apply the transform in all other cases.
|
||||
return true;
|
||||
}
|
||||
}
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
} // namespace
|
||||
|
||||
MergeReturn::MergeReturn() = default;
|
||||
|
||||
MergeReturn::~MergeReturn() = default;
|
||||
|
||||
bool MergeReturn::ShouldRun(const Program* program, const DataMap&) const {
|
||||
for (auto* func : program->AST().Functions()) {
|
||||
if (NeedsTransform(program, func)) {
|
||||
return true;
|
||||
}
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
namespace {
|
||||
|
||||
/// Internal class used to during the transform.
|
||||
class State {
|
||||
private:
|
||||
/// The clone context.
|
||||
CloneContext& ctx;
|
||||
|
||||
/// The program builder.
|
||||
ProgramBuilder& b;
|
||||
|
||||
/// The function.
|
||||
const ast::Function* function;
|
||||
|
||||
/// The symbol for the return flag variable.
|
||||
Symbol flag;
|
||||
|
||||
/// The symbol for the return value variable.
|
||||
Symbol retval;
|
||||
|
||||
/// Tracks whether we are currently inside a loop or switch statement.
|
||||
bool is_in_loop_or_switch = false;
|
||||
|
||||
public:
|
||||
/// Constructor
|
||||
/// @param context the clone context
|
||||
State(CloneContext& context, const ast::Function* func)
|
||||
: ctx(context), b(*ctx.dst), function(func) {}
|
||||
|
||||
/// Process a statement (recursively).
|
||||
void ProcessStatement(const ast::Statement* stmt) {
|
||||
if (stmt == nullptr || !HasBehavior(ctx.src, stmt, sem::Behavior::kReturn)) {
|
||||
return;
|
||||
}
|
||||
|
||||
Switch(
|
||||
stmt, [&](const ast::BlockStatement* block) { ProcessBlock(block); },
|
||||
[&](const ast::CaseStatement* c) { ProcessStatement(c->body); },
|
||||
[&](const ast::ForLoopStatement* f) {
|
||||
TINT_SCOPED_ASSIGNMENT(is_in_loop_or_switch, true);
|
||||
ProcessStatement(f->body);
|
||||
},
|
||||
[&](const ast::IfStatement* i) {
|
||||
ProcessStatement(i->body);
|
||||
ProcessStatement(i->else_statement);
|
||||
},
|
||||
[&](const ast::LoopStatement* l) {
|
||||
TINT_SCOPED_ASSIGNMENT(is_in_loop_or_switch, true);
|
||||
ProcessStatement(l->body);
|
||||
},
|
||||
[&](const ast::ReturnStatement* r) {
|
||||
utils::Vector<const ast::Statement*, 3> stmts;
|
||||
// Set the return flag to signal that we have hit a return.
|
||||
stmts.Push(b.Assign(b.Expr(flag), true));
|
||||
if (r->value) {
|
||||
// Set the return value if necessary.
|
||||
stmts.Push(b.Assign(b.Expr(retval), ctx.Clone(r->value)));
|
||||
}
|
||||
if (is_in_loop_or_switch) {
|
||||
// If we are in a loop or switch statement, break out of it.
|
||||
stmts.Push(b.Break());
|
||||
}
|
||||
ctx.Replace(r, b.Block(std::move(stmts)));
|
||||
},
|
||||
[&](const ast::SwitchStatement* s) {
|
||||
TINT_SCOPED_ASSIGNMENT(is_in_loop_or_switch, true);
|
||||
for (auto* c : s->body) {
|
||||
ProcessStatement(c);
|
||||
}
|
||||
},
|
||||
[&](const ast::WhileStatement* w) {
|
||||
TINT_SCOPED_ASSIGNMENT(is_in_loop_or_switch, true);
|
||||
ProcessStatement(w->body);
|
||||
},
|
||||
[&](Default) { TINT_ICE(Transform, b.Diagnostics()) << "unhandled statement type"; });
|
||||
}
|
||||
|
||||
void ProcessBlock(const ast::BlockStatement* block) {
|
||||
// We will rebuild the contents of the block statement.
|
||||
// We may introduce conditionals around statements that follow a statement with the
|
||||
// `Return` behavior, so build a stack of statement lists that represent the new
|
||||
// (potentially nested) conditional blocks.
|
||||
utils::Vector<utils::Vector<const ast::Statement*, 8>, 8> new_stmts({{}});
|
||||
|
||||
// Insert variables for the return flag and return value at the top of the function.
|
||||
if (block == function->body) {
|
||||
flag = b.Symbols().New("tint_return_flag");
|
||||
new_stmts[0].Push(b.Decl(b.Var(flag, b.ty.bool_())));
|
||||
|
||||
if (!function->return_type->Is<ast::Void>()) {
|
||||
retval = b.Symbols().New("tint_return_value");
|
||||
new_stmts[0].Push(b.Decl(b.Var(retval, ctx.Clone(function->return_type))));
|
||||
}
|
||||
}
|
||||
|
||||
for (auto* s : block->statements) {
|
||||
// Process the statement and add it to the current block.
|
||||
ProcessStatement(s);
|
||||
new_stmts.Back().Push(ctx.Clone(s));
|
||||
|
||||
// Check if the statement is or contains a return statement.
|
||||
// We need to make sure any statements that follow this one do not get executed if the
|
||||
// return flag has been set.
|
||||
if (HasBehavior(ctx.src, s, sem::Behavior::kReturn)) {
|
||||
if (is_in_loop_or_switch) {
|
||||
// We're in a loop/switch, and so we would have inserted a `break`.
|
||||
// If we've just come out of a loop/switch statement, we need to `break` again.
|
||||
if (s->IsAnyOf<ast::LoopStatement, ast::ForLoopStatement,
|
||||
ast::SwitchStatement>()) {
|
||||
// If the loop only has the 'Return' behavior, we can just unconditionally
|
||||
// break. Otherwise check the return flag.
|
||||
if (HasBehavior(ctx.src, s, sem::Behavior::kNext)) {
|
||||
new_stmts.Back().Push(
|
||||
b.If(b.Expr(flag), b.Block(utils::Vector{b.Break()})));
|
||||
} else {
|
||||
new_stmts.Back().Push(b.Break());
|
||||
}
|
||||
}
|
||||
} else {
|
||||
// Create a new list for any subsequent statements, which we will wrap in a
|
||||
// conditional block.
|
||||
new_stmts.Push({});
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// Descend the stack of new block statements, wrapping them in conditionals.
|
||||
while (new_stmts.Length() > 1) {
|
||||
const ast::IfStatement* i = nullptr;
|
||||
if (new_stmts.Back().Length() > 0) {
|
||||
i = b.If(b.Not(b.Expr(flag)), b.Block(new_stmts.Back()));
|
||||
}
|
||||
new_stmts.Pop();
|
||||
if (i) {
|
||||
new_stmts.Back().Push(i);
|
||||
}
|
||||
}
|
||||
|
||||
// Insert the final return statement at the end of the function body.
|
||||
if (block == function->body && retval.IsValid()) {
|
||||
new_stmts[0].Push(b.Return(b.Expr(retval)));
|
||||
}
|
||||
|
||||
ctx.Replace(block, b.Block(new_stmts[0]));
|
||||
}
|
||||
};
|
||||
|
||||
} // namespace
|
||||
|
||||
void MergeReturn::Run(CloneContext& ctx, const DataMap&, DataMap&) const {
|
||||
for (auto* func : ctx.src->AST().Functions()) {
|
||||
if (!NeedsTransform(ctx.src, func)) {
|
||||
continue;
|
||||
}
|
||||
|
||||
State state(ctx, func);
|
||||
state.ProcessStatement(func->body);
|
||||
}
|
||||
|
||||
ctx.Clone();
|
||||
}
|
||||
|
||||
} // namespace tint::transform
|
||||
|
|
@ -0,0 +1,47 @@
|
|||
// 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.
|
||||
|
||||
#ifndef SRC_TINT_TRANSFORM_MERGE_RETURN_H_
|
||||
#define SRC_TINT_TRANSFORM_MERGE_RETURN_H_
|
||||
|
||||
#include "src/tint/transform/transform.h"
|
||||
|
||||
namespace tint::transform {
|
||||
|
||||
/// Merge return statements into a single return at the end of the function.
|
||||
class MergeReturn final : public Castable<MergeReturn, Transform> {
|
||||
public:
|
||||
/// Constructor
|
||||
MergeReturn();
|
||||
/// Destructor
|
||||
~MergeReturn() override;
|
||||
|
||||
/// @param program the program to inspect
|
||||
/// @param data optional extra transform-specific input data
|
||||
/// @returns true if this transform should be run for the given program
|
||||
bool ShouldRun(const Program* program, const DataMap& data = {}) const override;
|
||||
|
||||
protected:
|
||||
/// Runs the transform using the CloneContext built for transforming a
|
||||
/// program. Run() is responsible for calling Clone() on the CloneContext.
|
||||
/// @param ctx the CloneContext primed with the input program and
|
||||
/// ProgramBuilder
|
||||
/// @param inputs optional extra transform-specific input data
|
||||
/// @param outputs optional extra transform-specific output data
|
||||
void Run(CloneContext& ctx, const DataMap& inputs, DataMap& outputs) const override;
|
||||
};
|
||||
|
||||
} // namespace tint::transform
|
||||
|
||||
#endif // SRC_TINT_TRANSFORM_MERGE_RETURN_H_
|
||||
|
|
@ -0,0 +1,862 @@
|
|||
// 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/merge_return.h"
|
||||
|
||||
#include <utility>
|
||||
|
||||
#include "src/tint/transform/test_helper.h"
|
||||
|
||||
namespace tint::transform {
|
||||
namespace {
|
||||
|
||||
using MergeReturnTest = TransformTest;
|
||||
|
||||
TEST_F(MergeReturnTest, ShouldRunEmptyModule) {
|
||||
auto* src = R"()";
|
||||
|
||||
EXPECT_FALSE(ShouldRun<MergeReturn>(src));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, ShouldRunOnlyEntryPoints) {
|
||||
auto* src = R"(
|
||||
@fragment
|
||||
fn foo() {}
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn bar() {}
|
||||
)";
|
||||
|
||||
EXPECT_FALSE(ShouldRun<MergeReturn>(src));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, ShouldRunUserFuncNoReturn) {
|
||||
auto* src = R"(
|
||||
fn foo() {}
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn bar() {
|
||||
foo();
|
||||
}
|
||||
)";
|
||||
|
||||
EXPECT_FALSE(ShouldRun<MergeReturn>(src));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, ShouldRunUserFuncSingleReturn) {
|
||||
auto* src = R"(
|
||||
fn foo() -> i32 {
|
||||
return 42;
|
||||
}
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn bar() {
|
||||
foo();
|
||||
}
|
||||
)";
|
||||
|
||||
EXPECT_FALSE(ShouldRun<MergeReturn>(src));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, ShouldRunUserFuncMultipleReturn) {
|
||||
auto* src = R"(
|
||||
fn foo(x : i32) -> i32 {
|
||||
if (x == 0) {
|
||||
return 42;
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
|
||||
@compute @workgroup_size(1)
|
||||
fn bar() {
|
||||
foo(7);
|
||||
}
|
||||
)";
|
||||
|
||||
EXPECT_TRUE(ShouldRun<MergeReturn>(src));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, EmptyModule) {
|
||||
auto* src = R"()";
|
||||
|
||||
auto* expect = src;
|
||||
|
||||
auto got = Run<MergeReturn>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, SingleExitPoint) {
|
||||
auto* src = R"(
|
||||
fn foo(x : i32) -> i32 {
|
||||
var retval = 0;
|
||||
if ((x == 0)) {
|
||||
retval = 42;
|
||||
} else {
|
||||
retval = 7;
|
||||
}
|
||||
return retval;
|
||||
}
|
||||
)";
|
||||
|
||||
auto* expect = src;
|
||||
|
||||
auto got = Run<MergeReturn>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, Basic) {
|
||||
auto* src = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
fn foo() {
|
||||
if (non_uniform_global == 0) {
|
||||
return;
|
||||
} else {
|
||||
return;
|
||||
}
|
||||
}
|
||||
)";
|
||||
|
||||
auto* expect = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
fn foo() {
|
||||
var tint_return_flag : bool;
|
||||
if ((non_uniform_global == 0)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
}
|
||||
} else {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
}
|
||||
}
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<MergeReturn>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, BasicRetval) {
|
||||
auto* src = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
if (non_uniform_global == 0) {
|
||||
return 1;
|
||||
} else {
|
||||
return 2;
|
||||
}
|
||||
}
|
||||
)";
|
||||
|
||||
auto* expect = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
var tint_return_flag : bool;
|
||||
var tint_return_value : i32;
|
||||
if ((non_uniform_global == 0)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 1;
|
||||
}
|
||||
} else {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 2;
|
||||
}
|
||||
}
|
||||
return tint_return_value;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<MergeReturn>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, CodeAfterIf) {
|
||||
auto* src = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
var<private> bar : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
if (non_uniform_global == 0) {
|
||||
return 1;
|
||||
}
|
||||
bar = 42;
|
||||
return 2;
|
||||
}
|
||||
)";
|
||||
|
||||
auto* expect = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
var<private> bar : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
var tint_return_flag : bool;
|
||||
var tint_return_value : i32;
|
||||
if ((non_uniform_global == 0)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 1;
|
||||
}
|
||||
}
|
||||
if (!(tint_return_flag)) {
|
||||
bar = 42;
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 2;
|
||||
}
|
||||
}
|
||||
return tint_return_value;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<MergeReturn>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, SequenceOfConditionalReturns) {
|
||||
auto* src = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
var<private> bar : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
if (non_uniform_global == 0) {
|
||||
return 1;
|
||||
}
|
||||
bar = 42;
|
||||
if (non_uniform_global == 1) {
|
||||
return 2;
|
||||
}
|
||||
bar = 43;
|
||||
if (non_uniform_global == 2) {
|
||||
return 3;
|
||||
}
|
||||
bar = 44;
|
||||
if (non_uniform_global == 3) {
|
||||
return 4;
|
||||
}
|
||||
bar = 45;
|
||||
return 0;
|
||||
}
|
||||
)";
|
||||
|
||||
auto* expect = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
var<private> bar : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
var tint_return_flag : bool;
|
||||
var tint_return_value : i32;
|
||||
if ((non_uniform_global == 0)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 1;
|
||||
}
|
||||
}
|
||||
if (!(tint_return_flag)) {
|
||||
bar = 42;
|
||||
if ((non_uniform_global == 1)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 2;
|
||||
}
|
||||
}
|
||||
if (!(tint_return_flag)) {
|
||||
bar = 43;
|
||||
if ((non_uniform_global == 2)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 3;
|
||||
}
|
||||
}
|
||||
if (!(tint_return_flag)) {
|
||||
bar = 44;
|
||||
if ((non_uniform_global == 3)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 4;
|
||||
}
|
||||
}
|
||||
if (!(tint_return_flag)) {
|
||||
bar = 45;
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 0;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
return tint_return_value;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<MergeReturn>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, NestedConditionals) {
|
||||
auto* src = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
var<private> bar : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
if (non_uniform_global == 0) {
|
||||
if (true) {
|
||||
return 1;
|
||||
} else {
|
||||
return 2;
|
||||
}
|
||||
} else {
|
||||
if (true) {
|
||||
bar = 42;
|
||||
} else if (false) {
|
||||
return 3;
|
||||
} else if (true) {
|
||||
return 4;
|
||||
}
|
||||
return 5;
|
||||
}
|
||||
}
|
||||
)";
|
||||
|
||||
auto* expect = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
var<private> bar : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
var tint_return_flag : bool;
|
||||
var tint_return_value : i32;
|
||||
if ((non_uniform_global == 0)) {
|
||||
if (true) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 1;
|
||||
}
|
||||
} else {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 2;
|
||||
}
|
||||
}
|
||||
} else {
|
||||
if (true) {
|
||||
bar = 42;
|
||||
} else if (false) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 3;
|
||||
}
|
||||
} else if (true) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 4;
|
||||
}
|
||||
}
|
||||
if (!(tint_return_flag)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 5;
|
||||
}
|
||||
}
|
||||
}
|
||||
return tint_return_value;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<MergeReturn>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, Loop) {
|
||||
auto* src = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
var<private> bar : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
var i = 0;
|
||||
loop {
|
||||
bar = i;
|
||||
if (non_uniform_global == 0) {
|
||||
return non_uniform_global;
|
||||
}
|
||||
if (i >= 10) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
)";
|
||||
|
||||
auto* expect = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
var<private> bar : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
var tint_return_flag : bool;
|
||||
var tint_return_value : i32;
|
||||
var i = 0;
|
||||
loop {
|
||||
bar = i;
|
||||
if ((non_uniform_global == 0)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = non_uniform_global;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if ((i >= 10)) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (!(tint_return_flag)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 0;
|
||||
}
|
||||
}
|
||||
return tint_return_value;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<MergeReturn>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, NestedLoop) {
|
||||
auto* src = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
var<private> bar : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
var i = 0;
|
||||
loop {
|
||||
loop {
|
||||
bar = i;
|
||||
if (non_uniform_global == 0) {
|
||||
return non_uniform_global;
|
||||
}
|
||||
if (i >= 10) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
bar = 10;
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
)";
|
||||
|
||||
auto* expect = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
var<private> bar : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
var tint_return_flag : bool;
|
||||
var tint_return_value : i32;
|
||||
var i = 0;
|
||||
loop {
|
||||
loop {
|
||||
bar = i;
|
||||
if ((non_uniform_global == 0)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = non_uniform_global;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if ((i >= 10)) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (tint_return_flag) {
|
||||
break;
|
||||
}
|
||||
bar = 10;
|
||||
}
|
||||
if (!(tint_return_flag)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 0;
|
||||
}
|
||||
}
|
||||
return tint_return_value;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<MergeReturn>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, NestedLoopAlwaysReturns) {
|
||||
auto* src = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
var<private> bar : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
var i = 0;
|
||||
loop {
|
||||
bar = i;
|
||||
loop {
|
||||
return non_uniform_global;
|
||||
}
|
||||
if (i == 10) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
)";
|
||||
|
||||
auto* expect = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
var<private> bar : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
var tint_return_flag : bool;
|
||||
var tint_return_value : i32;
|
||||
var i = 0;
|
||||
loop {
|
||||
bar = i;
|
||||
loop {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = non_uniform_global;
|
||||
break;
|
||||
}
|
||||
}
|
||||
break;
|
||||
if ((i == 10)) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (!(tint_return_flag)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 0;
|
||||
}
|
||||
}
|
||||
return tint_return_value;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<MergeReturn>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, ForLoop) {
|
||||
auto* src = R"(
|
||||
var<private> non_uniform_global : array<i32, 10>;
|
||||
|
||||
fn foo() -> i32 {
|
||||
for (var i = 0; i < 10; i++) {
|
||||
if (non_uniform_global[i] == 0) {
|
||||
return i;
|
||||
}
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
)";
|
||||
|
||||
auto* expect = R"(
|
||||
var<private> non_uniform_global : array<i32, 10>;
|
||||
|
||||
fn foo() -> i32 {
|
||||
var tint_return_flag : bool;
|
||||
var tint_return_value : i32;
|
||||
for(var i = 0; (i < 10); i++) {
|
||||
if ((non_uniform_global[i] == 0)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = i;
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
if (!(tint_return_flag)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 0;
|
||||
}
|
||||
}
|
||||
return tint_return_value;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<MergeReturn>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, WhileLoop) {
|
||||
auto* src = R"(
|
||||
var<private> non_uniform_global : array<i32, 10>;
|
||||
|
||||
fn foo() -> i32 {
|
||||
var i = 0;
|
||||
while (i < 10) {
|
||||
if (non_uniform_global[i] == 0) {
|
||||
return i;
|
||||
}
|
||||
i++;
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
)";
|
||||
|
||||
auto* expect = R"(
|
||||
var<private> non_uniform_global : array<i32, 10>;
|
||||
|
||||
fn foo() -> i32 {
|
||||
var tint_return_flag : bool;
|
||||
var tint_return_value : i32;
|
||||
var i = 0;
|
||||
while((i < 10)) {
|
||||
if ((non_uniform_global[i] == 0)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = i;
|
||||
break;
|
||||
}
|
||||
}
|
||||
i++;
|
||||
}
|
||||
if (!(tint_return_flag)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 0;
|
||||
}
|
||||
}
|
||||
return tint_return_value;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<MergeReturn>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, Switch) {
|
||||
auto* src = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
var<private> bar : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
switch (non_uniform_global) {
|
||||
case 0 {
|
||||
return 0;
|
||||
}
|
||||
case 1 {
|
||||
if (false) {
|
||||
return 1;
|
||||
}
|
||||
bar = 6;
|
||||
fallthrough;
|
||||
}
|
||||
case 2 {
|
||||
bar = 5;
|
||||
}
|
||||
default {
|
||||
return 42;
|
||||
}
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
)";
|
||||
|
||||
auto* expect = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
var<private> bar : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
var tint_return_flag : bool;
|
||||
var tint_return_value : i32;
|
||||
switch(non_uniform_global) {
|
||||
case 0: {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 0;
|
||||
break;
|
||||
}
|
||||
}
|
||||
case 1: {
|
||||
if (false) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 1;
|
||||
break;
|
||||
}
|
||||
}
|
||||
bar = 6;
|
||||
fallthrough;
|
||||
}
|
||||
case 2: {
|
||||
bar = 5;
|
||||
}
|
||||
default: {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 42;
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
if (!(tint_return_flag)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 0;
|
||||
}
|
||||
}
|
||||
return tint_return_value;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<MergeReturn>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
TEST_F(MergeReturnTest, ComplexNesting) {
|
||||
auto* src = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
var<private> bar : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
for (var i = 0; i < 10; i++) {
|
||||
bar = i;
|
||||
|
||||
switch (non_uniform_global) {
|
||||
case 0 {
|
||||
return 0;
|
||||
}
|
||||
case 1 {
|
||||
if (false) {
|
||||
loop {
|
||||
if (i == non_uniform_global) {
|
||||
return 1;
|
||||
}
|
||||
continuing {
|
||||
bar++;
|
||||
break if (i == 7);
|
||||
}
|
||||
}
|
||||
}
|
||||
bar = 0;
|
||||
}
|
||||
default {
|
||||
return 42;
|
||||
}
|
||||
}
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
)";
|
||||
|
||||
auto* expect = R"(
|
||||
var<private> non_uniform_global : i32;
|
||||
|
||||
var<private> bar : i32;
|
||||
|
||||
fn foo() -> i32 {
|
||||
var tint_return_flag : bool;
|
||||
var tint_return_value : i32;
|
||||
for(var i = 0; (i < 10); i++) {
|
||||
bar = i;
|
||||
switch(non_uniform_global) {
|
||||
case 0: {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 0;
|
||||
break;
|
||||
}
|
||||
}
|
||||
case 1: {
|
||||
if (false) {
|
||||
loop {
|
||||
if ((i == non_uniform_global)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 1;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
continuing {
|
||||
bar++;
|
||||
break if (i == 7);
|
||||
}
|
||||
}
|
||||
if (tint_return_flag) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
bar = 0;
|
||||
}
|
||||
default: {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 42;
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
if (tint_return_flag) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (!(tint_return_flag)) {
|
||||
{
|
||||
tint_return_flag = true;
|
||||
tint_return_value = 0;
|
||||
}
|
||||
}
|
||||
return tint_return_value;
|
||||
}
|
||||
)";
|
||||
|
||||
auto got = Run<MergeReturn>(src);
|
||||
|
||||
EXPECT_EQ(expect, str(got));
|
||||
}
|
||||
|
||||
} // namespace
|
||||
} // namespace tint::transform
|
||||
Loading…
Reference in New Issue