[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 <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index bff6c37..5970683 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -1254,6 +1254,8 @@
"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 e5e529b..17fa2d3 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -764,6 +764,8 @@
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 1daf458..7f2d9c3 100644
--- a/src/tint/ir/builder.cc
+++ b/src/tint/ir/builder.cc
@@ -205,6 +205,10 @@
return ir.values.Create<ir::Return>(func, args);
}
+ir::NextIteration* Builder::NextIteration(Loop* loop) {
+ return ir.values.Create<ir::NextIteration>(loop);
+}
+
ir::BreakIf* Builder::BreakIf(Value* condition, Loop* loop) {
return ir.values.Create<ir::BreakIf>(condition, loop);
}
diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h
index 8a4c021..b544cb2 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 @@
/// @returns the instruction
ir::Return* Return(Function* func, utils::VectorRef<Value*> 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 d39438a..75cc2aa 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 @@
[&](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 67c9f3f..4773f3f 100644
--- a/src/tint/ir/from_program.cc
+++ b/src/tint/ir/from_program.cc
@@ -575,7 +575,7 @@
}
// 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 @@
// 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 @@
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 f2b9c19..f9be668 100644
--- a/src/tint/ir/from_program_test.cc
+++ b/src/tint/ir/from_program_test.cc
@@ -381,7 +381,7 @@
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 @@
# Continuing block
%b3 = block {
- break_if false %b2
+ next_iteration %b2
}
# Merge block
@@ -513,7 +513,7 @@
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 @@
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 @@
# Continuing block
%b3 = block {
- break_if false %b2
+ next_iteration %b2
}
- # Merge block
- %b4 = block {
- ret
- }
-
}
}
)");
@@ -750,7 +745,7 @@
# Merge block
%b18 = block {
- break_if false %b5
+ next_iteration %b5
}
}
@@ -781,7 +776,7 @@
# Continuing block
%b3 = block {
- break_if false %b2
+ next_iteration %b2
}
# Merge block
@@ -812,7 +807,7 @@
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 @@
# Continuing block
%b3 = block {
- break_if false %b2
+ next_iteration %b2
}
# Merge block
@@ -873,7 +868,7 @@
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 @@
# 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 0000000..0c021eb
--- /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 0000000..f1211e3
--- /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<NextIteration, Branch> {
+ 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_
