[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>

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",

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

src/tint/ir/builder.cc

@@ -205,6 +205,10 @@ ir::Return* Builder::Return(Function* func, utils::VectorRef<Value*> args) {
     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);
 }

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<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

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());

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));
             }
         }
 

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 {
+          %b5 = block {
-            exit_if %b7
+            exit_if %b6
           }
 
         # Merge block
-        %b7 = block {
+        %b6 = block {
           continue %b3
         }
 
@@ -543,12 +543,7 @@ 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

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

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_