From 02025311594d7bbcca0fb0d0abd1c44c542c9379 Mon Sep 17 00:00:00 2001 From: dan sinclair Date: Fri, 26 May 2023 13:14:44 +0000 Subject: [PATCH] [ir] Add a NextIteration instruction. This CL adds a `NextIteration` instruction in order to branch from a continue block back to the start of a loop. The `next_iteration` is clearer then the `break-if false` pattern that was there previously. BreakIf is retained and used when an `ast::BreakIf` is encountered as it's clearer then the replaced `if` structure. Bug: tint:1718 Change-Id: Ie6ce0db51c244866e2e99118bc00e4cfd2b3dc74 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/134600 Kokoro: Kokoro Reviewed-by: Ben Clayton Commit-Queue: Ben Clayton Reviewed-by: James Price --- src/tint/BUILD.gn | 2 ++ src/tint/CMakeLists.txt | 2 ++ src/tint/ir/builder.cc | 4 +++ src/tint/ir/builder.h | 6 +++++ src/tint/ir/disassembler.cc | 4 +++ src/tint/ir/from_program.cc | 6 ++--- src/tint/ir/from_program_test.cc | 37 ++++++++++++-------------- src/tint/ir/next_iteration.cc | 31 ++++++++++++++++++++++ src/tint/ir/next_iteration.h | 45 ++++++++++++++++++++++++++++++++ 9 files changed, 113 insertions(+), 24 deletions(-) create mode 100644 src/tint/ir/next_iteration.cc create mode 100644 src/tint/ir/next_iteration.h diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn index bff6c37886..5970683500 100644 --- a/src/tint/BUILD.gn +++ b/src/tint/BUILD.gn @@ -1254,6 +1254,8 @@ if (tint_build_ir) { "ir/loop.h", "ir/module.cc", "ir/module.h", + "ir/next_iteration.cc", + "ir/next_iteration.h", "ir/return.cc", "ir/return.h", "ir/store.cc", diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt index e5e529b10b..17fa2d3fab 100644 --- a/src/tint/CMakeLists.txt +++ b/src/tint/CMakeLists.txt @@ -764,6 +764,8 @@ if(${TINT_BUILD_IR}) ir/loop.h ir/module.cc ir/module.h + ir/next_iteration.cc + ir/next_iteration.h ir/return.cc ir/return.h ir/store.cc diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc index 1daf4583ae..7f2d9c3462 100644 --- a/src/tint/ir/builder.cc +++ b/src/tint/ir/builder.cc @@ -205,6 +205,10 @@ ir::Return* Builder::Return(Function* func, utils::VectorRef args) { return ir.values.Create(func, args); } +ir::NextIteration* Builder::NextIteration(Loop* loop) { + return ir.values.Create(loop); +} + ir::BreakIf* Builder::BreakIf(Value* condition, Loop* loop) { return ir.values.Create(condition, loop); } diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h index 8a4c0213ed..b544cb285d 100644 --- a/src/tint/ir/builder.h +++ b/src/tint/ir/builder.h @@ -37,6 +37,7 @@ #include "src/tint/ir/load.h" #include "src/tint/ir/loop.h" #include "src/tint/ir/module.h" +#include "src/tint/ir/next_iteration.h" #include "src/tint/ir/return.h" #include "src/tint/ir/store.h" #include "src/tint/ir/switch.h" @@ -337,6 +338,11 @@ class Builder { /// @returns the instruction ir::Return* Return(Function* func, utils::VectorRef args = {}); + /// Creates a loop next iteration instruction + /// @param loop the loop being iterated + /// @returns the instruction + ir::NextIteration* NextIteration(Loop* loop); + /// Creates a loop break-if instruction /// @param condition the break condition /// @param loop the loop being iterated diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc index d39438aec3..75cc2aa860 100644 --- a/src/tint/ir/disassembler.cc +++ b/src/tint/ir/disassembler.cc @@ -33,6 +33,7 @@ #include "src/tint/ir/if.h" #include "src/tint/ir/load.h" #include "src/tint/ir/loop.h" +#include "src/tint/ir/next_iteration.h" #include "src/tint/ir/return.h" #include "src/tint/ir/store.h" #include "src/tint/ir/switch.h" @@ -428,6 +429,9 @@ void Disassembler::EmitBranch(const Branch* b) { [&](const ir::ExitIf* ei) { out_ << "exit_if %b" << IdOf(ei->If()->Merge()); }, [&](const ir::ExitSwitch* es) { out_ << "exit_switch %b" << IdOf(es->Switch()->Merge()); }, [&](const ir::ExitLoop* el) { out_ << "exit_loop %b" << IdOf(el->Loop()->Merge()); }, + [&](const ir::NextIteration* ni) { + out_ << "next_iteration %b" << IdOf(ni->Loop()->Start()); + }, [&](const ir::BreakIf* bi) { out_ << "break_if "; EmitValue(bi->Condition()); diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc index 67c9f3f53c..4773f3fcbb 100644 --- a/src/tint/ir/from_program.cc +++ b/src/tint/ir/from_program.cc @@ -575,7 +575,7 @@ class Impl { } // Branch back to the start node if the continue target didn't branch out already if (NeedBranch()) { - SetBranch(builder_.BreakIf(builder_.Constant(false), loop_inst)); + SetBranch(builder_.NextIteration(loop_inst)); } } } @@ -595,7 +595,7 @@ class Impl { // Continue is always empty, just go back to the start current_flow_block_ = loop_inst->Continuing(); - SetBranch(builder_.BreakIf(builder_.Constant(false), loop_inst)); + SetBranch(builder_.NextIteration(loop_inst)); { ControlStackScope scope(this, loop_inst); @@ -676,7 +676,7 @@ class Impl { if (stmt->continuing) { current_flow_block_ = loop_inst->Continuing(); EmitStatement(stmt->continuing); - SetBranch(builder_.BreakIf(builder_.Constant(false), loop_inst)); + SetBranch(builder_.NextIteration(loop_inst)); } } diff --git a/src/tint/ir/from_program_test.cc b/src/tint/ir/from_program_test.cc index f2b9c19aaf..f9be6680f0 100644 --- a/src/tint/ir/from_program_test.cc +++ b/src/tint/ir/from_program_test.cc @@ -381,7 +381,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) { EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length()); EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length()); - EXPECT_EQ(2u, loop_flow->Merge()->InboundBranches().Length()); + EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length()); @@ -411,7 +411,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) { # Continuing block %b3 = block { - break_if false %b2 + next_iteration %b2 } # Merge block @@ -513,7 +513,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) { EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length()); EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length()); - EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length()); + EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length()); @@ -521,21 +521,21 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) { EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { %b1 = block { - loop [s: %b2, c: %b3, m: %b4] + loop [s: %b2, c: %b3] %b2 = block { - if true [t: %b5, f: %b6, m: %b7] + if true [t: %b4, f: %b5, m: %b6] # True block - %b5 = block { + %b4 = block { ret } # False block - %b6 = block { - exit_if %b7 + %b5 = block { + exit_if %b6 } # Merge block - %b7 = block { + %b6 = block { continue %b3 } @@ -543,14 +543,9 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) { # Continuing block %b3 = block { - break_if false %b2 + next_iteration %b2 } - # Merge block - %b4 = block { - ret - } - } } )"); @@ -750,7 +745,7 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) { # Merge block %b18 = block { - break_if false %b5 + next_iteration %b5 } } @@ -781,7 +776,7 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) { # Continuing block %b3 = block { - break_if false %b2 + next_iteration %b2 } # Merge block @@ -812,7 +807,7 @@ TEST_F(IR_BuilderImplTest, While) { EXPECT_EQ(1u, flow->Start()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Continuing()->InboundBranches().Length()); - EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); + EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length()); @@ -842,7 +837,7 @@ TEST_F(IR_BuilderImplTest, While) { # Continuing block %b3 = block { - break_if false %b2 + next_iteration %b2 } # Merge block @@ -873,7 +868,7 @@ TEST_F(IR_BuilderImplTest, While_Return) { EXPECT_EQ(1u, flow->Start()->InboundBranches().Length()); EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length()); - EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); + EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length()); @@ -903,7 +898,7 @@ TEST_F(IR_BuilderImplTest, While_Return) { # Continuing block %b3 = block { - break_if false %b2 + next_iteration %b2 } # Merge block diff --git a/src/tint/ir/next_iteration.cc b/src/tint/ir/next_iteration.cc new file mode 100644 index 0000000000..0c021eb440 --- /dev/null +++ b/src/tint/ir/next_iteration.cc @@ -0,0 +1,31 @@ +// Copyright 2023 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/ir/next_iteration.h" + +#include "src/tint/ir/loop.h" + +TINT_INSTANTIATE_TYPEINFO(tint::ir::NextIteration); + +namespace tint::ir { + +NextIteration::NextIteration(ir::Loop* loop) : Base(utils::Empty), loop_(loop) { + TINT_ASSERT(IR, loop_); + loop_->AddUsage(this); + loop_->Start()->AddInboundBranch(this); +} + +NextIteration::~NextIteration() = default; + +} // namespace tint::ir diff --git a/src/tint/ir/next_iteration.h b/src/tint/ir/next_iteration.h new file mode 100644 index 0000000000..f1211e34c2 --- /dev/null +++ b/src/tint/ir/next_iteration.h @@ -0,0 +1,45 @@ +// Copyright 2023 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_IR_NEXT_ITERATION_H_ +#define SRC_TINT_IR_NEXT_ITERATION_H_ + +#include "src/tint/ir/branch.h" +#include "src/tint/utils/castable.h" + +// Forward declarations +namespace tint::ir { +class Loop; +} // namespace tint::ir + +namespace tint::ir { + +/// A next iteration instruction. +class NextIteration : public utils::Castable { + public: + /// Constructor + /// @param loop the loop being iterated + explicit NextIteration(ir::Loop* loop); + ~NextIteration() override; + + /// @returns the loop being iterated + const ir::Loop* Loop() const { return loop_; } + + private: + ir::Loop* loop_ = nullptr; +}; + +} // namespace tint::ir + +#endif // SRC_TINT_IR_NEXT_ITERATION_H_