Import Tint changes from Dawn
Changes:
- f75e5c07b0678bc2d9caaa3601867a6acbd41634 [tint][utils] Abstract absl dependency by Ben Clayton <bclayton@google.com>
- ca0b9ef49e9886a587851d0f459268b4f2e0ddb2 [tint][utils] Add more string helpers by Ben Clayton <bclayton@google.com>
- 0b82a3ba15e240fd5bfbf028910424d3c7a1229a [ir][spirv-writer] Implement loop instructions by James Price <jrprice@google.com>
- bcf4174c06eca0a3b19e10d087f276fd0dfc4641 [ir][spirv-writer] Emit builtin function calls by James Price <jrprice@google.com>
- 2ee63ffc0cfd384c5484405ab9fac7307469038d [ir] Emit initializer before var declaration by James Price <jrprice@google.com>
- 02025311594d7bbcca0fb0d0abd1c44c542c9379 [ir] Add a NextIteration instruction. by dan sinclair <dsinclair@chromium.org>
- bdbbffbdfb084cb832786ec2b05c5473c94468c5 [ir] Add the Exit instructions. by dan sinclair <dsinclair@chromium.org>
- e982520e7004c9a60f099b88612d62a359478ed7 [ir] Add BreakIf instruction. by dan sinclair <dsinclair@chromium.org>
GitOrigin-RevId: f75e5c07b0678bc2d9caaa3601867a6acbd41634
Change-Id: I968e96321df30a28b39adc70b3069b0301dd9862
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/134760
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index adf0e48..3305b8d 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -251,6 +251,8 @@
"utils/hashset.h",
"utils/map.h",
"utils/math.h",
+ "utils/parse_num.cc",
+ "utils/parse_num.h",
"utils/predicates.h",
"utils/scoped_assignment.h",
"utils/slice.h",
@@ -273,6 +275,8 @@
} else {
sources += [ "diagnostic/printer_other.cc" ]
}
+
+ deps = [ ":abseil" ]
}
libtint_source_set("libtint_clone_context_hdrs") {
@@ -1064,7 +1068,6 @@
]
deps = [
- ":abseil",
":libtint_ast_src",
":libtint_builtins_src",
":libtint_program_src",
@@ -1214,6 +1217,8 @@
"ir/block_param.h",
"ir/branch.cc",
"ir/branch.h",
+ "ir/break_if.cc",
+ "ir/break_if.h",
"ir/builder.cc",
"ir/builder.h",
"ir/builtin.cc",
@@ -1228,12 +1233,16 @@
"ir/continue.h",
"ir/convert.cc",
"ir/convert.h",
- "ir/debug.cc",
- "ir/debug.h",
"ir/disassembler.cc",
"ir/disassembler.h",
"ir/discard.cc",
"ir/discard.h",
+ "ir/exit_if.cc",
+ "ir/exit_if.h",
+ "ir/exit_loop.cc",
+ "ir/exit_loop.h",
+ "ir/exit_switch.cc",
+ "ir/exit_switch.h",
"ir/function.cc",
"ir/function.h",
"ir/function_param.cc",
@@ -1248,10 +1257,10 @@
"ir/loop.h",
"ir/module.cc",
"ir/module.h",
+ "ir/next_iteration.cc",
+ "ir/next_iteration.h",
"ir/return.cc",
"ir/return.h",
- "ir/root_terminator.cc",
- "ir/root_terminator.h",
"ir/store.cc",
"ir/store.h",
"ir/switch.cc",
@@ -1958,9 +1967,11 @@
if (tint_build_ir) {
sources += [
"writer/spirv/ir/generator_impl_ir_binary_test.cc",
+ "writer/spirv/ir/generator_impl_ir_builtin_test.cc",
"writer/spirv/ir/generator_impl_ir_constant_test.cc",
"writer/spirv/ir/generator_impl_ir_function_test.cc",
"writer/spirv/ir/generator_impl_ir_if_test.cc",
+ "writer/spirv/ir/generator_impl_ir_loop_test.cc",
"writer/spirv/ir/generator_impl_ir_test.cc",
"writer/spirv/ir/generator_impl_ir_type_test.cc",
"writer/spirv/ir/generator_impl_ir_var_test.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 39bc28b..d8fe22a 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -536,6 +536,8 @@
utils/hashset.h
utils/map.h
utils/math.h
+ utils/parse_num.cc
+ utils/parse_num.h
utils/predicates.h
utils/scoped_assignment.h
utils/slice.h
@@ -722,6 +724,8 @@
ir/block_param.h
ir/branch.cc
ir/branch.h
+ ir/break_if.cc
+ ir/break_if.h
ir/builder.cc
ir/builder.h
ir/builtin.cc
@@ -736,12 +740,16 @@
ir/continue.h
ir/convert.cc
ir/convert.h
- ir/debug.cc
- ir/debug.h
ir/disassembler.cc
ir/disassembler.h
ir/discard.cc
ir/discard.h
+ ir/exit_if.cc
+ ir/exit_if.h
+ ir/exit_loop.cc
+ ir/exit_loop.h
+ ir/exit_switch.cc
+ ir/exit_switch.h
ir/from_program.cc
ir/from_program.h
ir/function.cc
@@ -758,10 +766,10 @@
ir/loop.h
ir/module.cc
ir/module.h
+ ir/next_iteration.cc
+ ir/next_iteration.h
ir/return.cc
ir/return.h
- ir/root_terminator.cc
- ir/root_terminator.h
ir/store.cc
ir/store.h
ir/switch.cc
@@ -1248,9 +1256,11 @@
if(${TINT_BUILD_IR})
list(APPEND TINT_TEST_SRCS
writer/spirv/ir/generator_impl_ir_binary_test.cc
+ writer/spirv/ir/generator_impl_ir_builtin_test.cc
writer/spirv/ir/generator_impl_ir_constant_test.cc
writer/spirv/ir/generator_impl_ir_function_test.cc
writer/spirv/ir/generator_impl_ir_if_test.cc
+ writer/spirv/ir/generator_impl_ir_loop_test.cc
writer/spirv/ir/generator_impl_ir_test.cc
writer/spirv/ir/generator_impl_ir_type_test.cc
writer/spirv/ir/generator_impl_ir_var_test.cc
diff --git a/src/tint/ast/transform/renamer_test.cc b/src/tint/ast/transform/renamer_test.cc
index 4e97c6f..c93340d 100644
--- a/src/tint/ast/transform/renamer_test.cc
+++ b/src/tint/ast/transform/renamer_test.cc
@@ -22,6 +22,7 @@
#include "src/tint/ast/transform/test_helper.h"
#include "src/tint/builtin/builtin.h"
#include "src/tint/builtin/texel_format.h"
+#include "src/tint/utils/string.h"
namespace tint::ast::transform {
namespace {
diff --git a/src/tint/bench/benchmark.cc b/src/tint/bench/benchmark.cc
index 4804a71..115d0b2 100644
--- a/src/tint/bench/benchmark.cc
+++ b/src/tint/bench/benchmark.cc
@@ -19,6 +19,7 @@
#include <utility>
#include <vector>
+#include "src/tint/utils/string.h"
#include "src/tint/utils/string_stream.h"
namespace tint::bench {
diff --git a/src/tint/cmd/main.cc b/src/tint/cmd/main.cc
index bfdc41f..5d82479 100644
--- a/src/tint/cmd/main.cc
+++ b/src/tint/cmd/main.cc
@@ -49,7 +49,6 @@
#include "tint/tint.h"
#if TINT_BUILD_IR
-#include "src/tint/ir/debug.h" // nogncheck
#include "src/tint/ir/disassembler.h" // nogncheck
#include "src/tint/ir/from_program.h" // nogncheck
#include "src/tint/ir/module.h" // nogncheck
@@ -110,7 +109,6 @@
#if TINT_BUILD_IR
bool dump_ir = false;
- bool dump_ir_graph = false;
bool use_ir = false;
#endif // TINT_BUILD_IR
@@ -374,8 +372,6 @@
#if TINT_BUILD_IR
} else if (arg == "--dump-ir") {
opts->dump_ir = true;
- } else if (arg == "--dump-ir-graph") {
- opts->dump_ir_graph = true;
} else if (arg == "--use-ir") {
opts->use_ir = true;
#endif // TINT_BUILD_IR
@@ -1072,7 +1068,7 @@
#endif // TINT_BUILD_SYNTAX_TREE_WRITER
#if TINT_BUILD_IR
- if (options.dump_ir || options.dump_ir_graph) {
+ if (options.dump_ir) {
auto result = tint::ir::FromProgram(program.get());
if (!result) {
std::cerr << "Failed to build IR from program: " << result.Failure() << std::endl;
@@ -1082,10 +1078,6 @@
tint::ir::Disassembler d(mod);
std::cout << d.Disassemble() << std::endl;
}
- if (options.dump_ir_graph) {
- auto graph = tint::ir::Debug::AsDotGraph(&mod);
- WriteFile("tint.dot", "w", graph);
- }
}
}
#endif // TINT_BUILD_IR
diff --git a/src/tint/ir/block.h b/src/tint/ir/block.h
index 597ff2c..d40002d 100644
--- a/src/tint/ir/block.h
+++ b/src/tint/ir/block.h
@@ -46,18 +46,6 @@
return instructions_.Back()->As<ir::Branch>();
}
- /// @param target the block to see if we trampoline too
- /// @returns if this block just branches to the provided target.
- bool IsTrampoline(const Block* target) const {
- if (instructions_.Length() != 1) {
- return false;
- }
- if (auto* inst = instructions_.Front()->As<ir::Branch>()) {
- return inst->To() == target;
- }
- return false;
- }
-
/// Sets the instructions in the block
/// @param instructions the instructions to set
void SetInstructions(utils::VectorRef<const Instruction*> instructions) {
diff --git a/src/tint/ir/branch.cc b/src/tint/ir/branch.cc
index 0918962..191831f 100644
--- a/src/tint/ir/branch.cc
+++ b/src/tint/ir/branch.cc
@@ -28,13 +28,6 @@
}
}
-Branch::Branch(Block* to, utils::VectorRef<Value*> args) : Branch(args) {
- to_ = to;
-
- TINT_ASSERT(IR, to_);
- to_->AddInboundBranch(this);
-}
-
Branch::~Branch() = default;
} // namespace tint::ir
diff --git a/src/tint/ir/branch.h b/src/tint/ir/branch.h
index aa3d1a0..5c926e9 100644
--- a/src/tint/ir/branch.h
+++ b/src/tint/ir/branch.h
@@ -29,15 +29,8 @@
/// A branch instruction.
class Branch : public utils::Castable<Branch, Instruction> {
public:
- /// Constructor
- /// @param to the block to branch too
- /// @param args the branch arguments
- explicit Branch(Block* to, utils::VectorRef<Value*> args = {});
~Branch() override;
- /// @returns the block being branched too.
- const Block* To() const { return to_; }
-
/// @returns the branch arguments
utils::VectorRef<Value*> Args() const { return args_; }
@@ -47,7 +40,6 @@
explicit Branch(utils::VectorRef<Value*> args);
private:
- Block* to_ = nullptr;
utils::Vector<Value*, 2> args_;
};
diff --git a/src/tint/ir/break_if.cc b/src/tint/ir/break_if.cc
new file mode 100644
index 0000000..f19fb79
--- /dev/null
+++ b/src/tint/ir/break_if.cc
@@ -0,0 +1,35 @@
+// 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/break_if.h"
+
+#include "src/tint/ir/loop.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::BreakIf);
+
+namespace tint::ir {
+
+BreakIf::BreakIf(Value* condition, ir::Loop* loop)
+ : Base(utils::Empty), condition_(condition), loop_(loop) {
+ TINT_ASSERT(IR, condition_);
+ TINT_ASSERT(IR, loop_);
+ condition_->AddUsage(this);
+ loop_->AddUsage(this);
+ loop_->Start()->AddInboundBranch(this);
+ loop_->Merge()->AddInboundBranch(this);
+}
+
+BreakIf::~BreakIf() = default;
+
+} // namespace tint::ir
diff --git a/src/tint/ir/break_if.h b/src/tint/ir/break_if.h
new file mode 100644
index 0000000..47fd4e8
--- /dev/null
+++ b/src/tint/ir/break_if.h
@@ -0,0 +1,51 @@
+// 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_BREAK_IF_H_
+#define SRC_TINT_IR_BREAK_IF_H_
+
+#include "src/tint/ir/branch.h"
+#include "src/tint/ir/value.h"
+#include "src/tint/utils/castable.h"
+
+// Forward declarations
+namespace tint::ir {
+class Loop;
+} // namespace tint::ir
+
+namespace tint::ir {
+
+/// A break-if iteration instruction.
+class BreakIf : public utils::Castable<BreakIf, Branch> {
+ public:
+ /// Constructor
+ /// @param condition the break condition
+ /// @param loop the loop containing the break-if
+ BreakIf(Value* condition, ir::Loop* loop);
+ ~BreakIf() override;
+
+ /// @returns the break condition
+ const Value* Condition() const { return condition_; }
+
+ /// @returns the loop containing the break-if
+ const ir::Loop* Loop() const { return loop_; }
+
+ private:
+ Value* condition_ = nullptr;
+ ir::Loop* loop_ = nullptr;
+};
+
+} // namespace tint::ir
+
+#endif // SRC_TINT_IR_BREAK_IF_H_
diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc
index 7c94648..7f2d9c3 100644
--- a/src/tint/ir/builder.cc
+++ b/src/tint/ir/builder.cc
@@ -37,10 +37,6 @@
return ir.blocks.Create<Block>();
}
-RootTerminator* Builder::CreateRootTerminator() {
- return ir.blocks.Create<RootTerminator>();
-}
-
Function* Builder::CreateFunction(std::string_view name,
const type::Type* return_type,
Function::PipelineStage stage,
@@ -205,17 +201,32 @@
return ir.values.Create<ir::Var>(type);
}
-ir::Branch* Builder::Branch(Block* to, utils::VectorRef<Value*> args) {
- return ir.values.Create<ir::Branch>(to, args);
-}
-
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);
+}
+
ir::Continue* Builder::Continue(Loop* loop) {
return ir.values.Create<ir::Continue>(loop);
}
+ir::ExitSwitch* Builder::ExitSwitch(Switch* sw) {
+ return ir.values.Create<ir::ExitSwitch>(sw);
+}
+
+ir::ExitLoop* Builder::ExitLoop(Loop* loop) {
+ return ir.values.Create<ir::ExitLoop>(loop);
+}
+
+ir::ExitIf* Builder::ExitIf(If* i, utils::VectorRef<Value*> args) {
+ return ir.values.Create<ir::ExitIf>(i, args);
+}
ir::BlockParam* Builder::BlockParam(const type::Type* type) {
return ir.values.Create<ir::BlockParam>(type);
diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h
index a63515e..b544cb2 100644
--- a/src/tint/ir/builder.h
+++ b/src/tint/ir/builder.h
@@ -21,20 +21,24 @@
#include "src/tint/ir/binary.h"
#include "src/tint/ir/bitcast.h"
#include "src/tint/ir/block_param.h"
+#include "src/tint/ir/break_if.h"
#include "src/tint/ir/builtin.h"
#include "src/tint/ir/constant.h"
#include "src/tint/ir/construct.h"
#include "src/tint/ir/continue.h"
#include "src/tint/ir/convert.h"
#include "src/tint/ir/discard.h"
+#include "src/tint/ir/exit_if.h"
+#include "src/tint/ir/exit_loop.h"
+#include "src/tint/ir/exit_switch.h"
#include "src/tint/ir/function.h"
#include "src/tint/ir/function_param.h"
#include "src/tint/ir/if.h"
#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/root_terminator.h"
#include "src/tint/ir/store.h"
#include "src/tint/ir/switch.h"
#include "src/tint/ir/unary.h"
@@ -63,9 +67,6 @@
/// @returns a new block flow node
Block* CreateBlock();
- /// @returns a new root terminator flow node
- RootTerminator* CreateRootTerminator();
-
/// Creates a function flow node
/// @param name the function name
/// @param return_type the function return type
@@ -337,16 +338,37 @@
/// @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
+ /// @returns the instruction
+ ir::BreakIf* BreakIf(Value* condition, Loop* loop);
+
/// Creates a continue instruction
/// @param loop the loop being continued
/// @returns the instruction
ir::Continue* Continue(Loop* loop);
- /// Creates a branch declaration
- /// @param to the node being branched too
+ /// Creates an exit switch instruction
+ /// @param sw the switch being exited
+ /// @returns the instruction
+ ir::ExitSwitch* ExitSwitch(Switch* sw);
+
+ /// Creates an exit loop instruction
+ /// @param loop the loop being exited
+ /// @returns the instruction
+ ir::ExitLoop* ExitLoop(Loop* loop);
+
+ /// Creates an exit if instruction
+ /// @param i the if being exited
/// @param args the branch arguments
/// @returns the instruction
- ir::Branch* Branch(Block* to, utils::VectorRef<Value*> args = {});
+ ir::ExitIf* ExitIf(If* i, utils::VectorRef<Value*> args = {});
/// Creates a new `BlockParam`
/// @param type the parameter type
diff --git a/src/tint/ir/debug.cc b/src/tint/ir/debug.cc
deleted file mode 100644
index c719698..0000000
--- a/src/tint/ir/debug.cc
+++ /dev/null
@@ -1,95 +0,0 @@
-// 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/ir/debug.h"
-
-#include <unordered_map>
-#include <unordered_set>
-
-#include "src/tint/ir/block.h"
-#include "src/tint/ir/continue.h"
-#include "src/tint/ir/if.h"
-#include "src/tint/ir/loop.h"
-#include "src/tint/ir/return.h"
-#include "src/tint/ir/switch.h"
-#include "src/tint/switch.h"
-#include "src/tint/utils/string_stream.h"
-
-namespace tint::ir {
-
-// static
-std::string Debug::AsDotGraph(const Module* mod) {
- size_t block_count = 0;
-
- std::unordered_set<const Block*> visited;
- std::unordered_set<const Block*> merge_blocks;
- std::unordered_map<const Block*, std::string> block_to_name;
- utils::StringStream out;
-
- auto name_for = [&](const Block* blk) -> std::string {
- if (block_to_name.count(blk) > 0) {
- return block_to_name[blk];
- }
-
- std::string name = "blk_" + std::to_string(block_count);
- block_count += 1;
-
- block_to_name[blk] = name;
- return name;
- };
-
- std::function<void(const Block*)> Graph = [&](const Block* blk) {
- if (visited.count(blk) > 0) {
- return;
- }
- visited.insert(blk);
-
- tint::Switch(blk, //
- [&](const ir::Block* b) {
- if (block_to_name.count(b) == 0) {
- out << name_for(b) << R"( [label="block"])" << std::endl;
- }
- out << name_for(b) << " -> " << name_for(b->Branch()->To());
-
- // Dashed lines to merge blocks
- if (merge_blocks.count(b->Branch()->To()) != 0) {
- out << " [style=dashed]";
- }
-
- out << std::endl;
-
- if (b->Branch()->Is<ir::Return>()) {
- return;
- } else if (auto* cont = b->Branch()->As<ir::Continue>()) {
- Graph(cont->Loop()->Continuing());
- } else {
- Graph(b->Branch()->To());
- }
- });
- };
-
- out << "digraph G {" << std::endl;
- for (const auto* func : mod->functions) {
- // Cluster each function to label and draw a box around it.
- out << "subgraph cluster_" << mod->NameOf(func).Name() << " {" << std::endl;
- out << R"(label=")" << mod->NameOf(func).Name() << R"(")" << std::endl;
- out << name_for(func->StartTarget()) << R"( [label="start"])" << std::endl;
- Graph(func->StartTarget());
- out << "}" << std::endl;
- }
- out << "}";
- return out.str();
-}
-
-} // namespace tint::ir
diff --git a/src/tint/ir/debug.h b/src/tint/ir/debug.h
deleted file mode 100644
index 2363776..0000000
--- a/src/tint/ir/debug.h
+++ /dev/null
@@ -1,40 +0,0 @@
-// 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_IR_DEBUG_H_
-#define SRC_TINT_IR_DEBUG_H_
-
-#include <string>
-
-#include "src/tint/ir/module.h"
-
-namespace tint::ir {
-
-/// Helper class to debug IR.
-class Debug {
- public:
- /// Returns the module as a dot graph
- /// @param mod the module to emit
- /// @returns the dot graph for the given module
- static std::string AsDotGraph(const Module* mod);
-
- /// Returns the module as a string
- /// @param mod the module to emit
- /// @returns the string representation of the module
- static std::string AsString(const Module* mod);
-};
-
-} // namespace tint::ir
-
-#endif // SRC_TINT_IR_DEBUG_H_
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index ae89902..75cc2aa 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -21,16 +21,20 @@
#include "src/tint/ir/binary.h"
#include "src/tint/ir/bitcast.h"
#include "src/tint/ir/block.h"
+#include "src/tint/ir/break_if.h"
#include "src/tint/ir/builtin.h"
#include "src/tint/ir/construct.h"
#include "src/tint/ir/continue.h"
#include "src/tint/ir/convert.h"
#include "src/tint/ir/discard.h"
+#include "src/tint/ir/exit_if.h"
+#include "src/tint/ir/exit_loop.h"
+#include "src/tint/ir/exit_switch.h"
#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/root_terminator.h"
#include "src/tint/ir/store.h"
#include "src/tint/ir/switch.h"
#include "src/tint/ir/user_call.h"
@@ -69,7 +73,6 @@
for (const auto* inst : b->Instructions()) {
Indent();
EmitInstruction(inst);
- out_ << std::endl;
}
}
@@ -91,8 +94,8 @@
std::string Disassembler::Disassemble() {
if (mod_.root_block) {
Indent() << "# Root block" << std::endl;
- Walk(mod_.root_block);
- Walk(mod_.root_block->Branch()->To());
+ WalkInternal(mod_.root_block);
+ out_ << std::endl;
}
for (auto* func : mod_.functions) {
@@ -107,41 +110,33 @@
}
visited_.Add(blk);
- tint::Switch(
- blk,
- [&](const ir::RootTerminator* t) {
- TINT_ASSERT(IR, !in_function_);
- Indent() << "%b" << IdOf(t) << " = root_terminator" << std::endl << std::endl;
- },
- [&](const ir::Block* b) {
- // If this block is dead, nothing to do
- if (!b->HasBranchTarget()) {
- return;
- }
+ // If this block is dead, nothing to do
+ if (!blk->HasBranchTarget()) {
+ return;
+ }
- Indent() << "%b" << IdOf(b) << " = block";
- if (!b->Params().IsEmpty()) {
- out_ << " (";
- for (auto* p : b->Params()) {
- if (p != b->Params().Front()) {
- out_ << ", ";
- }
- EmitValue(p);
- }
- out_ << ")";
- }
+ WalkInternal(blk);
+}
- out_ << " {" << std::endl;
- {
- ScopedIndent si(indent_size_);
- EmitBlockInstructions(b);
+void Disassembler::WalkInternal(const Block* blk) {
+ Indent() << "%b" << IdOf(blk) << " = block";
+ if (!blk->Params().IsEmpty()) {
+ out_ << " (";
+ for (auto* p : blk->Params()) {
+ if (p != blk->Params().Front()) {
+ out_ << ", ";
}
- Indent() << "}" << std::endl;
+ EmitValue(p);
+ }
+ out_ << ")";
+ }
- if (!b->Branch()->Is<ir::Return>()) {
- out_ << std::endl;
- }
- });
+ out_ << " {" << std::endl;
+ {
+ ScopedIndent si(indent_size_);
+ EmitBlockInstructions(blk);
+ }
+ Indent() << "}" << std::endl;
}
void Disassembler::EmitFunction(const Function* func) {
@@ -257,33 +252,39 @@
EmitValueWithType(b);
out_ << " = bitcast ";
EmitArgs(b);
+ out_ << std::endl;
},
- [&](const ir::Discard*) { out_ << "discard"; },
+ [&](const ir::Discard*) { out_ << "discard" << std::endl; },
[&](const ir::Builtin* b) {
EmitValueWithType(b);
out_ << " = " << builtin::str(b->Func()) << " ";
EmitArgs(b);
+ out_ << std::endl;
},
[&](const ir::Construct* c) {
EmitValueWithType(c);
out_ << " = construct ";
EmitArgs(c);
+ out_ << std::endl;
},
[&](const ir::Convert* c) {
EmitValueWithType(c);
out_ << " = convert " << c->FromType()->FriendlyName() << ", ";
EmitArgs(c);
+ out_ << std::endl;
},
[&](const ir::Load* l) {
EmitValueWithType(l);
out_ << " = load ";
EmitValue(l->From());
+ out_ << std::endl;
},
[&](const ir::Store* s) {
out_ << "store ";
EmitValue(s->To());
out_ << ", ";
EmitValue(s->From());
+ out_ << std::endl;
},
[&](const ir::UserCall* uc) {
EmitValueWithType(uc);
@@ -292,6 +293,7 @@
out_ << ", ";
}
EmitArgs(uc);
+ out_ << std::endl;
},
[&](const ir::Var* v) {
EmitValueWithType(v);
@@ -300,6 +302,7 @@
out_ << ", ";
EmitValue(v->Initializer());
}
+ out_ << std::endl;
},
[&](const ir::Branch* b) { EmitBranch(b); },
[&](Default) { out_ << "Unknown instruction: " << inst->TypeInfo().name; });
@@ -331,15 +334,18 @@
ScopedIndent si(indent_size_);
Indent() << "# True block" << std::endl;
Walk(i->True());
+ out_ << std::endl;
}
if (has_false) {
ScopedIndent si(indent_size_);
Indent() << "# False block" << std::endl;
Walk(i->False());
+ out_ << std::endl;
}
if (i->Merge()->HasBranchTarget()) {
Indent() << "# Merge block" << std::endl;
Walk(i->Merge());
+ out_ << std::endl;
}
}
@@ -357,16 +363,19 @@
{
ScopedIndent si(indent_size_);
Walk(l->Start());
+ out_ << std::endl;
}
if (l->Continuing()->HasBranchTarget()) {
ScopedIndent si(indent_size_);
Indent() << "# Continuing block" << std::endl;
Walk(l->Continuing());
+ out_ << std::endl;
}
if (l->Merge()->HasBranchTarget()) {
Indent() << "# Merge block" << std::endl;
Walk(l->Merge());
+ out_ << std::endl;
}
}
@@ -401,25 +410,34 @@
ScopedIndent si(indent_size_);
Indent() << "# Case block" << std::endl;
Walk(c.Start());
+ out_ << std::endl;
}
if (s->Merge()->HasBranchTarget()) {
Indent() << "# Merge block" << std::endl;
Walk(s->Merge());
+ out_ << std::endl;
}
}
void Disassembler::EmitBranch(const Branch* b) {
- std::string suffix = "";
- if (b->Is<ir::Return>()) {
- out_ << "ret";
- } else if (auto* cont = b->As<ir::Continue>()) {
- out_ << "continue %b" << IdOf(cont->Loop()->Continuing());
- } else {
- out_ << "br %b" << IdOf(b->To());
- if (b->To()->Is<RootTerminator>()) {
- suffix = "root_end";
- }
- }
+ tint::Switch(
+ b, //
+ [&](const ir::Return*) { out_ << "ret"; },
+ [&](const ir::Continue* cont) {
+ out_ << "continue %b" << IdOf(cont->Loop()->Continuing());
+ },
+ [&](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());
+ out_ << " %b" << IdOf(bi->Loop()->Start());
+ },
+ [&](Default) { out_ << "Unknown branch " << b->TypeInfo().name; });
if (!b->Args().IsEmpty()) {
out_ << " ";
@@ -430,9 +448,7 @@
EmitValue(v);
}
}
- if (!suffix.empty()) {
- out_ << " # " << suffix;
- }
+ out_ << std::endl;
}
void Disassembler::EmitArgs(const Call* call) {
@@ -503,6 +519,7 @@
EmitValue(b->LHS());
out_ << ", ";
EmitValue(b->RHS());
+ out_ << std::endl;
}
void Disassembler::EmitUnary(const Unary* u) {
@@ -518,6 +535,7 @@
}
out_ << " ";
EmitValue(u->Val());
+ out_ << std::endl;
}
} // namespace tint::ir
diff --git a/src/tint/ir/disassembler.h b/src/tint/ir/disassembler.h
index f171031..06f5b08 100644
--- a/src/tint/ir/disassembler.h
+++ b/src/tint/ir/disassembler.h
@@ -57,6 +57,7 @@
std::string_view IdOf(const Value* node);
void Walk(const Block* blk);
+ void WalkInternal(const Block* blk);
void EmitFunction(const Function* func);
void EmitInstruction(const Instruction* inst);
void EmitValueWithType(const Value* val);
diff --git a/src/tint/ir/root_terminator.cc b/src/tint/ir/exit_if.cc
similarity index 63%
rename from src/tint/ir/root_terminator.cc
rename to src/tint/ir/exit_if.cc
index bfccf46..8b7de7f 100644
--- a/src/tint/ir/root_terminator.cc
+++ b/src/tint/ir/exit_if.cc
@@ -1,4 +1,4 @@
-// Copyright 2022 The Tint Authors.
+// 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.
@@ -12,14 +12,20 @@
// See the License for the specific language governing permissions and
// limitations under the License.
-#include "src/tint/ir/root_terminator.h"
+#include "src/tint/ir/exit_if.h"
-TINT_INSTANTIATE_TYPEINFO(tint::ir::RootTerminator);
+#include "src/tint/ir/if.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::ExitIf);
namespace tint::ir {
-RootTerminator::RootTerminator() : Base() {}
+ExitIf::ExitIf(ir::If* i, utils::VectorRef<Value*> args) : Base(args), if_(i) {
+ TINT_ASSERT(IR, if_);
+ if_->AddUsage(this);
+ if_->Merge()->AddInboundBranch(this);
+}
-RootTerminator::~RootTerminator() = default;
+ExitIf::~ExitIf() = default;
} // namespace tint::ir
diff --git a/src/tint/ir/exit_if.h b/src/tint/ir/exit_if.h
new file mode 100644
index 0000000..9ba1421
--- /dev/null
+++ b/src/tint/ir/exit_if.h
@@ -0,0 +1,46 @@
+// 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_EXIT_IF_H_
+#define SRC_TINT_IR_EXIT_IF_H_
+
+#include "src/tint/ir/branch.h"
+#include "src/tint/utils/castable.h"
+
+// Forward declarations
+namespace tint::ir {
+class If;
+} // namespace tint::ir
+
+namespace tint::ir {
+
+/// A exit if instruction.
+class ExitIf : public utils::Castable<ExitIf, Branch> {
+ public:
+ /// Constructor
+ /// @param i the if being exited
+ /// @param args the branch arguments
+ explicit ExitIf(ir::If* i, utils::VectorRef<Value*> args = {});
+ ~ExitIf() override;
+
+ /// @returns the if being exited
+ const ir::If* If() const { return if_; }
+
+ private:
+ ir::If* if_ = nullptr;
+};
+
+} // namespace tint::ir
+
+#endif // SRC_TINT_IR_EXIT_IF_H_
diff --git a/src/tint/ir/root_terminator.cc b/src/tint/ir/exit_loop.cc
similarity index 63%
copy from src/tint/ir/root_terminator.cc
copy to src/tint/ir/exit_loop.cc
index bfccf46..5fe3910 100644
--- a/src/tint/ir/root_terminator.cc
+++ b/src/tint/ir/exit_loop.cc
@@ -1,4 +1,4 @@
-// Copyright 2022 The Tint Authors.
+// 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.
@@ -12,14 +12,20 @@
// See the License for the specific language governing permissions and
// limitations under the License.
-#include "src/tint/ir/root_terminator.h"
+#include "src/tint/ir/exit_loop.h"
-TINT_INSTANTIATE_TYPEINFO(tint::ir::RootTerminator);
+#include "src/tint/ir/loop.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::ExitLoop);
namespace tint::ir {
-RootTerminator::RootTerminator() : Base() {}
+ExitLoop::ExitLoop(ir::Loop* loop) : Base(utils::Empty), loop_(loop) {
+ TINT_ASSERT(IR, loop_);
+ loop_->AddUsage(this);
+ loop_->Merge()->AddInboundBranch(this);
+}
-RootTerminator::~RootTerminator() = default;
+ExitLoop::~ExitLoop() = default;
} // namespace tint::ir
diff --git a/src/tint/ir/exit_loop.h b/src/tint/ir/exit_loop.h
new file mode 100644
index 0000000..1df1119
--- /dev/null
+++ b/src/tint/ir/exit_loop.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_EXIT_LOOP_H_
+#define SRC_TINT_IR_EXIT_LOOP_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 exit loop instruction.
+class ExitLoop : public utils::Castable<ExitLoop, Branch> {
+ public:
+ /// Constructor
+ /// @param loop the loop being exited
+ explicit ExitLoop(ir::Loop* loop);
+ ~ExitLoop() override;
+
+ /// @returns the loop being exited
+ const ir::Loop* Loop() const { return loop_; }
+
+ private:
+ ir::Loop* loop_ = nullptr;
+};
+
+} // namespace tint::ir
+
+#endif // SRC_TINT_IR_EXIT_LOOP_H_
diff --git a/src/tint/ir/root_terminator.cc b/src/tint/ir/exit_switch.cc
similarity index 61%
copy from src/tint/ir/root_terminator.cc
copy to src/tint/ir/exit_switch.cc
index bfccf46..ba6a178 100644
--- a/src/tint/ir/root_terminator.cc
+++ b/src/tint/ir/exit_switch.cc
@@ -1,4 +1,4 @@
-// Copyright 2022 The Tint Authors.
+// 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.
@@ -12,14 +12,20 @@
// See the License for the specific language governing permissions and
// limitations under the License.
-#include "src/tint/ir/root_terminator.h"
+#include "src/tint/ir/exit_switch.h"
-TINT_INSTANTIATE_TYPEINFO(tint::ir::RootTerminator);
+#include "src/tint/ir/switch.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::ExitSwitch);
namespace tint::ir {
-RootTerminator::RootTerminator() : Base() {}
+ExitSwitch::ExitSwitch(ir::Switch* sw) : Base(utils::Empty), switch_(sw) {
+ TINT_ASSERT(IR, switch_);
+ switch_->AddUsage(this);
+ switch_->Merge()->AddInboundBranch(this);
+}
-RootTerminator::~RootTerminator() = default;
+ExitSwitch::~ExitSwitch() = default;
} // namespace tint::ir
diff --git a/src/tint/ir/exit_switch.h b/src/tint/ir/exit_switch.h
new file mode 100644
index 0000000..6b406fe
--- /dev/null
+++ b/src/tint/ir/exit_switch.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_EXIT_SWITCH_H_
+#define SRC_TINT_IR_EXIT_SWITCH_H_
+
+#include "src/tint/ir/branch.h"
+#include "src/tint/utils/castable.h"
+
+// Forward declarations
+namespace tint::ir {
+class Switch;
+} // namespace tint::ir
+
+namespace tint::ir {
+
+/// A exit switch instruction.
+class ExitSwitch : public utils::Castable<ExitSwitch, Branch> {
+ public:
+ /// Constructor
+ /// @param sw the switch being exited
+ explicit ExitSwitch(ir::Switch* sw);
+ ~ExitSwitch() override;
+
+ /// @returns the switch being exited
+ const ir::Switch* Switch() const { return switch_; }
+
+ private:
+ ir::Switch* switch_ = nullptr;
+};
+
+} // namespace tint::ir
+
+#endif // SRC_TINT_IR_EXIT_SWITCH_H_
diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc
index 444950e..20555d0 100644
--- a/src/tint/ir/from_program.cc
+++ b/src/tint/ir/from_program.cc
@@ -62,6 +62,9 @@
#include "src/tint/ast/while_statement.h"
#include "src/tint/ir/block_param.h"
#include "src/tint/ir/builder.h"
+#include "src/tint/ir/exit_if.h"
+#include "src/tint/ir/exit_loop.h"
+#include "src/tint/ir/exit_switch.h"
#include "src/tint/ir/function.h"
#include "src/tint/ir/if.h"
#include "src/tint/ir/loop.h"
@@ -98,11 +101,8 @@
using ResultType = utils::Result<Module, diag::List>;
-// For an `if` and `switch` block, the merge has a registered incoming branch instruction of the
-// `if` and `switch. So, to determine if the merge is connected to any of the branches that happend
-// in the `if` or `switch` we need a `count` value that is larger then 1.
-bool IsConnected(const Block* b, uint32_t count) {
- return b->InboundBranches().Length() > count;
+bool IsConnected(const Block* b) {
+ return b->InboundBranches().Length() > 0;
}
/// Impl is the private-implementation of FromProgram().
@@ -176,21 +176,6 @@
current_flow_block_ = nullptr;
}
- void BranchTo(Block* node, utils::VectorRef<Value*> args = {}) {
- TINT_ASSERT(IR, current_flow_block_);
- TINT_ASSERT(IR, !current_flow_block_->HasBranchTarget());
-
- current_flow_block_->Instructions().Push(builder_.Branch(node, args));
- current_flow_block_ = nullptr;
- }
-
- void BranchToIfNeeded(Block* node) {
- if (!NeedBranch()) {
- return;
- }
- BranchTo(node);
- }
-
Branch* FindEnclosingControl(ControlFlags flags) {
for (auto it = control_stack_.rbegin(); it != control_stack_.rend(); ++it) {
if ((*it)->Is<Loop>()) {
@@ -238,11 +223,6 @@
});
}
- // Add the root terminator if needed
- if (mod.root_block) {
- mod.root_block->Instructions().Push(builder_.Branch(builder_.CreateRootTerminator()));
- }
-
if (diagnostics_.contains_errors()) {
return ResultType(std::move(diagnostics_));
}
@@ -541,7 +521,9 @@
EmitBlock(stmt->body);
// If the true branch did not execute control flow, then go to the Merge().target
- BranchToIfNeeded(if_inst->Merge());
+ if (NeedBranch()) {
+ SetBranch(builder_.ExitIf(if_inst));
+ }
current_flow_block_ = if_inst->False();
if (stmt->else_statement) {
@@ -549,14 +531,16 @@
}
// If the false branch did not execute control flow, then go to the Merge().target
- BranchToIfNeeded(if_inst->Merge());
+ if (NeedBranch()) {
+ SetBranch(builder_.ExitIf(if_inst));
+ }
}
current_flow_block_ = nullptr;
// If both branches went somewhere, then they both returned, continued or broke. So,
// there is no need for the if merge-block and there is nothing to branch to the merge
// block anyway.
- if (IsConnected(if_inst->Merge(), 1)) {
+ if (IsConnected(if_inst->Merge())) {
current_flow_block_ = if_inst->Merge();
}
}
@@ -580,7 +564,7 @@
SetBranch(builder_.Continue(loop_inst));
}
- if (IsConnected(loop_inst->Continuing(), 0)) {
+ if (IsConnected(loop_inst->Continuing())) {
// Note, even if there is no continuing block, we may have branched into the
// continue so we have to set the current block and then emit the branch if needed
// below otherwise empty continuing blocks will fail to branch back to the start
@@ -590,7 +574,9 @@
EmitBlock(stmt->continuing);
}
// Branch back to the start node if the continue target didn't branch out already
- BranchToIfNeeded(loop_inst->Start());
+ if (NeedBranch()) {
+ SetBranch(builder_.NextIteration(loop_inst));
+ }
}
}
@@ -598,7 +584,7 @@
// target branches, eventually, to the merge, but nothing branched to the
// Continuing() block.
current_flow_block_ = loop_inst->Merge();
- if (!IsConnected(loop_inst->Merge(), 0)) {
+ if (!IsConnected(loop_inst->Merge())) {
current_flow_block_ = nullptr;
}
}
@@ -608,7 +594,8 @@
current_flow_block_->Instructions().Push(loop_inst);
// Continue is always empty, just go back to the start
- loop_inst->Continuing()->Instructions().Push(builder_.Branch(loop_inst->Start()));
+ current_flow_block_ = loop_inst->Continuing();
+ SetBranch(builder_.NextIteration(loop_inst));
{
ControlStackScope scope(this, loop_inst);
@@ -623,10 +610,14 @@
// Create an `if (cond) {} else {break;}` control flow
auto* if_inst = builder_.CreateIf(reg.Get());
- if_inst->True()->Instructions().Push(builder_.Branch(if_inst->Merge()));
- if_inst->False()->Instructions().Push(builder_.Branch(loop_inst->Merge()));
current_flow_block_->Instructions().Push(if_inst);
+ current_flow_block_ = if_inst->True();
+ SetBranch(builder_.ExitIf(if_inst));
+
+ current_flow_block_ = if_inst->False();
+ SetBranch(builder_.ExitLoop(loop_inst));
+
current_flow_block_ = if_inst->Merge();
EmitBlock(stmt->body);
@@ -666,10 +657,14 @@
// Create an `if (cond) {} else {break;}` control flow
auto* if_inst = builder_.CreateIf(reg.Get());
- if_inst->True()->Instructions().Push(builder_.Branch(if_inst->Merge()));
- if_inst->False()->Instructions().Push(builder_.Branch(loop_inst->Merge()));
current_flow_block_->Instructions().Push(if_inst);
+ current_flow_block_ = if_inst->True();
+ SetBranch(builder_.ExitIf(if_inst));
+
+ current_flow_block_ = if_inst->False();
+ SetBranch(builder_.ExitLoop(loop_inst));
+
current_flow_block_ = if_inst->Merge();
}
@@ -681,7 +676,7 @@
if (stmt->continuing) {
current_flow_block_ = loop_inst->Continuing();
EmitStatement(stmt->continuing);
- loop_inst->Continuing()->Instructions().Push(builder_.Branch(loop_inst->Start()));
+ SetBranch(builder_.NextIteration(loop_inst));
}
}
@@ -716,12 +711,14 @@
current_flow_block_ = builder_.CreateCase(switch_inst, selectors);
EmitBlock(c->Body()->Declaration());
- BranchToIfNeeded(switch_inst->Merge());
+ if (NeedBranch()) {
+ SetBranch(builder_.ExitSwitch(switch_inst));
+ }
}
}
current_flow_block_ = nullptr;
- if (IsConnected(switch_inst->Merge(), 1)) {
+ if (IsConnected(switch_inst->Merge())) {
current_flow_block_ = switch_inst->Merge();
}
}
@@ -743,9 +740,9 @@
TINT_ASSERT(IR, current_control);
if (auto* c = current_control->As<Loop>()) {
- BranchTo(c->Merge());
+ SetBranch(builder_.ExitLoop(c));
} else if (auto* s = current_control->As<Switch>()) {
- BranchTo(s->Merge());
+ SetBranch(builder_.ExitSwitch(s));
} else {
TINT_UNREACHABLE(IR, diagnostics_);
}
@@ -772,31 +769,14 @@
}
void EmitBreakIf(const ast::BreakIfStatement* stmt) {
+ auto* current_control = FindEnclosingControl(ControlFlags::kExcludeSwitch);
+
// Emit the break-if condition into the end of the preceding block
- auto reg = EmitExpression(stmt->condition);
- if (!reg) {
+ auto cond = EmitExpression(stmt->condition);
+ if (!cond) {
return;
}
- auto* if_inst = builder_.CreateIf(reg.Get());
- current_flow_block_->Instructions().Push(if_inst);
-
- auto* current_control = FindEnclosingControl(ControlFlags::kExcludeSwitch);
- TINT_ASSERT(IR, current_control);
- TINT_ASSERT(IR, current_control->Is<Loop>());
-
- auto* loop = current_control->As<Loop>();
-
- current_flow_block_ = if_inst->True();
- BranchTo(loop->Merge());
-
- current_flow_block_ = if_inst->False();
- BranchTo(if_inst->Merge());
-
- current_flow_block_ = if_inst->Merge();
-
- // The `break-if` has to be the last item in the continuing block. The false branch of
- // the `break-if` will always take us back to the start of the loop.
- BranchTo(loop->Start());
+ SetBranch(builder_.BreakIf(cond.Get(), current_control->As<ir::Loop>()));
}
utils::Result<Value*> EmitExpression(const ast::Expression* expr) {
@@ -862,8 +842,6 @@
ref->Access());
auto* val = builder_.Declare(ty);
- current_flow_block_->Instructions().Push(val);
-
if (v->initializer) {
auto init = EmitExpression(v->initializer);
if (!init) {
@@ -871,6 +849,8 @@
}
val->SetInitializer(init.Get());
}
+ current_flow_block_->Instructions().Push(val);
+
// Store the declaration so we can get the instruction to store too
scopes_.Set(v->name->symbol, val);
@@ -978,14 +958,14 @@
// If the lhs is false, then that is the result we want to pass to the merge
// block as our argument
current_flow_block_ = if_inst->False();
- BranchTo(if_inst->Merge(), std::move(alt_args));
+ SetBranch(builder_.ExitIf(if_inst, std::move(alt_args)));
current_flow_block_ = if_inst->True();
} else {
// If the lhs is true, then that is the result we want to pass to the merge
// block as our argument
current_flow_block_ = if_inst->True();
- BranchTo(if_inst->Merge(), std::move(alt_args));
+ SetBranch(builder_.ExitIf(if_inst, std::move(alt_args)));
current_flow_block_ = if_inst->False();
}
@@ -997,7 +977,7 @@
utils::Vector<Value*, 1> args;
args.Push(rhs.Get());
- BranchTo(if_inst->Merge(), std::move(args));
+ SetBranch(builder_.ExitIf(if_inst, std::move(args)));
}
current_flow_block_ = if_inst->Merge();
diff --git a/src/tint/ir/from_program_binary_test.cc b/src/tint/ir/from_program_binary_test.cc
index 1a81f56..8cc5604 100644
--- a/src/tint/ir/from_program_binary_test.cc
+++ b/src/tint/ir/from_program_binary_test.cc
@@ -60,13 +60,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
%3:u32 = load %v1
%4:u32 = add %3, 1u
store %v1, %4
@@ -87,13 +84,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
%3:u32 = load %v1
%4:u32 = add %3, 1u
store %v1, %4
@@ -137,13 +131,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, i32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
%3:i32 = load %v1
%4:i32 = sub %3, 1i
store %v1, %4
@@ -164,13 +155,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
%3:u32 = load %v1
%4:u32 = sub %3, 1u
store %v1, %4
@@ -214,13 +202,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
%3:u32 = load %v1
%4:u32 = mul %3, 1u
store %v1, %4
@@ -264,13 +249,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
%3:u32 = load %v1
%4:u32 = div %3, 1u
store %v1, %4
@@ -314,13 +296,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
%3:u32 = load %v1
%4:u32 = mod %3, 1u
store %v1, %4
@@ -364,13 +343,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, bool, read_write> = var
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
%3:bool = load %v1
%4:bool = and %3, false
store %v1, %4
@@ -414,13 +390,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, bool, read_write> = var
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
%3:bool = load %v1
%4:bool = or %3, false
store %v1, %4
@@ -464,13 +437,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
%3:u32 = load %v1
%4:u32 = xor %3, 1u
store %v1, %4
@@ -499,12 +469,12 @@
if %3 [t: %b3, f: %b4, m: %b5]
# True block
%b3 = block {
- br %b5 false
+ exit_if %b5 false
}
# False block
%b4 = block {
- br %b5 %3
+ exit_if %b5 %3
}
# Merge block
@@ -512,12 +482,12 @@
if %4:bool [t: %b6, f: %b7, m: %b8]
# True block
%b6 = block {
- br %b8
+ exit_if %b8
}
# False block
%b7 = block {
- br %b8
+ exit_if %b8
}
# Merge block
@@ -527,9 +497,7 @@
}
-
}
-
}
)");
}
@@ -553,12 +521,12 @@
if %3 [t: %b3, f: %b4, m: %b5]
# True block
%b3 = block {
- br %b5 %3
+ exit_if %b5 %3
}
# False block
%b4 = block {
- br %b5 true
+ exit_if %b5 true
}
# Merge block
@@ -566,12 +534,12 @@
if %4:bool [t: %b6, f: %b7, m: %b8]
# True block
%b6 = block {
- br %b8
+ exit_if %b8
}
# False block
%b7 = block {
- br %b8
+ exit_if %b8
}
# Merge block
@@ -581,9 +549,7 @@
}
-
}
-
}
)");
}
@@ -760,13 +726,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
%3:u32 = load %v1
%4:u32 = shiftl %3, 1u
store %v1, %4
@@ -810,13 +773,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
%3:u32 = load %v1
%4:u32 = shiftr %3, 1u
store %v1, %4
@@ -853,12 +813,12 @@
%7:f32 = mul 2.29999995231628417969f, %6
%8:f32 = div %5, %7
%9:bool = gt 2.5f, %8
- br %b5 %9
+ exit_if %b5 %9
}
# False block
%b4 = block {
- br %b5 %4
+ exit_if %b5 %4
}
# Merge block
@@ -867,7 +827,6 @@
}
}
-
}
)");
}
diff --git a/src/tint/ir/from_program_builtin_test.cc b/src/tint/ir/from_program_builtin_test.cc
index 3993fbb..acdbc8a 100644
--- a/src/tint/ir/from_program_builtin_test.cc
+++ b/src/tint/ir/from_program_builtin_test.cc
@@ -37,13 +37,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%i:ptr<private, f32, read_write> = var, 1.0f
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
%3:f32 = load %i
%tint_symbol:f32 = asin %3
ret
diff --git a/src/tint/ir/from_program_call_test.cc b/src/tint/ir/from_program_call_test.cc
index 43a9709..155e42c 100644
--- a/src/tint/ir/from_program_call_test.cc
+++ b/src/tint/ir/from_program_call_test.cc
@@ -102,13 +102,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%i:ptr<private, i32, read_write> = var, 1i
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
%3:i32 = load %i
%tint_symbol:f32 = convert i32, %3
ret
@@ -127,11 +124,8 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%i:ptr<private, vec3<f32>, read_write> = var, vec3<f32> 0.0f
- br %b2 # root_end
}
-%b2 = root_terminator
-
)");
}
@@ -146,13 +140,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%i:ptr<private, f32, read_write> = var, 1.0f
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
%3:f32 = load %i
%tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %3
ret
diff --git a/src/tint/ir/from_program_store_test.cc b/src/tint/ir/from_program_store_test.cc
index 2d3f3d7..38efea0 100644
--- a/src/tint/ir/from_program_store_test.cc
+++ b/src/tint/ir/from_program_store_test.cc
@@ -38,13 +38,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%a:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
store %a, 4u
ret
}
diff --git a/src/tint/ir/from_program_test.cc b/src/tint/ir/from_program_test.cc
index 0b23ec5..f9be668 100644
--- a/src/tint/ir/from_program_test.cc
+++ b/src/tint/ir/from_program_test.cc
@@ -141,7 +141,7 @@
EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
- EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -149,12 +149,12 @@
if true [t: %b2, f: %b3, m: %b4]
# True block
%b2 = block {
- br %b4
+ exit_if %b4
}
# False block
%b3 = block {
- br %b4
+ exit_if %b4
}
# Merge block
@@ -163,7 +163,6 @@
}
}
-
}
)");
}
@@ -182,7 +181,7 @@
EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
- EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -192,9 +191,10 @@
%b2 = block {
ret
}
+
# False block
%b3 = block {
- br %b4
+ exit_if %b4
}
# Merge block
@@ -203,7 +203,6 @@
}
}
-
}
)");
}
@@ -222,7 +221,7 @@
EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
- EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -230,20 +229,20 @@
if true [t: %b2, f: %b3, m: %b4]
# True block
%b2 = block {
- br %b4
+ exit_if %b4
}
# False block
%b3 = block {
ret
}
+
# Merge block
%b4 = block {
ret
}
}
-
}
)");
}
@@ -262,7 +261,7 @@
EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
- EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(0u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -272,13 +271,13 @@
%b2 = block {
ret
}
+
# False block
%b3 = block {
ret
}
}
-
}
)");
}
@@ -306,20 +305,19 @@
%b2 = block {
loop [s: %b5, m: %b6]
%b5 = block {
- br %b6
+ exit_loop %b6
}
# Merge block
%b6 = block {
- br %b4
+ exit_if %b4
}
-
}
# False block
%b3 = block {
- br %b4
+ exit_if %b4
}
# Merge block
@@ -328,7 +326,6 @@
}
}
-
}
)");
}
@@ -345,7 +342,7 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(0u, flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
@@ -354,7 +351,7 @@
%b1 = block {
loop [s: %b2, m: %b3]
%b2 = block {
- br %b3
+ exit_loop %b3
}
# Merge block
@@ -363,7 +360,6 @@
}
}
-
}
)");
}
@@ -383,12 +379,12 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
+ 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(1u, if_flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length());
- EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -398,12 +394,12 @@
if true [t: %b5, f: %b6, m: %b7]
# True block
%b5 = block {
- br %b4
+ exit_loop %b4
}
# False block
%b6 = block {
- br %b7
+ exit_if %b7
}
# Merge block
@@ -411,12 +407,11 @@
continue %b3
}
-
}
# Continuing block
%b3 = block {
- br %b2
+ next_iteration %b2
}
# Merge block
@@ -425,7 +420,6 @@
}
}
-
}
)");
}
@@ -440,16 +434,12 @@
auto m = res.Move();
auto* loop_flow = FindSingleValue<ir::Loop>(m);
- auto* break_if_flow = FindSingleValue<ir::If>(m);
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
+ 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(1u, break_if_flow->True()->InboundBranches().Length());
- EXPECT_EQ(1u, break_if_flow->False()->InboundBranches().Length());
- EXPECT_EQ(2u, break_if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -461,23 +451,7 @@
# Continuing block
%b3 = block {
- if true [t: %b5, f: %b6, m: %b7]
- # True block
- %b5 = block {
- br %b4
- }
-
- # False block
- %b6 = block {
- br %b7
- }
-
- # Merge block
- %b7 = block {
- br %b2
- }
-
-
+ break_if true %b2
}
# Merge block
@@ -486,7 +460,6 @@
}
}
-
}
)");
}
@@ -511,23 +484,7 @@
# Continuing block
%b3 = block {
- if true [t: %b5, f: %b6, m: %b7]
- # True block
- %b5 = block {
- br %b4
- }
-
- # False block
- %b6 = block {
- br %b7
- }
-
- # Merge block
- %b7 = block {
- br %b2
- }
-
-
+ break_if true %b2
}
# Merge block
@@ -536,7 +493,6 @@
}
}
-
}
)");
}
@@ -555,12 +511,12 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
EXPECT_EQ(1u, loop_flow->Continuing()->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(2u, if_flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -572,9 +528,10 @@
%b4 = block {
ret
}
+
# False block
%b5 = block {
- br %b6
+ exit_if %b6
}
# Merge block
@@ -582,17 +539,14 @@
continue %b3
}
-
}
# Continuing block
%b3 = block {
- br %b2
+ next_iteration %b2
}
-
}
-
}
)");
}
@@ -609,7 +563,7 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(0u, loop_flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length());
@@ -622,7 +576,6 @@
}
}
-
}
)");
}
@@ -648,7 +601,7 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(0u, loop_flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length());
@@ -661,7 +614,6 @@
}
}
-
}
)");
}
@@ -680,12 +632,12 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(0u, loop_flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(2u, 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());
+ EXPECT_EQ(0u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -695,15 +647,14 @@
if true [t: %b4, f: %b5]
# True block
%b4 = block {
- br %b3
+ exit_loop %b3
}
# False block
%b5 = block {
- br %b3
+ exit_loop %b3
}
-
}
# Merge block
@@ -712,7 +663,6 @@
}
}
-
}
)");
}
@@ -744,12 +694,12 @@
if true [t: %b8, f: %b9, m: %b10]
# True block
%b8 = block {
- br %b7
+ exit_loop %b7
}
# False block
%b9 = block {
- br %b10
+ exit_if %b10
}
# Merge block
@@ -762,7 +712,7 @@
# False block
%b12 = block {
- br %b13
+ exit_if %b13
}
# Merge block
@@ -770,17 +720,15 @@
continue %b6
}
-
}
-
}
# Continuing block
%b6 = block {
loop [s: %b14, m: %b15]
%b14 = block {
- br %b15
+ exit_loop %b15
}
# Merge block
@@ -792,63 +740,43 @@
# Continuing block
%b17 = block {
- if true [t: %b19, f: %b20, m: %b21]
- # True block
- %b19 = block {
- br %b18
- }
-
- # False block
- %b20 = block {
- br %b21
- }
-
- # Merge block
- %b21 = block {
- br %b16
- }
-
-
+ break_if true %b16
}
# Merge block
%b18 = block {
- br %b5
+ next_iteration %b5
}
-
}
-
}
# Merge block
%b7 = block {
- if true [t: %b22, f: %b23, m: %b24]
+ if true [t: %b19, f: %b20, m: %b21]
# True block
- %b22 = block {
- br %b4
+ %b19 = block {
+ exit_loop %b4
}
# False block
- %b23 = block {
- br %b24
+ %b20 = block {
+ exit_if %b21
}
# Merge block
- %b24 = block {
+ %b21 = block {
continue %b3
}
-
}
-
}
# Continuing block
%b3 = block {
- br %b2
+ next_iteration %b2
}
# Merge block
@@ -857,7 +785,6 @@
}
}
-
}
)");
}
@@ -878,12 +805,12 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(2u, flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Start()->InboundBranches().Length());
EXPECT_EQ(1u, flow->Continuing()->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(2u, if_flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -893,12 +820,12 @@
if false [t: %b5, f: %b6, m: %b7]
# True block
%b5 = block {
- br %b7
+ exit_if %b7
}
# False block
%b6 = block {
- br %b4
+ exit_loop %b4
}
# Merge block
@@ -906,12 +833,11 @@
continue %b3
}
-
}
# Continuing block
%b3 = block {
- br %b2
+ next_iteration %b2
}
# Merge block
@@ -920,7 +846,6 @@
}
}
-
}
)");
}
@@ -941,12 +866,12 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(2u, flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, flow->Continuing()->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(2u, if_flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -956,12 +881,12 @@
if true [t: %b5, f: %b6, m: %b7]
# True block
%b5 = block {
- br %b7
+ exit_if %b7
}
# False block
%b6 = block {
- br %b4
+ exit_loop %b4
}
# Merge block
@@ -973,7 +898,7 @@
# Continuing block
%b3 = block {
- br %b2
+ next_iteration %b2
}
# Merge block
@@ -982,7 +907,6 @@
}
}
-
}
)");
}
@@ -1021,7 +945,7 @@
EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length());
- EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), R"()");
}
@@ -1038,7 +962,7 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(0u, flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
@@ -1047,7 +971,7 @@
%b1 = block {
loop [s: %b2, m: %b3]
%b2 = block {
- br %b3
+ exit_loop %b3
}
# Merge block
@@ -1056,7 +980,6 @@
}
}
-
}
)");
}
@@ -1095,7 +1018,7 @@
EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length());
EXPECT_EQ(1u, cases[2].Start()->InboundBranches().Length());
- EXPECT_EQ(4u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1103,17 +1026,17 @@
switch 1i [c: (0i, %b2), c: (1i, %b3), c: (default, %b4), m: %b5]
# Case block
%b2 = block {
- br %b5
+ exit_switch %b5
}
# Case block
%b3 = block {
- br %b5
+ exit_switch %b5
}
# Case block
%b4 = block {
- br %b5
+ exit_switch %b5
}
# Merge block
@@ -1122,7 +1045,6 @@
}
}
-
}
)");
}
@@ -1157,7 +1079,7 @@
EXPECT_TRUE(cases[0].selectors[2].IsDefault());
EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
- EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1165,7 +1087,7 @@
switch 1i [c: (0i 1i default, %b2), m: %b3]
# Case block
%b2 = block {
- br %b3
+ exit_switch %b3
}
# Merge block
@@ -1174,7 +1096,6 @@
}
}
-
}
)");
}
@@ -1197,7 +1118,7 @@
EXPECT_TRUE(cases[0].selectors[0].IsDefault());
EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
- EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1205,7 +1126,7 @@
switch 1i [c: (default, %b2), m: %b3]
# Case block
%b2 = block {
- br %b3
+ exit_switch %b3
}
# Merge block
@@ -1214,7 +1135,6 @@
}
}
-
}
)");
}
@@ -1245,7 +1165,7 @@
EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length());
- EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
// This is 1 because the if is dead-code eliminated and the return doesn't happen.
EXPECT_EQ(Disassemble(m),
@@ -1254,12 +1174,12 @@
switch 1i [c: (0i, %b2), c: (default, %b3), m: %b4]
# Case block
%b2 = block {
- br %b4
+ exit_switch %b4
}
# Case block
%b3 = block {
- br %b4
+ exit_switch %b4
}
# Merge block
@@ -1268,7 +1188,6 @@
}
}
-
}
)");
}
@@ -1302,7 +1221,7 @@
EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length());
- EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(0u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1312,13 +1231,13 @@
%b2 = block {
ret
}
+
# Case block
%b3 = block {
ret
}
}
-
}
)");
}
diff --git a/src/tint/ir/from_program_unary_test.cc b/src/tint/ir/from_program_unary_test.cc
index 774be8c..bb58c02 100644
--- a/src/tint/ir/from_program_unary_test.cc
+++ b/src/tint/ir/from_program_unary_test.cc
@@ -107,13 +107,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v2:ptr<private, i32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
ret
}
}
@@ -134,13 +131,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v3:ptr<private, i32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_terminator
-
-%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
- %b3 = block {
+%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
+ %b2 = block {
store %v3, 42i
ret
}
diff --git a/src/tint/ir/from_program_var_test.cc b/src/tint/ir/from_program_var_test.cc
index 29377e4..c533738 100644
--- a/src/tint/ir/from_program_var_test.cc
+++ b/src/tint/ir/from_program_var_test.cc
@@ -35,11 +35,8 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%a:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_terminator
-
)");
}
@@ -53,11 +50,8 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%a:ptr<private, u32, read_write> = var, 2u
- br %b2 # root_end
}
-%b2 = root_terminator
-
)");
}
@@ -78,7 +72,7 @@
)");
}
-TEST_F(IR_BuilderImplTest, Emit_Var_Init) {
+TEST_F(IR_BuilderImplTest, Emit_Var_Init_Constant) {
auto* expr = Expr(2_u);
auto* a = Var("a", ty.u32(), builtin::AddressSpace::kFunction, expr);
WrapInFunction(a);
@@ -95,5 +89,27 @@
}
)");
}
+
+TEST_F(IR_BuilderImplTest, Emit_Var_Init_NonConstant) {
+ auto* a = Var("a", ty.u32(), builtin::AddressSpace::kFunction);
+ auto* b = Var("b", ty.u32(), builtin::AddressSpace::kFunction, Add("a", 2_u));
+ WrapInFunction(a, b);
+
+ auto m = Build();
+ ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+ EXPECT_EQ(Disassemble(m.Get()),
+ R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
+ %b1 = block {
+ %a:ptr<function, u32, read_write> = var
+ %3:u32 = load %a
+ %4:u32 = add %3, 2u
+ %b:ptr<function, u32, read_write> = var, %4
+ ret
+ }
+}
+)");
+}
+
} // namespace
} // namespace tint::ir
diff --git a/src/tint/ir/if.cc b/src/tint/ir/if.cc
index a89a51f..2069c6d 100644
--- a/src/tint/ir/if.cc
+++ b/src/tint/ir/if.cc
@@ -19,7 +19,7 @@
namespace tint::ir {
If::If(Value* cond, Block* t, Block* f, Block* m)
- : Base(m), condition_(cond), true_(t), false_(f), merge_(m) {
+ : Base(utils::Empty), condition_(cond), true_(t), false_(f), merge_(m) {
TINT_ASSERT(IR, true_);
TINT_ASSERT(IR, false_);
TINT_ASSERT(IR, merge_);
diff --git a/src/tint/ir/loop.cc b/src/tint/ir/loop.cc
index 0bbb710..fe34283 100644
--- a/src/tint/ir/loop.cc
+++ b/src/tint/ir/loop.cc
@@ -18,7 +18,8 @@
namespace tint::ir {
-Loop::Loop(Block* s, Block* c, Block* m) : Base(s), start_(s), continuing_(c), merge_(m) {
+Loop::Loop(Block* s, Block* c, Block* m)
+ : Base(utils::Empty), start_(s), continuing_(c), merge_(m) {
TINT_ASSERT(IR, start_);
TINT_ASSERT(IR, continuing_);
TINT_ASSERT(IR, merge_);
diff --git a/src/tint/ir/root_terminator.cc b/src/tint/ir/next_iteration.cc
similarity index 61%
copy from src/tint/ir/root_terminator.cc
copy to src/tint/ir/next_iteration.cc
index bfccf46..0c021eb 100644
--- a/src/tint/ir/root_terminator.cc
+++ b/src/tint/ir/next_iteration.cc
@@ -1,4 +1,4 @@
-// Copyright 2022 The Tint Authors.
+// 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.
@@ -12,14 +12,20 @@
// See the License for the specific language governing permissions and
// limitations under the License.
-#include "src/tint/ir/root_terminator.h"
+#include "src/tint/ir/next_iteration.h"
-TINT_INSTANTIATE_TYPEINFO(tint::ir::RootTerminator);
+#include "src/tint/ir/loop.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::NextIteration);
namespace tint::ir {
-RootTerminator::RootTerminator() : Base() {}
+NextIteration::NextIteration(ir::Loop* loop) : Base(utils::Empty), loop_(loop) {
+ TINT_ASSERT(IR, loop_);
+ loop_->AddUsage(this);
+ loop_->Start()->AddInboundBranch(this);
+}
-RootTerminator::~RootTerminator() = default;
+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_
diff --git a/src/tint/ir/root_terminator.h b/src/tint/ir/root_terminator.h
deleted file mode 100644
index 4a52b32..0000000
--- a/src/tint/ir/root_terminator.h
+++ /dev/null
@@ -1,32 +0,0 @@
-// 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_IR_ROOT_TERMINATOR_H_
-#define SRC_TINT_IR_ROOT_TERMINATOR_H_
-
-#include "src/tint/ir/block.h"
-
-namespace tint::ir {
-
-/// Block used as the end of a root block. There are no instructions in this block.
-class RootTerminator : public utils::Castable<RootTerminator, Block> {
- public:
- /// Constructor
- RootTerminator();
- ~RootTerminator() override;
-};
-
-} // namespace tint::ir
-
-#endif // SRC_TINT_IR_ROOT_TERMINATOR_H_
diff --git a/src/tint/ir/switch.cc b/src/tint/ir/switch.cc
index a28666a..003feb2 100644
--- a/src/tint/ir/switch.cc
+++ b/src/tint/ir/switch.cc
@@ -18,7 +18,7 @@
namespace tint::ir {
-Switch::Switch(Value* cond, Block* m) : Base(m), condition_(cond), merge_(m) {
+Switch::Switch(Value* cond, Block* m) : Base(utils::Empty), condition_(cond), merge_(m) {
TINT_ASSERT(IR, condition_);
TINT_ASSERT(IR, merge_);
condition_->AddUsage(this);
diff --git a/src/tint/ir/to_program.cc b/src/tint/ir/to_program.cc
index af4a957..08436f3 100644
--- a/src/tint/ir/to_program.cc
+++ b/src/tint/ir/to_program.cc
@@ -20,6 +20,7 @@
#include "src/tint/ir/block.h"
#include "src/tint/ir/call.h"
#include "src/tint/ir/constant.h"
+#include "src/tint/ir/exit_if.h"
#include "src/tint/ir/if.h"
#include "src/tint/ir/instruction.h"
#include "src/tint/ir/load.h"
@@ -121,46 +122,27 @@
while (block) {
TINT_ASSERT(IR, block->HasBranchTarget());
- enum Status { kContinue, kStop, kError };
-
- Status status = tint::Switch(
- block,
-
- [&](const ir::Block* blk) {
- for (auto* inst : blk->Instructions()) {
- auto stmt = Stmt(inst);
- if (TINT_UNLIKELY(!stmt)) {
- return kError;
- }
- if (auto* s = stmt.Get()) {
- stmts.Push(s);
- }
- }
- if (auto* if_ = blk->Branch()->As<ir::If>()) {
- if (if_->Merge()->HasBranchTarget()) {
- block = if_->Merge();
- return kContinue;
- }
- } else if (auto* switch_ = blk->Branch()->As<ir::Switch>()) {
- if (switch_->Merge()->HasBranchTarget()) {
- block = switch_->Merge();
- return kContinue;
- }
- }
- return kStop;
- },
-
- [&](Default) {
- UNHANDLED_CASE(block);
- return kError;
- });
-
- if (TINT_UNLIKELY(status == kError)) {
- return nullptr;
+ for (auto* inst : block->Instructions()) {
+ auto stmt = Stmt(inst);
+ if (TINT_UNLIKELY(!stmt)) {
+ return nullptr;
+ }
+ if (auto* s = stmt.Get()) {
+ stmts.Push(s);
+ }
}
- if (status == kStop) {
- break;
+ if (auto* if_ = block->Branch()->As<ir::If>()) {
+ if (if_->Merge()->HasBranchTarget()) {
+ block = if_->Merge();
+ continue;
+ }
+ } else if (auto* switch_ = block->Branch()->As<ir::Switch>()) {
+ if (switch_->Merge()->HasBranchTarget()) {
+ block = switch_->Merge();
+ continue;
+ }
}
+ break;
}
return b.Block(std::move(stmts));
@@ -174,16 +156,20 @@
return nullptr;
}
- if (!IsEmpty(i->False(), i->Merge())) {
+ auto* false_blk = i->False();
+ if (false_blk->Instructions().Length() > 1 ||
+ (false_blk->Instructions().Length() == 1 && false_blk->HasBranchTarget() &&
+ !false_blk->Branch()->Is<ir::ExitIf>())) {
// If the else target is an `if` which has a merge target that just bounces to the outer
// if merge target then emit an 'else if' instead of a block statement for the else.
- if (auto* inst = i->False()->Instructions().Front()->As<ir::If>();
- inst && inst->Merge()->IsTrampoline(i->Merge())) {
- auto* f = If(inst);
- if (!f) {
- return nullptr;
+ if (auto* inst = i->False()->Instructions().Front()->As<ir::If>()) {
+ if (auto* br = inst->Merge()->Branch()->As<ir::ExitIf>(); br && br->If() == i) {
+ auto* f = If(inst);
+ if (!f) {
+ return nullptr;
+ }
+ return b.If(cond, t, b.Else(f));
}
- return b.If(cond, t, b.Else(f));
} else {
auto* f = BlockGraph(i->False());
if (!f) {
@@ -192,7 +178,6 @@
return b.If(cond, t, b.Else(f));
}
}
-
return b.If(cond, t);
}
@@ -265,17 +250,6 @@
return b.Return(val);
}
- /// @return true if there are no instructions between @p node and and @p stop_at
- bool IsEmpty(const ir::Block* node, const ir::Block* stop_at) {
- if (node->Instructions().IsEmpty()) {
- return true;
- }
- if (auto* br = node->Instructions().Front()->As<Branch>()) {
- return !br->Is<ir::Return>() && br->To() == stop_at;
- }
- return false;
- }
-
utils::Result<const ast::Statement*> Stmt(const ir::Instruction* inst) {
return tint::Switch<utils::Result<const ast::Statement*>>(
inst, //
diff --git a/src/tint/reader/wgsl/lexer.cc b/src/tint/reader/wgsl/lexer.cc
index 1e1dad4..4ed4786 100644
--- a/src/tint/reader/wgsl/lexer.cc
+++ b/src/tint/reader/wgsl/lexer.cc
@@ -25,9 +25,9 @@
#include <type_traits>
#include <utility>
-#include "absl/strings/charconv.h"
#include "src/tint/debug.h"
#include "src/tint/number.h"
+#include "src/tint/utils/parse_num.h"
#include "src/tint/utils/unicode.h"
namespace tint::reader::wgsl {
@@ -414,12 +414,13 @@
end_ptr = &at(length() - 1) + 1;
}
- double value = 0;
- auto ret = absl::from_chars(&at(start), end_ptr, value);
- bool overflow = ret.ec != std::errc();
+ auto ret = utils::ParseDouble(std::string_view(&at(start), end - start));
+ double value = ret ? ret.Get() : 0.0;
+ bool overflow = !ret && ret.Failure() == utils::ParseNumberError::kResultOutOfRange;
- // Value didn't fit in a double, check for underflow as that is 0.0 in WGSL and not an error.
- if (ret.ec == std::errc::result_out_of_range) {
+ // If the value didn't fit in a double, check for underflow as that is 0.0 in WGSL and not an
+ // error.
+ if (overflow) {
// The exponent is negative, so treat as underflow
if (negative_exponent) {
overflow = false;
@@ -446,7 +447,6 @@
}
}
- TINT_ASSERT(Reader, end_ptr == ret.ptr);
advance(end - start);
if (has_f_suffix) {
diff --git a/src/tint/resolver/dependency_graph.cc b/src/tint/resolver/dependency_graph.cc
index 5ba9ee2..504e8ac 100644
--- a/src/tint/resolver/dependency_graph.cc
+++ b/src/tint/resolver/dependency_graph.cc
@@ -66,6 +66,7 @@
#include "src/tint/utils/defer.h"
#include "src/tint/utils/map.h"
#include "src/tint/utils/scoped_assignment.h"
+#include "src/tint/utils/string.h"
#include "src/tint/utils/string_stream.h"
#include "src/tint/utils/unique_vector.h"
diff --git a/src/tint/resolver/resolver.cc b/src/tint/resolver/resolver.cc
index 09d8464..fb06435 100644
--- a/src/tint/resolver/resolver.cc
+++ b/src/tint/resolver/resolver.cc
@@ -3761,8 +3761,9 @@
} else {
utils::StringStream ss;
ss << "unrecognized diagnostic rule 'chromium." << name << "'\n";
- utils::SuggestAlternatives(name, builtin::kChromiumDiagnosticRuleStrings, ss,
- "chromium.");
+ utils::SuggestAlternativeOptions opts;
+ opts.prefix = "chromium.";
+ utils::SuggestAlternatives(name, builtin::kChromiumDiagnosticRuleStrings, ss, opts);
AddWarning(ss.str(), control.rule_name->source);
}
}
diff --git a/src/tint/utils/parse_num.cc b/src/tint/utils/parse_num.cc
new file mode 100644
index 0000000..fa04344
--- /dev/null
+++ b/src/tint/utils/parse_num.cc
@@ -0,0 +1,98 @@
+// 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/utils/parse_num.h"
+
+#include <charconv>
+
+#include "absl/strings/charconv.h"
+
+namespace tint::utils {
+
+namespace {
+
+template <typename T>
+Result<T, ParseNumberError> Parse(std::string_view number) {
+ T val = 0;
+ if constexpr (std::is_floating_point_v<T>) {
+ auto result = absl::from_chars(number.data(), number.data() + number.size(), val);
+ if (result.ec == std::errc::result_out_of_range) {
+ return ParseNumberError::kResultOutOfRange;
+ }
+ if (result.ec != std::errc() || result.ptr != number.data() + number.size()) {
+ return ParseNumberError::kUnparsable;
+ }
+ } else {
+ auto result = std::from_chars(number.data(), number.data() + number.size(), val);
+ if (result.ec == std::errc::result_out_of_range) {
+ return ParseNumberError::kResultOutOfRange;
+ }
+ if (result.ec != std::errc() || result.ptr != number.data() + number.size()) {
+ return ParseNumberError::kUnparsable;
+ }
+ }
+ return val;
+}
+
+} // namespace
+
+Result<float, ParseNumberError> ParseFloat(std::string_view str) {
+ return Parse<float>(str);
+}
+
+Result<double, ParseNumberError> ParseDouble(std::string_view str) {
+ return Parse<double>(str);
+}
+
+Result<int, ParseNumberError> ParseInt(std::string_view str) {
+ return Parse<int>(str);
+}
+
+Result<unsigned int, ParseNumberError> ParseUint(std::string_view str) {
+ return Parse<unsigned int>(str);
+}
+
+Result<int64_t, ParseNumberError> ParseInt64(std::string_view str) {
+ return Parse<int64_t>(str);
+}
+
+Result<uint64_t, ParseNumberError> ParseUint64(std::string_view str) {
+ return Parse<uint64_t>(str);
+}
+
+Result<int32_t, ParseNumberError> ParseInt32(std::string_view str) {
+ return Parse<int32_t>(str);
+}
+
+Result<uint32_t, ParseNumberError> ParseUint32(std::string_view str) {
+ return Parse<uint32_t>(str);
+}
+
+Result<int16_t, ParseNumberError> ParseInt16(std::string_view str) {
+ return Parse<int16_t>(str);
+}
+
+Result<uint16_t, ParseNumberError> ParseUint16(std::string_view str) {
+ return Parse<uint16_t>(str);
+}
+
+Result<int8_t, ParseNumberError> ParseInt8(std::string_view str) {
+ return Parse<int8_t>(str);
+}
+
+Result<uint8_t, ParseNumberError> ParseUint8(std::string_view str) {
+ return Parse<uint8_t>(str);
+}
+
+} // namespace tint::utils
diff --git a/src/tint/utils/parse_num.h b/src/tint/utils/parse_num.h
new file mode 100644
index 0000000..6d4fcb4
--- /dev/null
+++ b/src/tint/utils/parse_num.h
@@ -0,0 +1,126 @@
+// 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_UTILS_PARSE_NUM_H_
+#define SRC_TINT_UTILS_PARSE_NUM_H_
+
+#include <optional>
+#include <string>
+
+#include "src/tint/utils/result.h"
+
+namespace tint::utils {
+
+/// Error returned by the number parsing functions
+enum class ParseNumberError {
+ /// The number was unparsable
+ kUnparsable,
+ /// The parsed number is not representable by the target datatype
+ kResultOutOfRange,
+};
+
+/// @param str the string
+/// @returns the string @p str parsed as a float
+Result<float, ParseNumberError> ParseFloat(std::string_view str);
+
+/// @param str the string
+/// @returns the string @p str parsed as a double
+Result<double, ParseNumberError> ParseDouble(std::string_view str);
+
+/// @param str the string
+/// @returns the string @p str parsed as a int
+Result<int, ParseNumberError> ParseInt(std::string_view str);
+
+/// @param str the string
+/// @returns the string @p str parsed as a unsigned int
+Result<unsigned int, ParseNumberError> ParseUint(std::string_view str);
+
+/// @param str the string
+/// @returns the string @p str parsed as a int64_t
+Result<int64_t, ParseNumberError> ParseInt64(std::string_view str);
+
+/// @param str the string
+/// @returns the string @p str parsed as a uint64_t
+Result<uint64_t, ParseNumberError> ParseUint64(std::string_view str);
+
+/// @param str the string
+/// @returns the string @p str parsed as a int32_t
+Result<int32_t, ParseNumberError> ParseInt32(std::string_view str);
+
+/// @param str the string
+/// @returns the string @p str parsed as a uint32_t
+Result<uint32_t, ParseNumberError> ParseUint32(std::string_view str);
+
+/// @param str the string
+/// @returns the string @p str parsed as a int16_t
+Result<int16_t, ParseNumberError> ParseInt16(std::string_view str);
+
+/// @param str the string
+/// @returns the string @p str parsed as a uint16_t
+Result<uint16_t, ParseNumberError> ParseUint16(std::string_view str);
+
+/// @param str the string
+/// @returns the string @p str parsed as a int8_t
+Result<int8_t, ParseNumberError> ParseInt8(std::string_view str);
+
+/// @param str the string
+/// @returns the string @p str parsed as a uint8_t
+Result<uint8_t, ParseNumberError> ParseUint8(std::string_view str);
+
+/// @param str the string
+/// @returns the string @p str parsed as a the number @p T
+template <typename T>
+inline Result<T, ParseNumberError> ParseNumber(std::string_view str) {
+ if constexpr (std::is_same_v<T, float>) {
+ return ParseFloat(str);
+ }
+ if constexpr (std::is_same_v<T, double>) {
+ return ParseDouble(str);
+ }
+ if constexpr (std::is_same_v<T, int>) {
+ return ParseInt(str);
+ }
+ if constexpr (std::is_same_v<T, unsigned int>) {
+ return ParseUint(str);
+ }
+ if constexpr (std::is_same_v<T, int64_t>) {
+ return ParseInt64(str);
+ }
+ if constexpr (std::is_same_v<T, uint64_t>) {
+ return ParseUint64(str);
+ }
+ if constexpr (std::is_same_v<T, int32_t>) {
+ return ParseInt32(str);
+ }
+ if constexpr (std::is_same_v<T, uint32_t>) {
+ return ParseUint32(str);
+ }
+ if constexpr (std::is_same_v<T, int16_t>) {
+ return ParseInt16(str);
+ }
+ if constexpr (std::is_same_v<T, uint16_t>) {
+ return ParseUint16(str);
+ }
+ if constexpr (std::is_same_v<T, int8_t>) {
+ return ParseInt8(str);
+ }
+ if constexpr (std::is_same_v<T, uint8_t>) {
+ return ParseUint8(str);
+ }
+ return ParseNumberError::kUnparsable;
+}
+
+} // namespace tint::utils
+
+#endif // SRC_TINT_UTILS_PARSE_NUM_H_
diff --git a/src/tint/utils/string.cc b/src/tint/utils/string.cc
index 67eaf2b..00b51a8 100644
--- a/src/tint/utils/string.cc
+++ b/src/tint/utils/string.cc
@@ -15,6 +15,7 @@
#include <algorithm>
#include "src/tint/utils/string.h"
+#include "src/tint/utils/transform.h"
#include "src/tint/utils/vector.h"
namespace tint::utils {
@@ -51,33 +52,46 @@
void SuggestAlternatives(std::string_view got,
Slice<char const* const> strings,
utils::StringStream& ss,
- std::string_view prefix /* = "" */) {
+ const SuggestAlternativeOptions& options /* = {} */) {
+ auto views = Transform<8>(strings, [](char const* const str) { return std::string_view(str); });
+ SuggestAlternatives(got, views.Slice(), ss, options);
+}
+
+void SuggestAlternatives(std::string_view got,
+ Slice<std::string_view> strings,
+ utils::StringStream& ss,
+ const SuggestAlternativeOptions& options /* = {} */) {
// If the string typed was within kSuggestionDistance of one of the possible enum values,
// suggest that. Don't bother with suggestions if the string was extremely long.
constexpr size_t kSuggestionDistance = 5;
constexpr size_t kSuggestionMaxLength = 64;
if (!got.empty() && got.size() < kSuggestionMaxLength) {
size_t candidate_dist = kSuggestionDistance;
- const char* candidate = nullptr;
- for (auto* str : strings) {
+ std::string_view candidate;
+ for (auto str : strings) {
auto dist = utils::Distance(str, got);
if (dist < candidate_dist) {
candidate = str;
candidate_dist = dist;
}
}
- if (candidate) {
- ss << "Did you mean '" << prefix << candidate << "'?\n";
+ if (!candidate.empty()) {
+ ss << "Did you mean '" << options.prefix << candidate << "'?";
+ if (options.list_possible_values) {
+ ss << "\n";
+ }
}
}
- // List all the possible enumerator values
- ss << "Possible values: ";
- for (auto* str : strings) {
- if (str != strings[0]) {
- ss << ", ";
+ if (options.list_possible_values) {
+ // List all the possible enumerator values
+ ss << "Possible values: ";
+ for (auto str : strings) {
+ if (str != strings[0]) {
+ ss << ", ";
+ }
+ ss << "'" << options.prefix << str << "'";
}
- ss << "'" << prefix << str << "'";
}
}
diff --git a/src/tint/utils/string.h b/src/tint/utils/string.h
index e731365..99c930a 100644
--- a/src/tint/utils/string.h
+++ b/src/tint/utils/string.h
@@ -20,6 +20,7 @@
#include "src/tint/utils/slice.h"
#include "src/tint/utils/string_stream.h"
+#include "src/tint/utils/vector.h"
namespace tint::utils {
@@ -38,6 +39,12 @@
return str;
}
+/// @param value the boolean value to be printed as a string
+/// @returns value printed as a string via the stream `<<` operator
+inline std::string ToString(bool value) {
+ return value ? "true" : "false";
+}
+
/// @param value the value to be printed as a string
/// @returns value printed as a string via the stream `<<` operator
template <typename T>
@@ -75,15 +82,33 @@
/// @returns the Levenshtein distance between @p a and @p b
size_t Distance(std::string_view a, std::string_view b);
+/// Options for SuggestAlternatives()
+struct SuggestAlternativeOptions {
+ /// The prefix to apply to the strings when printing
+ std::string_view prefix;
+ /// List all the possible values
+ bool list_possible_values = true;
+};
+
/// Suggest alternatives for an unrecognized string from a list of possible values.
/// @param got the unrecognized string
/// @param strings the list of possible values
/// @param ss the stream to write the suggest and list of possible values to
-/// @param prefix the prefix to apply to the strings when printing (optional)
+/// @param options options for the suggestion
void SuggestAlternatives(std::string_view got,
Slice<char const* const> strings,
utils::StringStream& ss,
- std::string_view prefix = "");
+ const SuggestAlternativeOptions& options = {});
+
+/// Suggest alternatives for an unrecognized string from a list of possible values.
+/// @param got the unrecognized string
+/// @param strings the list of possible values
+/// @param ss the stream to write the suggest and list of possible values to
+/// @param options options for the suggestion
+void SuggestAlternatives(std::string_view got,
+ Slice<std::string_view> strings,
+ utils::StringStream& ss,
+ const SuggestAlternativeOptions& options = {});
/// @param str the input string
/// @param pred the predicate function
@@ -150,6 +175,51 @@
return Trim(str, IsSpace);
}
+/// @param str the input string
+/// @param delimiter the delimiter
+/// @return @p str split at each occurrence of @p delimiter
+inline utils::Vector<std::string_view, 8> Split(std::string_view str, std::string_view delimiter) {
+ utils::Vector<std::string_view, 8> out;
+ while (str.length() > delimiter.length()) {
+ auto pos = str.find(delimiter);
+ if (pos == std::string_view::npos) {
+ break;
+ }
+ out.Push(str.substr(0, pos));
+ str = str.substr(pos + delimiter.length());
+ }
+ out.Push(str);
+ return out;
+}
+
+/// @returns @p str quoted with <code>'</code>
+inline std::string Quote(std::string_view str) {
+ return "'" + std::string(str) + "'";
+}
+
+/// @param parts the input parts
+/// @param delimiter the delimiter
+/// @return @p parts joined as a string, delimited with @p delimiter
+template <typename T>
+inline std::string Join(utils::VectorRef<T> parts, std::string_view delimiter) {
+ utils::StringStream s;
+ for (auto& part : parts) {
+ if (part != parts.Front()) {
+ s << delimiter;
+ }
+ s << part;
+ }
+ return s.str();
+}
+
+/// @param parts the input parts
+/// @param delimiter the delimiter
+/// @return @p parts joined as a string, delimited with @p delimiter
+template <typename T, size_t N>
+inline std::string Join(const utils::Vector<T, N>& parts, std::string_view delimiter) {
+ return Join(utils::VectorRef<T>(parts), delimiter);
+}
+
} // namespace tint::utils
#endif // SRC_TINT_UTILS_STRING_H_
diff --git a/src/tint/utils/string_stream.h b/src/tint/utils/string_stream.h
index 7cbce09..ecb88f7 100644
--- a/src/tint/utils/string_stream.h
+++ b/src/tint/utils/string_stream.h
@@ -24,6 +24,7 @@
#include <utility>
#include "src/tint/utils/unicode.h"
+#include "src/tint/utils/vector.h"
namespace tint::utils {
@@ -189,6 +190,44 @@
/// @returns out so calls can be chained
utils::StringStream& operator<<(utils::StringStream& out, CodePoint codepoint);
+/// Prints the vector @p vec to @p o
+/// @param o the stream to write to
+/// @param vec the vector
+/// @return the stream so calls can be chained
+template <typename T, size_t N>
+inline utils::StringStream& operator<<(utils::StringStream& o, const utils::Vector<T, N>& vec) {
+ o << "[";
+ bool first = true;
+ for (auto& el : vec) {
+ if (!first) {
+ o << ", ";
+ }
+ first = false;
+ o << el;
+ }
+ o << "]";
+ return o;
+}
+
+/// Prints the vector @p vec to @p o
+/// @param o the stream to write to
+/// @param vec the vector reference
+/// @return the stream so calls can be chained
+template <typename T>
+inline utils::StringStream& operator<<(utils::StringStream& o, utils::VectorRef<T> vec) {
+ o << "[";
+ bool first = true;
+ for (auto& el : vec) {
+ if (!first) {
+ o << ", ";
+ }
+ first = false;
+ o << el;
+ }
+ o << "]";
+ return o;
+}
+
} // namespace tint::utils
#endif // SRC_TINT_UTILS_STRING_STREAM_H_
diff --git a/src/tint/utils/string_test.cc b/src/tint/utils/string_test.cc
index 9cf8370..676c341 100644
--- a/src/tint/utils/string_test.cc
+++ b/src/tint/utils/string_test.cc
@@ -14,7 +14,7 @@
#include "src/tint/utils/string.h"
-#include "gtest/gtest.h"
+#include "gmock/gmock.h"
#include "src/tint/utils/string_stream.h"
namespace tint::utils {
@@ -34,6 +34,8 @@
}
TEST(StringTest, ToString) {
+ EXPECT_EQ("true", ToString(true));
+ EXPECT_EQ("false", ToString(false));
EXPECT_EQ("123", ToString(123));
EXPECT_EQ("hello", ToString("hello"));
}
@@ -82,6 +84,23 @@
SuggestAlternatives("hello world", alternatives, ss);
EXPECT_EQ(ss.str(), R"(Possible values: 'foobar', 'something else')");
}
+ {
+ const char* alternatives[] = {"hello world", "Hello World"};
+ utils::StringStream ss;
+ SuggestAlternativeOptions opts;
+ opts.prefix = "$";
+ SuggestAlternatives("hello wordl", alternatives, ss, opts);
+ EXPECT_EQ(ss.str(), R"(Did you mean '$hello world'?
+Possible values: '$hello world', '$Hello World')");
+ }
+ {
+ const char* alternatives[] = {"hello world", "Hello World"};
+ utils::StringStream ss;
+ SuggestAlternativeOptions opts;
+ opts.list_possible_values = false;
+ SuggestAlternatives("hello world", alternatives, ss, opts);
+ EXPECT_EQ(ss.str(), R"(Did you mean 'hello world'?)");
+ }
}
TEST(StringTest, TrimLeft) {
@@ -153,5 +172,27 @@
EXPECT_EQ(TrimSpace(""), "");
}
+TEST(StringTest, Quote) {
+ EXPECT_EQ("'meow'", Quote("meow"));
+}
+
+#if 0 // Enable when moved to C++20 (https://github.com/google/googletest/issues/3081)
+TEST(StringTest, Split) {
+ EXPECT_THAT(Split("", ","), testing::ElementsAre(""));
+ EXPECT_THAT(Split("cat", ","), testing::ElementsAre("cat"));
+ EXPECT_THAT(Split("cat,", ","), testing::ElementsAre("cat", ""));
+ EXPECT_THAT(Split(",cat", ","), testing::ElementsAre("", "cat"));
+ EXPECT_THAT(Split("cat,dog,fish", ","), testing::ElementsAre("cat", "dog", "fish"));
+ EXPECT_THAT(Split("catdogfish", "dog"), testing::ElementsAre("cat", "fish"));
+}
+#endif
+
+TEST(StringTest, Join) {
+ EXPECT_EQ(Join(utils::Vector<int, 1>{}, ","), "");
+ EXPECT_EQ(Join(utils::Vector{1, 2, 3}, ","), "1,2,3");
+ EXPECT_EQ(Join(utils::Vector{"cat"}, ","), "cat");
+ EXPECT_EQ(Join(utils::Vector{"cat", "dog"}, ","), "cat,dog");
+}
+
} // namespace
} // namespace tint::utils
diff --git a/src/tint/utils/transform.h b/src/tint/utils/transform.h
index 9615471..7ecd941 100644
--- a/src/tint/utils/transform.h
+++ b/src/tint/utils/transform.h
@@ -85,6 +85,36 @@
return result;
}
+/// Transform performs an element-wise transformation of a slice.
+/// @param in the input slice.
+/// @param transform the transformation function with signature: `OUT(IN)`
+/// @tparam N the small-array size of the returned Vector
+/// @returns a new vector with each element of the source vector transformed by `transform`.
+template <size_t N, typename IN, typename TRANSFORMER>
+auto Transform(Slice<IN> in, TRANSFORMER&& transform) -> Vector<decltype(transform(in[0])), N> {
+ Vector<decltype(transform(in[0])), N> result;
+ result.Reserve(in.len);
+ for (size_t i = 0; i < in.len; ++i) {
+ result.Push(transform(in[i]));
+ }
+ return result;
+}
+
+/// Transform performs an element-wise transformation of a slice.
+/// @param in the input slice.
+/// @param transform the transformation function with signature: `OUT(IN, size_t)`
+/// @tparam N the small-array size of the returned Vector
+/// @returns a new vector with each element of the source vector transformed by `transform`.
+template <size_t N, typename IN, typename TRANSFORMER>
+auto Transform(Slice<IN> in, TRANSFORMER&& transform) -> Vector<decltype(transform(in[0], 1u)), N> {
+ Vector<decltype(transform(in[0], 1u)), N> result;
+ result.Reserve(in.len);
+ for (size_t i = 0; i < in.len; ++i) {
+ result.Push(transform(in[i], i));
+ }
+ return result;
+}
+
/// Transform performs an element-wise transformation of a vector reference.
/// @param in the input vector.
/// @param transform the transformation function with signature: `OUT(IN)`
@@ -92,13 +122,7 @@
/// @returns a new vector with each element of the source vector transformed by `transform`.
template <size_t N, typename IN, typename TRANSFORMER>
auto Transform(VectorRef<IN> in, TRANSFORMER&& transform) -> Vector<decltype(transform(in[0])), N> {
- const auto count = in.Length();
- Vector<decltype(transform(in[0])), N> result;
- result.Reserve(count);
- for (size_t i = 0; i < count; ++i) {
- result.Push(transform(in[i]));
- }
- return result;
+ return Transform<N>(in.Slice(), std::forward<TRANSFORMER>(transform));
}
/// Transform performs an element-wise transformation of a vector reference.
@@ -109,13 +133,7 @@
template <size_t N, typename IN, typename TRANSFORMER>
auto Transform(VectorRef<IN> in, TRANSFORMER&& transform)
-> Vector<decltype(transform(in[0], 1u)), N> {
- const auto count = in.Length();
- Vector<decltype(transform(in[0], 1u)), N> result;
- result.Reserve(count);
- for (size_t i = 0; i < count; ++i) {
- result.Push(transform(in[i], i));
- }
- return result;
+ return Transform<N>(in.Slice(), std::forward<TRANSFORMER>(transform));
}
/// TransformN performs an element-wise transformation of a vector, transforming and returning at
diff --git a/src/tint/utils/vector.h b/src/tint/utils/vector.h
index 65595e7..acb0e39 100644
--- a/src/tint/utils/vector.h
+++ b/src/tint/utils/vector.h
@@ -26,8 +26,6 @@
#include "src/tint/utils/bitcast.h"
#include "src/tint/utils/compiler_macros.h"
#include "src/tint/utils/slice.h"
-#include "src/tint/utils/string.h"
-#include "src/tint/utils/string_stream.h"
namespace tint::utils {
@@ -587,12 +585,9 @@
/// Aside from this move pattern, a VectorRef provides an immutable reference to the Vector.
template <typename T>
class VectorRef {
- /// The slice type used by this vector reference
- using Slice = utils::Slice<T>;
-
/// @returns an empty slice.
- static Slice& EmptySlice() {
- static Slice empty;
+ static utils::Slice<T>& EmptySlice() {
+ static utils::Slice<T> empty;
return empty;
}
@@ -608,7 +603,7 @@
/// Constructor from a Slice
/// @param slice the slice
- VectorRef(Slice& slice) // NOLINT(runtime/explicit)
+ VectorRef(utils::Slice<T>& slice) // NOLINT(runtime/explicit)
: slice_(slice) {}
/// Constructor from a Vector
@@ -621,7 +616,7 @@
/// @param vector the vector to create a reference of
template <size_t N>
VectorRef(const Vector<T, N>& vector) // NOLINT(runtime/explicit)
- : slice_(const_cast<Slice&>(vector.impl_.slice)) {}
+ : slice_(const_cast<utils::Slice<T>&>(vector.impl_.slice)) {}
/// Constructor from a moved Vector
/// @param vector the vector being moved
@@ -689,6 +684,9 @@
return {slice_.template Reinterpret<U, ReinterpretMode::kUnsafe>()};
}
+ /// @returns the internal slice of the vector
+ utils::Slice<T> Slice() { return slice_; }
+
/// @returns true if the vector is empty.
bool IsEmpty() const { return slice_.len == 0; }
@@ -724,7 +722,7 @@
friend class VectorRef;
/// The slice of the vector being referenced.
- Slice& slice_;
+ utils::Slice<T>& slice_;
/// Whether the slice data is passed by r-value reference, and can be moved.
bool can_move_ = false;
};
@@ -753,44 +751,6 @@
return out;
}
-/// Prints the vector @p vec to @p o
-/// @param o the stream to write to
-/// @param vec the vector
-/// @return the stream so calls can be chained
-template <typename T, size_t N>
-inline utils::StringStream& operator<<(utils::StringStream& o, const utils::Vector<T, N>& vec) {
- o << "[";
- bool first = true;
- for (auto& el : vec) {
- if (!first) {
- o << ", ";
- }
- first = false;
- o << ToString(el);
- }
- o << "]";
- return o;
-}
-
-/// Prints the vector @p vec to @p o
-/// @param o the stream to write to
-/// @param vec the vector reference
-/// @return the stream so calls can be chained
-template <typename T>
-inline utils::StringStream& operator<<(utils::StringStream& o, utils::VectorRef<T> vec) {
- o << "[";
- bool first = true;
- for (auto& el : vec) {
- if (!first) {
- o << ", ";
- }
- first = false;
- o << ToString(el);
- }
- o << "]";
- return o;
-}
-
} // namespace tint::utils
#endif // SRC_TINT_UTILS_VECTOR_H_
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir.cc b/src/tint/writer/spirv/ir/generator_impl_ir.cc
index 29c89e6..89a1a1c 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir.cc
@@ -16,12 +16,20 @@
#include <utility>
+#include "spirv/unified1/GLSL.std.450.h"
#include "spirv/unified1/spirv.h"
#include "src/tint/ir/binary.h"
#include "src/tint/ir/block.h"
+#include "src/tint/ir/break_if.h"
+#include "src/tint/ir/builtin.h"
+#include "src/tint/ir/continue.h"
+#include "src/tint/ir/exit_if.h"
+#include "src/tint/ir/exit_loop.h"
#include "src/tint/ir/if.h"
#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/transform/add_empty_entry_point.h"
@@ -329,7 +337,12 @@
auto result = Switch(
inst, //
[&](const ir::Binary* b) { return EmitBinary(b); },
+ [&](const ir::Builtin* b) { return EmitBuiltin(b); },
[&](const ir::Load* l) { return EmitLoad(l); },
+ [&](const ir::Loop* l) {
+ EmitLoop(l);
+ return 0u;
+ },
[&](const ir::Store* s) {
EmitStore(s);
return 0u;
@@ -354,25 +367,41 @@
}
void GeneratorImplIr::EmitBranch(const ir::Branch* b) {
- if (b->Is<ir::Return>()) {
- if (!b->Args().IsEmpty()) {
- TINT_ASSERT(Writer, b->Args().Length() == 1u);
- OperandList operands;
- operands.push_back(Value(b->Args()[0]));
- current_function_.push_inst(spv::Op::OpReturnValue, operands);
- } else {
- current_function_.push_inst(spv::Op::OpReturn, {});
- }
- return;
- }
-
- Switch(
- b->To(),
- [&](const ir::Block* blk) { current_function_.push_inst(spv::Op::OpBranch, {Label(blk)}); },
+ tint::Switch( //
+ b, //
+ [&](const ir::Return*) {
+ if (!b->Args().IsEmpty()) {
+ TINT_ASSERT(Writer, b->Args().Length() == 1u);
+ OperandList operands;
+ operands.push_back(Value(b->Args()[0]));
+ current_function_.push_inst(spv::Op::OpReturnValue, operands);
+ } else {
+ current_function_.push_inst(spv::Op::OpReturn, {});
+ }
+ return;
+ },
+ [&](const ir::BreakIf* breakif) {
+ current_function_.push_inst(spv::Op::OpBranchConditional,
+ {
+ Value(breakif->Condition()),
+ Label(breakif->Loop()->Merge()),
+ Label(breakif->Loop()->Start()),
+ });
+ },
+ [&](const ir::Continue* cont) {
+ current_function_.push_inst(spv::Op::OpBranch, {Label(cont->Loop()->Continuing())});
+ },
+ [&](const ir::ExitIf* if_) {
+ current_function_.push_inst(spv::Op::OpBranch, {Label(if_->If()->Merge())});
+ },
+ [&](const ir::ExitLoop* loop) {
+ current_function_.push_inst(spv::Op::OpBranch, {Label(loop->Loop()->Merge())});
+ },
+ [&](const ir::NextIteration* loop) {
+ current_function_.push_inst(spv::Op::OpBranch, {Label(loop->Loop()->Start())});
+ },
[&](Default) {
- // A block may not have an outward branch (e.g. an unreachable merge
- // block).
- current_function_.push_inst(spv::Op::OpUnreachable, {});
+ TINT_ICE(Writer, diagnostics_) << "unimplemented branch: " << b->TypeInfo().name;
});
}
@@ -388,10 +417,12 @@
uint32_t merge_label = Label(merge_block);
uint32_t true_label = merge_label;
uint32_t false_label = merge_label;
- if (true_block->Instructions().Length() > 1 || true_block->Branch()->To() != merge_block) {
+ if (true_block->Instructions().Length() > 1 ||
+ (true_block->HasBranchTarget() && !true_block->Branch()->Is<ir::ExitIf>())) {
true_label = Label(true_block);
}
- if (false_block->Instructions().Length() > 1 || false_block->Branch()->To() != merge_block) {
+ if (false_block->Instructions().Length() > 1 ||
+ (false_block->HasBranchTarget() && !false_block->Branch()->Is<ir::ExitIf>())) {
false_label = Label(false_block);
}
@@ -516,12 +547,110 @@
return id;
}
+uint32_t GeneratorImplIr::EmitBuiltin(const ir::Builtin* builtin) {
+ auto id = module_.NextId();
+ auto* result_ty = builtin->Type();
+
+ spv::Op op = spv::Op::Max;
+ OperandList operands = {Type(result_ty), id};
+
+ // Helper to set up the opcode and operand list for a GLSL extended instruction.
+ auto glsl_ext_inst = [&](enum GLSLstd450 inst) {
+ constexpr const char* kGLSLstd450 = "GLSL.std.450";
+ op = spv::Op::OpExtInst;
+ operands.push_back(imports_.GetOrCreate(kGLSLstd450, [&]() {
+ // Import the instruction set the first time it is requested.
+ auto import = module_.NextId();
+ module_.PushExtImport(spv::Op::OpExtInstImport, {import, Operand(kGLSLstd450)});
+ return import;
+ }));
+ operands.push_back(U32Operand(inst));
+ };
+
+ // Determine the opcode.
+ switch (builtin->Func()) {
+ case builtin::Function::kAbs:
+ if (result_ty->is_float_scalar_or_vector()) {
+ glsl_ext_inst(GLSLstd450FAbs);
+ } else if (result_ty->is_signed_integer_scalar_or_vector()) {
+ glsl_ext_inst(GLSLstd450SAbs);
+ } else if (result_ty->is_unsigned_integer_scalar_or_vector()) {
+ // abs() is a no-op for unsigned integers.
+ return Value(builtin->Args()[0]);
+ }
+ break;
+ case builtin::Function::kMax:
+ if (result_ty->is_float_scalar_or_vector()) {
+ glsl_ext_inst(GLSLstd450FMax);
+ } else if (result_ty->is_signed_integer_scalar_or_vector()) {
+ glsl_ext_inst(GLSLstd450SMax);
+ } else if (result_ty->is_unsigned_integer_scalar_or_vector()) {
+ glsl_ext_inst(GLSLstd450UMax);
+ }
+ break;
+ case builtin::Function::kMin:
+ if (result_ty->is_float_scalar_or_vector()) {
+ glsl_ext_inst(GLSLstd450FMin);
+ } else if (result_ty->is_signed_integer_scalar_or_vector()) {
+ glsl_ext_inst(GLSLstd450SMin);
+ } else if (result_ty->is_unsigned_integer_scalar_or_vector()) {
+ glsl_ext_inst(GLSLstd450UMin);
+ }
+ break;
+ default:
+ TINT_ICE(Writer, diagnostics_) << "unimplemented builtin function: " << builtin->Func();
+ return 0u;
+ }
+ TINT_ASSERT(Writer, op != spv::Op::Max);
+
+ // Add the arguments to the builtin call.
+ for (auto* arg : builtin->Args()) {
+ operands.push_back(Value(arg));
+ }
+
+ // Emit the instruction.
+ current_function_.push_inst(op, operands);
+
+ return id;
+}
+
uint32_t GeneratorImplIr::EmitLoad(const ir::Load* load) {
auto id = module_.NextId();
current_function_.push_inst(spv::Op::OpLoad, {Type(load->Type()), id, Value(load->From())});
return id;
}
+void GeneratorImplIr::EmitLoop(const ir::Loop* loop) {
+ auto header_label = module_.NextId();
+ auto body_label = Label(loop->Start());
+ auto continuing_label = Label(loop->Continuing());
+ auto merge_label = Label(loop->Merge());
+
+ // Branch to and emit the loop header, which contains OpLoopMerge and OpBranch instructions.
+ current_function_.push_inst(spv::Op::OpBranch, {header_label});
+ current_function_.push_inst(spv::Op::OpLabel, {header_label});
+ current_function_.push_inst(
+ spv::Op::OpLoopMerge, {merge_label, continuing_label, U32Operand(SpvLoopControlMaskNone)});
+ current_function_.push_inst(spv::Op::OpBranch, {body_label});
+
+ // Emit the loop body.
+ EmitBlock(loop->Start());
+
+ // Emit the loop continuing block.
+ // The back-edge needs to go to the loop header, so update the label for the start block.
+ block_labels_.Replace(loop->Start(), header_label);
+ if (loop->Continuing()->HasBranchTarget()) {
+ EmitBlock(loop->Continuing());
+ } else {
+ // We still need to emit a continuing block with a back-edge, even if it is unreachable.
+ current_function_.push_inst(spv::Op::OpLabel, {continuing_label});
+ current_function_.push_inst(spv::Op::OpBranch, {header_label});
+ }
+
+ // Emit the loop merge block.
+ EmitBlock(loop->Merge());
+}
+
void GeneratorImplIr::EmitStore(const ir::Store* store) {
current_function_.push_inst(spv::Op::OpStore, {Value(store->To()), Value(store->From())});
}
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir.h b/src/tint/writer/spirv/ir/generator_impl_ir.h
index f4ababa..5293fea 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir.h
+++ b/src/tint/writer/spirv/ir/generator_impl_ir.h
@@ -32,9 +32,11 @@
class Binary;
class Block;
class Branch;
+class Builtin;
class If;
class Function;
class Load;
+class Loop;
class Module;
class Store;
class UserCall;
@@ -110,11 +112,20 @@
/// @returns the result ID of the instruction
uint32_t EmitBinary(const ir::Binary* binary);
+ /// Emit a builtin function call instruction.
+ /// @param call the builtin call instruction to emit
+ /// @returns the result ID of the instruction
+ uint32_t EmitBuiltin(const ir::Builtin* call);
+
/// Emit a load instruction.
/// @param load the load instruction to emit
/// @returns the result ID of the instruction
uint32_t EmitLoad(const ir::Load* load);
+ /// Emit a loop instruction.
+ /// @param loop the loop instruction to emit
+ void EmitLoop(const ir::Loop* loop);
+
/// Emit a store instruction.
/// @param store the store instruction to emit
void EmitStore(const ir::Store* store);
@@ -184,6 +195,9 @@
/// The map of blocks to the IDs of their label instructions.
utils::Hashmap<const ir::Block*, uint32_t, 8> block_labels_;
+ /// The map of extended instruction set names to their result IDs.
+ utils::Hashmap<std::string_view, uint32_t, 2> imports_;
+
/// The current function that is being emitted.
Function current_function_;
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc
index cfe3dc4..f6e5015 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc
@@ -22,102 +22,17 @@
namespace tint::writer::spirv {
namespace {
-/// The element type of a test.
-enum Type {
- kBool,
- kI32,
- kU32,
- kF32,
- kF16,
-};
-
/// A parameterized test case.
struct BinaryTestCase {
/// The element type to test.
- Type type;
+ TestElementType type;
/// The binary operation.
enum ir::Binary::Kind kind;
/// The expected SPIR-V instruction.
std::string spirv_inst;
};
-/// A helper class for parameterized binary instruction tests.
-class BinaryInstructionTest : public SpvGeneratorImplTestWithParam<BinaryTestCase> {
- protected:
- /// Helper to make a scalar type corresponding to the element type `ty`.
- /// @param ty the element type
- /// @returns the scalar type
- const type::Type* MakeScalarType(Type ty) {
- switch (ty) {
- case kBool:
- return mod.Types().bool_();
- case kI32:
- return mod.Types().i32();
- case kU32:
- return mod.Types().u32();
- case kF32:
- return mod.Types().f32();
- case kF16:
- return mod.Types().f16();
- }
- return nullptr;
- }
-
- /// Helper to make a vector type corresponding to the element type `ty`.
- /// @param ty the element type
- /// @returns the vector type
- const type::Type* MakeVectorType(Type ty) { return mod.Types().vec2(MakeScalarType(ty)); }
-
- /// Helper to make a scalar value with the scalar type `ty`.
- /// @param ty the element type
- /// @returns the scalar value
- ir::Value* MakeScalarValue(Type ty) {
- switch (ty) {
- case kBool:
- return b.Constant(true);
- case kI32:
- return b.Constant(1_i);
- case kU32:
- return b.Constant(1_u);
- case kF32:
- return b.Constant(1_f);
- case kF16:
- return b.Constant(1_h);
- }
- return nullptr;
- }
-
- /// Helper to make a vector value with an element type of `ty`.
- /// @param ty the element type
- /// @returns the vector value
- ir::Value* MakeVectorValue(Type ty) {
- switch (ty) {
- case kBool:
- return b.Constant(b.ir.constant_values.Composite(
- MakeVectorType(ty), utils::Vector{b.ir.constant_values.Get(true),
- b.ir.constant_values.Get(false)}));
- case kI32:
- return b.Constant(b.ir.constant_values.Composite(
- MakeVectorType(ty), utils::Vector{b.ir.constant_values.Get(42_i),
- b.ir.constant_values.Get(-10_i)}));
- case kU32:
- return b.Constant(b.ir.constant_values.Composite(
- MakeVectorType(ty),
- utils::Vector{b.ir.constant_values.Get(42_u), b.ir.constant_values.Get(10_u)}));
- case kF32:
- return b.Constant(b.ir.constant_values.Composite(
- MakeVectorType(ty), utils::Vector{b.ir.constant_values.Get(42_f),
- b.ir.constant_values.Get(-0.5_f)}));
- case kF16:
- return b.Constant(b.ir.constant_values.Composite(
- MakeVectorType(ty), utils::Vector{b.ir.constant_values.Get(42_h),
- b.ir.constant_values.Get(-0.5_h)}));
- }
- return nullptr;
- }
-};
-
-using Arithmetic = BinaryInstructionTest;
+using Arithmetic = SpvGeneratorImplTestWithParam<BinaryTestCase>;
TEST_P(Arithmetic, Scalar) {
auto params = GetParam();
@@ -164,7 +79,7 @@
BinaryTestCase{kF16, ir::Binary::Kind::kSubtract,
"OpFSub"}));
-using Bitwise = BinaryInstructionTest;
+using Bitwise = SpvGeneratorImplTestWithParam<BinaryTestCase>;
TEST_P(Bitwise, Scalar) {
auto params = GetParam();
@@ -203,7 +118,7 @@
BinaryTestCase{kU32, ir::Binary::Kind::kOr, "OpBitwiseOr"},
BinaryTestCase{kU32, ir::Binary::Kind::kXor, "OpBitwiseXor"}));
-using Comparison = BinaryInstructionTest;
+using Comparison = SpvGeneratorImplTestWithParam<BinaryTestCase>;
TEST_P(Comparison, Scalar) {
auto params = GetParam();
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_builtin_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_builtin_test.cc
new file mode 100644
index 0000000..f62848e
--- /dev/null
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_builtin_test.cc
@@ -0,0 +1,145 @@
+// 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/writer/spirv/ir/test_helper_ir.h"
+
+#include "gmock/gmock.h"
+#include "src/tint/builtin/function.h"
+
+using namespace tint::number_suffixes; // NOLINT
+
+namespace tint::writer::spirv {
+namespace {
+
+/// A parameterized builtin function test case.
+struct BuiltinTestCase {
+ /// The element type to test.
+ TestElementType type;
+ /// The builtin function.
+ enum builtin::Function function;
+ /// The expected SPIR-V instruction string.
+ std::string spirv_inst;
+};
+
+// Tests for builtins with the signature: T = func(T)
+using Builtin_1arg = SpvGeneratorImplTestWithParam<BuiltinTestCase>;
+TEST_P(Builtin_1arg, Scalar) {
+ auto params = GetParam();
+
+ auto* func = b.CreateFunction("foo", mod.Types().void_());
+ func->StartTarget()->SetInstructions(
+ utils::Vector{b.Builtin(MakeScalarType(params.type), params.function,
+ utils::Vector{MakeScalarValue(params.type)}),
+ b.Return(func)});
+
+ generator_.EmitFunction(func);
+ EXPECT_THAT(DumpModule(generator_.Module()), ::testing::HasSubstr(params.spirv_inst));
+}
+TEST_P(Builtin_1arg, Vector) {
+ auto params = GetParam();
+
+ auto* func = b.CreateFunction("foo", mod.Types().void_());
+ func->StartTarget()->SetInstructions(
+ utils::Vector{b.Builtin(MakeVectorType(params.type), params.function,
+ utils::Vector{MakeVectorValue(params.type)}),
+
+ b.Return(func)});
+
+ generator_.EmitFunction(func);
+ EXPECT_THAT(DumpModule(generator_.Module()), ::testing::HasSubstr(params.spirv_inst));
+}
+INSTANTIATE_TEST_SUITE_P(SpvGeneratorImplTest,
+ Builtin_1arg,
+ testing::Values(BuiltinTestCase{kI32, builtin::Function::kAbs, "SAbs"},
+ BuiltinTestCase{kF32, builtin::Function::kAbs, "FAbs"}));
+
+// Test that abs of an unsigned value just folds away.
+TEST_F(SpvGeneratorImplTest, Builtin_Abs_u32) {
+ auto* result = b.Builtin(MakeScalarType(kU32), builtin::Function::kAbs,
+ utils::Vector{MakeScalarValue(kU32)});
+ auto* func = b.CreateFunction("foo", MakeScalarType(kU32));
+ func->StartTarget()->SetInstructions(
+ utils::Vector{result, b.Return(func, utils::Vector{result})});
+
+ generator_.EmitFunction(func);
+ EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeInt 32 0
+%3 = OpTypeFunction %2
+%6 = OpConstant %2 1
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+OpReturnValue %6
+OpFunctionEnd
+)");
+}
+TEST_F(SpvGeneratorImplTest, Builtin_Abs_vec2u) {
+ auto* result = b.Builtin(MakeVectorType(kU32), builtin::Function::kAbs,
+ utils::Vector{MakeVectorValue(kU32)});
+ auto* func = b.CreateFunction("foo", MakeVectorType(kU32));
+ func->StartTarget()->SetInstructions(
+ utils::Vector{result, b.Return(func, utils::Vector{result})});
+
+ generator_.EmitFunction(func);
+ EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%3 = OpTypeInt 32 0
+%2 = OpTypeVector %3 2
+%4 = OpTypeFunction %2
+%8 = OpConstant %3 42
+%9 = OpConstant %3 10
+%7 = OpConstantComposite %2 %8 %9
+%1 = OpFunction %2 None %4
+%5 = OpLabel
+OpReturnValue %7
+OpFunctionEnd
+)");
+}
+
+// Tests for builtins with the signature: T = func(T, T)
+using Builtin_2arg = SpvGeneratorImplTestWithParam<BuiltinTestCase>;
+TEST_P(Builtin_2arg, Scalar) {
+ auto params = GetParam();
+
+ auto* func = b.CreateFunction("foo", mod.Types().void_());
+ func->StartTarget()->SetInstructions(utils::Vector{
+ b.Builtin(MakeScalarType(params.type), params.function,
+ utils::Vector{MakeScalarValue(params.type), MakeScalarValue(params.type)}),
+ b.Return(func)});
+
+ generator_.EmitFunction(func);
+ EXPECT_THAT(DumpModule(generator_.Module()), ::testing::HasSubstr(params.spirv_inst));
+}
+TEST_P(Builtin_2arg, Vector) {
+ auto params = GetParam();
+
+ auto* func = b.CreateFunction("foo", mod.Types().void_());
+ func->StartTarget()->SetInstructions(utils::Vector{
+ b.Builtin(MakeVectorType(params.type), params.function,
+ utils::Vector{MakeVectorValue(params.type), MakeVectorValue(params.type)}),
+
+ b.Return(func)});
+
+ generator_.EmitFunction(func);
+ EXPECT_THAT(DumpModule(generator_.Module()), ::testing::HasSubstr(params.spirv_inst));
+}
+INSTANTIATE_TEST_SUITE_P(SpvGeneratorImplTest,
+ Builtin_2arg,
+ testing::Values(BuiltinTestCase{kF32, builtin::Function::kMax, "FMax"},
+ BuiltinTestCase{kI32, builtin::Function::kMax, "SMax"},
+ BuiltinTestCase{kU32, builtin::Function::kMax, "UMax"},
+ BuiltinTestCase{kF32, builtin::Function::kMin, "FMin"},
+ BuiltinTestCase{kI32, builtin::Function::kMin, "SMin"},
+ BuiltinTestCase{kU32, builtin::Function::kMin, "UMin"}));
+
+} // namespace
+} // namespace tint::writer::spirv
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc
index b526c91..f43a09a 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc
@@ -23,8 +23,8 @@
auto* func = b.CreateFunction("foo", mod.Types().void_());
auto* i = b.CreateIf(b.Constant(true));
- i->True()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
- i->False()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
+ i->True()->SetInstructions(utils::Vector{b.ExitIf(i)});
+ i->False()->SetInstructions(utils::Vector{b.ExitIf(i)});
i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
func->StartTarget()->SetInstructions(utils::Vector{i});
@@ -49,12 +49,12 @@
auto* func = b.CreateFunction("foo", mod.Types().void_());
auto* i = b.CreateIf(b.Constant(true));
- i->False()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
+ i->False()->SetInstructions(utils::Vector{b.ExitIf(i)});
i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
auto* true_block = i->True();
- true_block->SetInstructions(utils::Vector{
- b.Add(mod.Types().i32(), b.Constant(1_i), b.Constant(1_i)), b.Branch(i->Merge())});
+ true_block->SetInstructions(
+ utils::Vector{b.Add(mod.Types().i32(), b.Constant(1_i), b.Constant(1_i)), b.ExitIf(i)});
func->StartTarget()->SetInstructions(utils::Vector{i});
@@ -83,12 +83,12 @@
auto* func = b.CreateFunction("foo", mod.Types().void_());
auto* i = b.CreateIf(b.Constant(true));
- i->True()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
+ i->True()->SetInstructions(utils::Vector{b.ExitIf(i)});
i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
auto* false_block = i->False();
- false_block->SetInstructions(utils::Vector{
- b.Add(mod.Types().i32(), b.Constant(1_i), b.Constant(1_i)), b.Branch(i->Merge())});
+ false_block->SetInstructions(
+ utils::Vector{b.Add(mod.Types().i32(), b.Constant(1_i), b.Constant(1_i)), b.ExitIf(i)});
func->StartTarget()->SetInstructions(utils::Vector{i});
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_loop_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_loop_test.cc
new file mode 100644
index 0000000..8ec7a9c
--- /dev/null
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_loop_test.cc
@@ -0,0 +1,334 @@
+// 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/writer/spirv/ir/test_helper_ir.h"
+
+using namespace tint::number_suffixes; // NOLINT
+
+namespace tint::writer::spirv {
+namespace {
+
+TEST_F(SpvGeneratorImplTest, Loop_BreakIf) {
+ auto* func = b.CreateFunction("foo", mod.Types().void_());
+
+ auto* loop = b.CreateLoop();
+
+ loop->Start()->Instructions().Push(b.Continue(loop));
+ loop->Continuing()->Instructions().Push(b.BreakIf(b.Constant(true), loop));
+ loop->Merge()->Instructions().Push(b.Return(func));
+
+ func->StartTarget()->Instructions().Push(loop);
+
+ generator_.EmitFunction(func);
+ EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%10 = OpTypeBool
+%9 = OpConstantTrue %10
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+OpBranch %5
+%5 = OpLabel
+OpLoopMerge %8 %7 None
+OpBranch %6
+%6 = OpLabel
+OpBranch %7
+%7 = OpLabel
+OpBranchConditional %9 %8 %5
+%8 = OpLabel
+OpReturn
+OpFunctionEnd
+)");
+}
+
+// Test that we still emit the continuing block with a back-edge, even when it is unreachable.
+TEST_F(SpvGeneratorImplTest, Loop_UnconditionalBreakInBody) {
+ auto* func = b.CreateFunction("foo", mod.Types().void_());
+
+ auto* loop = b.CreateLoop();
+
+ loop->Start()->Instructions().Push(b.ExitLoop(loop));
+ loop->Merge()->Instructions().Push(b.Return(func));
+
+ func->StartTarget()->Instructions().Push(loop);
+
+ generator_.EmitFunction(func);
+ EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+OpBranch %5
+%5 = OpLabel
+OpLoopMerge %8 %7 None
+OpBranch %6
+%6 = OpLabel
+OpBranch %8
+%7 = OpLabel
+OpBranch %5
+%8 = OpLabel
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(SpvGeneratorImplTest, Loop_ConditionalBreakInBody) {
+ auto* func = b.CreateFunction("foo", mod.Types().void_());
+
+ auto* loop = b.CreateLoop();
+
+ auto* cond_break = b.CreateIf(b.Constant(true));
+ cond_break->True()->Instructions().Push(b.ExitLoop(loop));
+ cond_break->False()->Instructions().Push(b.ExitIf(cond_break));
+ cond_break->Merge()->Instructions().Push(b.Continue(loop));
+
+ loop->Start()->Instructions().Push(cond_break);
+ loop->Continuing()->Instructions().Push(b.NextIteration(loop));
+ loop->Merge()->Instructions().Push(b.Return(func));
+
+ func->StartTarget()->Instructions().Push(loop);
+
+ generator_.EmitFunction(func);
+ EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%12 = OpTypeBool
+%11 = OpConstantTrue %12
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+OpBranch %5
+%5 = OpLabel
+OpLoopMerge %8 %7 None
+OpBranch %6
+%6 = OpLabel
+OpSelectionMerge %9 None
+OpBranchConditional %11 %10 %9
+%10 = OpLabel
+OpBranch %8
+%9 = OpLabel
+OpBranch %7
+%7 = OpLabel
+OpBranch %5
+%8 = OpLabel
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(SpvGeneratorImplTest, Loop_ConditionalContinueInBody) {
+ auto* func = b.CreateFunction("foo", mod.Types().void_());
+
+ auto* loop = b.CreateLoop();
+
+ auto* cond_break = b.CreateIf(b.Constant(true));
+ cond_break->True()->Instructions().Push(b.Continue(loop));
+ cond_break->False()->Instructions().Push(b.ExitIf(cond_break));
+ cond_break->Merge()->Instructions().Push(b.ExitLoop(loop));
+
+ loop->Start()->Instructions().Push(cond_break);
+ loop->Continuing()->Instructions().Push(b.NextIteration(loop));
+ loop->Merge()->Instructions().Push(b.Return(func));
+
+ func->StartTarget()->Instructions().Push(loop);
+
+ generator_.EmitFunction(func);
+ EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%12 = OpTypeBool
+%11 = OpConstantTrue %12
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+OpBranch %5
+%5 = OpLabel
+OpLoopMerge %8 %7 None
+OpBranch %6
+%6 = OpLabel
+OpSelectionMerge %9 None
+OpBranchConditional %11 %10 %9
+%10 = OpLabel
+OpBranch %7
+%9 = OpLabel
+OpBranch %8
+%7 = OpLabel
+OpBranch %5
+%8 = OpLabel
+OpReturn
+OpFunctionEnd
+)");
+}
+
+// Test that we still emit the continuing block with a back-edge, and the merge block, even when
+// they are unreachable.
+TEST_F(SpvGeneratorImplTest, Loop_UnconditionalReturnInBody) {
+ auto* func = b.CreateFunction("foo", mod.Types().void_());
+
+ auto* loop = b.CreateLoop();
+
+ loop->Start()->Instructions().Push(b.Return(func));
+
+ func->StartTarget()->Instructions().Push(loop);
+
+ generator_.EmitFunction(func);
+ EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+OpBranch %5
+%5 = OpLabel
+OpLoopMerge %8 %7 None
+OpBranch %6
+%6 = OpLabel
+OpReturn
+%7 = OpLabel
+OpBranch %5
+%8 = OpLabel
+OpUnreachable
+OpFunctionEnd
+)");
+}
+
+TEST_F(SpvGeneratorImplTest, Loop_UseResultFromBodyInContinuing) {
+ auto* func = b.CreateFunction("foo", mod.Types().void_());
+
+ auto* loop = b.CreateLoop();
+
+ auto* result = b.Equal(mod.Types().i32(), b.Constant(1_i), b.Constant(2_i));
+
+ loop->Start()->Instructions().Push(result);
+ loop->Continuing()->Instructions().Push(b.BreakIf(result, loop));
+ loop->Merge()->Instructions().Push(b.Return(func));
+
+ func->StartTarget()->Instructions().Push(loop);
+
+ generator_.EmitFunction(func);
+ EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%10 = OpTypeInt 32 1
+%11 = OpConstant %10 1
+%12 = OpConstant %10 2
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+OpBranch %5
+%5 = OpLabel
+OpLoopMerge %8 %7 None
+OpBranch %6
+%6 = OpLabel
+%9 = OpIEqual %10 %11 %12
+%7 = OpLabel
+OpBranchConditional %9 %8 %5
+%8 = OpLabel
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(SpvGeneratorImplTest, Loop_NestedLoopInBody) {
+ auto* func = b.CreateFunction("foo", mod.Types().void_());
+
+ auto* outer_loop = b.CreateLoop();
+ auto* inner_loop = b.CreateLoop();
+
+ inner_loop->Start()->Instructions().Push(b.ExitLoop(inner_loop));
+ inner_loop->Continuing()->Instructions().Push(b.NextIteration(inner_loop));
+ inner_loop->Merge()->Instructions().Push(b.Continue(outer_loop));
+
+ outer_loop->Start()->Instructions().Push(inner_loop);
+ outer_loop->Continuing()->Instructions().Push(b.BreakIf(b.Constant(true), outer_loop));
+ outer_loop->Merge()->Instructions().Push(b.Return(func));
+
+ func->StartTarget()->Instructions().Push(outer_loop);
+
+ generator_.EmitFunction(func);
+ EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%14 = OpTypeBool
+%13 = OpConstantTrue %14
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+OpBranch %5
+%5 = OpLabel
+OpLoopMerge %8 %7 None
+OpBranch %6
+%6 = OpLabel
+OpBranch %9
+%9 = OpLabel
+OpLoopMerge %12 %11 None
+OpBranch %10
+%10 = OpLabel
+OpBranch %12
+%11 = OpLabel
+OpBranch %9
+%12 = OpLabel
+OpBranch %7
+%7 = OpLabel
+OpBranchConditional %13 %8 %5
+%8 = OpLabel
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(SpvGeneratorImplTest, Loop_NestedLoopInContinuing) {
+ auto* func = b.CreateFunction("foo", mod.Types().void_());
+
+ auto* outer_loop = b.CreateLoop();
+ auto* inner_loop = b.CreateLoop();
+
+ inner_loop->Start()->Instructions().Push(b.Continue(inner_loop));
+ inner_loop->Continuing()->Instructions().Push(b.BreakIf(b.Constant(true), inner_loop));
+ inner_loop->Merge()->Instructions().Push(b.BreakIf(b.Constant(true), outer_loop));
+
+ outer_loop->Start()->Instructions().Push(b.Continue(outer_loop));
+ outer_loop->Continuing()->Instructions().Push(inner_loop);
+ outer_loop->Merge()->Instructions().Push(b.Return(func));
+
+ func->StartTarget()->Instructions().Push(outer_loop);
+
+ generator_.EmitFunction(func);
+ EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%14 = OpTypeBool
+%13 = OpConstantTrue %14
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+OpBranch %5
+%5 = OpLabel
+OpLoopMerge %8 %7 None
+OpBranch %6
+%6 = OpLabel
+OpBranch %7
+%7 = OpLabel
+OpBranch %9
+%9 = OpLabel
+OpLoopMerge %12 %11 None
+OpBranch %10
+%10 = OpLabel
+OpBranch %11
+%11 = OpLabel
+OpBranchConditional %13 %12 %9
+%12 = OpLabel
+OpBranchConditional %13 %8 %5
+%8 = OpLabel
+OpReturn
+OpFunctionEnd
+)");
+}
+
+} // namespace
+} // namespace tint::writer::spirv
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc
index d66df98..574a0b0 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc
@@ -100,7 +100,7 @@
v->SetInitializer(b.Constant(42_i));
auto* i = b.CreateIf(b.Constant(true));
- i->True()->SetInstructions(utils::Vector{v, b.Branch(i->Merge())});
+ i->True()->SetInstructions(utils::Vector{v, b.ExitIf(i)});
i->False()->SetInstructions(utils::Vector{b.Return(func)});
i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
diff --git a/src/tint/writer/spirv/ir/test_helper_ir.h b/src/tint/writer/spirv/ir/test_helper_ir.h
index 9509b42..add9eaf 100644
--- a/src/tint/writer/spirv/ir/test_helper_ir.h
+++ b/src/tint/writer/spirv/ir/test_helper_ir.h
@@ -24,6 +24,15 @@
namespace tint::writer::spirv {
+/// The element type of a test.
+enum TestElementType {
+ kBool,
+ kI32,
+ kU32,
+ kF32,
+ kF16,
+};
+
/// Base helper class for testing the SPIR-V generator implementation.
template <typename BASE>
class SpvGeneratorTestHelperBase : public BASE {
@@ -41,6 +50,85 @@
/// @returns the disassembled types from the generated module.
std::string DumpTypes() { return DumpInstructions(generator_.Module().Types()); }
+
+ /// Helper to make a scalar type corresponding to the element type `ty`.
+ /// @param ty the element type
+ /// @returns the scalar type
+ const type::Type* MakeScalarType(TestElementType ty) {
+ switch (ty) {
+ case kBool:
+ return mod.Types().bool_();
+ case kI32:
+ return mod.Types().i32();
+ case kU32:
+ return mod.Types().u32();
+ case kF32:
+ return mod.Types().f32();
+ case kF16:
+ return mod.Types().f16();
+ }
+ return nullptr;
+ }
+
+ /// Helper to make a vector type corresponding to the element type `ty`.
+ /// @param ty the element type
+ /// @returns the vector type
+ const type::Type* MakeVectorType(TestElementType ty) {
+ return mod.Types().vec2(MakeScalarType(ty));
+ }
+
+ /// Helper to make a scalar value with the scalar type `ty`.
+ /// @param ty the element type
+ /// @returns the scalar value
+ ir::Value* MakeScalarValue(TestElementType ty) {
+ switch (ty) {
+ case kBool:
+ return b.Constant(true);
+ case kI32:
+ return b.Constant(i32(1));
+ case kU32:
+ return b.Constant(u32(1));
+ case kF32:
+ return b.Constant(f32(1));
+ case kF16:
+ return b.Constant(f16(1));
+ }
+ return nullptr;
+ }
+
+ /// Helper to make a vector value with an element type of `ty`.
+ /// @param ty the element type
+ /// @returns the vector value
+ ir::Value* MakeVectorValue(TestElementType ty) {
+ switch (ty) {
+ case kBool:
+ return b.Constant(mod.constant_values.Composite(
+ MakeVectorType(ty),
+ utils::Vector<const constant::Value*, 2>{mod.constant_values.Get(true),
+ mod.constant_values.Get(false)}));
+ case kI32:
+ return b.Constant(mod.constant_values.Composite(
+ MakeVectorType(ty),
+ utils::Vector<const constant::Value*, 2>{mod.constant_values.Get(i32(42)),
+ mod.constant_values.Get(i32(-10))}));
+ case kU32:
+ return b.Constant(mod.constant_values.Composite(
+ MakeVectorType(ty),
+ utils::Vector<const constant::Value*, 2>{mod.constant_values.Get(u32(42)),
+ mod.constant_values.Get(u32(10))}));
+ case kF32:
+ return b.Constant(mod.constant_values.Composite(
+ MakeVectorType(ty),
+ utils::Vector<const constant::Value*, 2>{mod.constant_values.Get(f32(42)),
+ mod.constant_values.Get(f32(-0.5))}));
+ case kF16:
+ return b.Constant(mod.constant_values.Composite(
+ MakeVectorType(ty),
+ utils::Vector<const constant::Value*, 2>{mod.constant_values.Get(f16(42)),
+ mod.constant_values.Get(f16(-0.5))}));
+ }
+ return nullptr;
+ }
};
using SpvGeneratorImplTest = SpvGeneratorTestHelperBase<testing::Test>;
