[spirv-reader][ir] Add OpPhi support.
Add support for `OpPhi`.
Bug: 398889961
Change-Id: I749266373eefc82784cb7bad9aff1874ca956322
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/243894
Auto-Submit: dan sinclair <dsinclair@chromium.org>
Commit-Queue: James Price <jrprice@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/tint/lang/core/ir/operand_instruction.h b/src/tint/lang/core/ir/operand_instruction.h
index 1daf3f4..301df24 100644
--- a/src/tint/lang/core/ir/operand_instruction.h
+++ b/src/tint/lang/core/ir/operand_instruction.h
@@ -83,10 +83,14 @@
/// Appends a new operand
/// @param operand the new operand
- void PushOperand(ir::Value* operand) {
- auto idx = operands_.Length();
+ /// @returns the operand index
+ uint32_t PushOperand(ir::Value* operand) {
+ uint32_t idx = static_cast<uint32_t>(operands_.Length());
operands_.Push(operand);
- operand->AddUsage({this, static_cast<uint32_t>(idx)});
+ if (operand) {
+ operand->AddUsage({this, idx});
+ }
+ return idx;
}
/// Removes all operands from the instruction
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.bazel b/src/tint/lang/spirv/reader/parser/BUILD.bazel
index 1f7b7c4..48335aa 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.bazel
+++ b/src/tint/lang/spirv/reader/parser/BUILD.bazel
@@ -105,6 +105,7 @@
"memory_test.cc",
"misc_test.cc",
"name_test.cc",
+ "phi_test.cc",
"struct_test.cc",
"unary_test.cc",
"var_test.cc",
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.cmake b/src/tint/lang/spirv/reader/parser/BUILD.cmake
index 66dc68d..c832152 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.cmake
+++ b/src/tint/lang/spirv/reader/parser/BUILD.cmake
@@ -115,6 +115,7 @@
lang/spirv/reader/parser/memory_test.cc
lang/spirv/reader/parser/misc_test.cc
lang/spirv/reader/parser/name_test.cc
+ lang/spirv/reader/parser/phi_test.cc
lang/spirv/reader/parser/struct_test.cc
lang/spirv/reader/parser/unary_test.cc
lang/spirv/reader/parser/var_test.cc
diff --git a/src/tint/lang/spirv/reader/parser/BUILD.gn b/src/tint/lang/spirv/reader/parser/BUILD.gn
index a5e2017..949d77a 100644
--- a/src/tint/lang/spirv/reader/parser/BUILD.gn
+++ b/src/tint/lang/spirv/reader/parser/BUILD.gn
@@ -112,6 +112,7 @@
"memory_test.cc",
"misc_test.cc",
"name_test.cc",
+ "phi_test.cc",
"struct_test.cc",
"unary_test.cc",
"var_test.cc",
diff --git a/src/tint/lang/spirv/reader/parser/helper_test.h b/src/tint/lang/spirv/reader/parser/helper_test.h
index 5a0959d..297ef85 100644
--- a/src/tint/lang/spirv/reader/parser/helper_test.h
+++ b/src/tint/lang/spirv/reader/parser/helper_test.h
@@ -73,11 +73,12 @@
// Validate the IR module against the capabilities supported by the SPIR-V dialect.
auto validated =
- core::ir::Validate(parsed.Get(), core::ir::Capabilities{
- core::ir::Capability::kAllowMultipleEntryPoints,
- core::ir::Capability::kAllowOverrides,
- core::ir::Capability::kAllowVectorElementPointer,
- });
+ ValidateAndDumpIfNeeded(parsed.Get(), "spirv.test",
+ core::ir::Capabilities{
+ core::ir::Capability::kAllowMultipleEntryPoints,
+ core::ir::Capability::kAllowOverrides,
+ core::ir::Capability::kAllowVectorElementPointer,
+ });
if (validated != Success) {
return validated.Failure();
}
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index d718701..84a3a22 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -67,6 +67,18 @@
namespace {
+// Stores information for operands which need to be calculated after a block is complete. Because
+// a phi can store values which come after it, we can't calculate the value when the `OpPhi` is
+// seen.
+struct ReplacementValue {
+ // The terminator instruction the operand belongs to
+ core::ir::Terminator* terminator;
+ // The operand index in `terminator` to replace
+ uint32_t idx;
+ // The SPIR-V value id to create the value from
+ uint32_t value_id;
+};
+
/// The SPIR-V environment that we validate against.
constexpr auto kTargetEnv = SPV_ENV_VULKAN_1_1;
@@ -716,53 +728,127 @@
return functions_.GetOrAdd(id, [&] { return b_.Function(ty_.void_()); });
}
+ // Passes a value up through control flow to make it visible in an outer scope. Because SPIR-V
+ // allows an id to be referenced as long as it's dominated, you can access a variable which is
+ // defined inside an if branch for example. In order for that to be accessed in IR, we have to
+ // propagate the value as a return of the control instruction (like the if).
+ //
+ // e.g. if the IR is similar to the following:
+ // ```
+ // if (b) {
+ // %a:i32 = let 4;
+ // exit_if
+ // }
+ // %c:i32 = %a + %a;
+ // ```
+ //
+ // We propagate to something like:
+ // ```
+ // %d:i32 = if (b) {
+ // %a:i32 = let 4; // The spir-v ID refers to %a at this point
+ // exit_if %a
+ // }
+ // %c:i32 = %d + %d // The spir-v ID will now refer to %d instead of %a
+ // ```
+ //
+ // We can end up propagating up through multiple levels, so we can end up with something like:
+ // ```
+ // %k:i32 = if (true) {
+ // %l:i32 = if (false) {
+ // %m:i32 = if (true) {
+ // %n:i32 = switch 4 {
+ // default: {
+ // %o:i32 = loop {
+ // %a:i32 = let 4;
+ // exit_loop %a
+ // }
+ // exit_switch %o
+ // }
+ // }
+ // exit_if %n
+ // }
+ // exit_if %m
+ // }
+ // exit_if %l
+ // }
+ // %b:i32 = %k + %k
+ // ```
+ //
+ // @param id the spir-v ID to propagate up
+ // @param src the source value being propagated
core::ir::Value* Propagate(uint32_t id, core::ir::Value* src) {
- auto* src_res = src->As<core::ir::InstructionResult>();
- TINT_ASSERT(src_res);
+ // Function params are always in scope so we should never need to propagate.
+ if (src->Is<core::ir::FunctionParam>()) {
+ return src;
+ }
- auto* blk = src_res->Instruction()->Block();
+ auto* blk = tint::Switch(
+ src, //
+ [&](core::ir::BlockParam* bp) { return bp->Block(); },
+ [&](core::ir::InstructionResult* res) { return res->Instruction()->Block(); },
+ TINT_ICE_ON_NO_MATCH);
+
+ // Walk up the set of control instructions from the current `blk`. We'll update the `src`
+ // instruction as the new result which is to be used for the given SPIR-V `id`. At each
+ // control instruction we'll add the current `src` as a result of each exit from the control
+ // instruction, making a new result which is available in the parent scope.
while (blk) {
if (InBlock(blk)) {
break;
}
- TINT_ASSERT(blk->Terminator());
+ core::ir::ControlInstruction* ctrl = nullptr;
+ if (auto* mb = blk->As<core::ir::MultiInBlock>()) {
+ ctrl = mb->Parent()->As<core::ir::ControlInstruction>();
+ TINT_ASSERT(ctrl);
- core::ir::ControlInstruction* ctrl = tint::Switch(
- blk->Terminator(), //
- [&](core::ir::ExitIf* ei) { return ei->If(); },
- [&](core::ir::ExitSwitch* es) { return es->Switch(); },
- [&](core::ir::ExitLoop* el) { return el->Loop(); },
- [&](core::ir::Continue* cont) {
- // The propagation is going through a `continue`. This means
- // this is the only path to the continuing block, but it also
- // means we're current in the continuing block. We can't do
- // normal propagation here, we have to pass a block param
- // instead.
+ for (auto exit : ctrl->Exits()) {
+ tint::Switch(
+ exit.Value(), //
+ [&](core::ir::ExitLoop* el) { el->PushOperand(src); },
+ [&](core::ir::BreakIf* bi) { bi->PushOperand(src); }, //
+ TINT_ICE_ON_NO_MATCH);
+ }
+ } else {
+ TINT_ASSERT(blk->Terminator());
- auto* param = b_.BlockParam(src->Type());
+ ctrl = tint::Switch(
+ blk->Terminator(), //
+ [&](core::ir::ExitIf* ei) { return ei->If(); },
+ [&](core::ir::ExitSwitch* es) { return es->Switch(); },
+ [&](core::ir::ExitLoop* el) { return el->Loop(); },
+ [&](core::ir::NextIteration* ni) { return ni->Loop(); },
+ [&](core::ir::Continue* cont) {
+ // The propagation is going through a `continue`. This means
+ // this is the only path to the continuing block, but it also
+ // means we're current in the continuing block. We can't do
+ // normal propagation here, we have to pass a block param
+ // instead.
- // We're in the continuing block, so make the block param available in the
- // scope.
- id_stack_.back().insert(id);
+ auto* param = b_.BlockParam(src->Type());
- auto* loop = cont->Loop();
- loop->Continuing()->AddParam(param);
+ // We're in the continuing block, so make the block param available in the
+ // scope.
+ id_stack_.back().insert(id);
- cont->PushOperand(src);
+ auto* loop = cont->Loop();
+ loop->Continuing()->AddParam(param);
- // Set `src` as the `param` so it's returned as the new value
- src = param;
- return nullptr;
- }, //
- [&](core::ir::Unreachable*) { return blk->Parent(); }, //
- TINT_ICE_ON_NO_MATCH);
- if (!ctrl) {
- break;
- }
+ cont->PushOperand(src);
- for (auto& exit : ctrl->Exits()) {
- exit->PushOperand(src);
+ // Set `src` as the `param` so it's returned as the new value
+ src = param;
+ return nullptr;
+ }, //
+ TINT_ICE_ON_NO_MATCH);
+
+ if (!ctrl) {
+ break;
+ }
+
+ for (auto& exit : ctrl->Exits()) {
+ exit->PushOperand(src);
+ }
}
// Add a new result to the control instruction
@@ -786,26 +872,16 @@
return false;
}
+ /// Attempts to retrieve the current Tint IR value for `id`. This ignores scoping for the
+ /// variable, if it exists it's returned (or if it's constant it's created). The value will not
+ /// propagate up through control instructions.
+ ///
/// @param id a SPIR-V result ID
/// @returns a Tint value object
- core::ir::Value* Value(uint32_t id) {
+ core::ir::Value* ValueNoPropagate(uint32_t id) {
auto v = values_.Get(id);
if (v) {
- if (!(*v)->Is<core::ir::InstructionResult>()) {
- return *v;
- }
- if (IdIsInScope(id)) {
- return *v;
- }
-
- // The Value is not in scope, so we need to find the originating Value, and then
- // propagate it up through the control instructions. That will then change the
- // `Value` which is returned so, set it into the values map as the new "Value" and
- // return it.
-
- auto* new_v = Propagate(id, *v);
- values_.Replace(id, new_v);
- return new_v;
+ return *v;
}
if (auto* c = spirv_context_->get_constant_mgr()->FindDeclaredConstant(id)) {
@@ -813,9 +889,30 @@
values_.Add(id, val);
return val;
}
+
TINT_UNREACHABLE() << "missing value for result ID " << id;
}
+ /// Attempts to retrieve the current Tint IR value for `id`. If the value exists and is not in
+ /// scope it will propagate the value up through the control instructions.
+ ///
+ /// @param id a SPIR-V result ID
+ /// @returns a Tint value object
+ core::ir::Value* Value(uint32_t id) {
+ auto v = ValueNoPropagate(id);
+ TINT_ASSERT(v);
+
+ if (v->Is<core::ir::Constant>() || IdIsInScope(id)) {
+ return v;
+ }
+
+ auto* new_v = Propagate(id, v);
+ values_.Replace(id, new_v);
+ return new_v;
+ }
+
+ /// Creates the Tint IR constant for the SPIR-V `constant` value.
+ ///
/// @param constant a SPIR-V constant object
/// @returns a Tint constant value
const core::constant::Value* Constant(const spvtools::opt::analysis::Constant* constant) {
@@ -1057,7 +1154,7 @@
// A block parent is a container for a scope, like a `{}`d section in code. It controls the
// block addition to the current blocks and the ID stack entry for the block.
- void EmitBlockParent(core::ir::Block* dst, const spvtools::opt::BasicBlock& src) {
+ void EmitBlockParent(core::ir::Block* dst, spvtools::opt::BasicBlock& src) {
TINT_ASSERT(!InBlock(dst));
id_stack_.emplace_back();
@@ -1072,9 +1169,11 @@
/// Emit the contents of SPIR-V block @p src into Tint IR block @p dst.
/// @param dst the Tint IR block to append to
/// @param src the SPIR-V block to emit
- void EmitBlock(core::ir::Block* dst, const spvtools::opt::BasicBlock& src) {
+ void EmitBlock(core::ir::Block* dst, spvtools::opt::BasicBlock& src) {
TINT_SCOPED_ASSIGNMENT(current_block_, dst);
+ values_to_replace_.push_back({});
+
auto* loop_merge_inst = src.GetLoopMergeInst();
// This is a loop merge block, so we need to treat it as a Loop.
if (loop_merge_inst) {
@@ -1094,6 +1193,8 @@
current_block_ = loop->Body();
}
+ spirv_id_to_block_.insert({src.id(), current_block_});
+
for (auto& inst : src) {
switch (inst.opcode()) {
case spv::Op::OpNop:
@@ -1498,6 +1599,9 @@
case spv::Op::OpImageWrite:
EmitImageWrite(inst);
break;
+ case spv::Op::OpPhi:
+ EmitPhi(inst);
+ break;
default:
TINT_UNIMPLEMENTED()
<< "unhandled SPIR-V instruction: " << static_cast<uint32_t>(inst.opcode());
@@ -1510,31 +1614,312 @@
auto* loop = StopWalkingAt(src.id())->As<core::ir::Loop>();
TINT_ASSERT(loop);
- auto continue_id = loop_merge_inst->GetSingleWordInOperand(1);
- if (continue_id != src.id()) {
- const auto& bb_continue = current_spirv_function_->FindBlock(continue_id);
-
- // Emit the continuing block.
- EmitBlockParent(loop->Continuing(), *bb_continue);
- }
-
// Add the body terminator if necessary
if (!loop->Body()->Terminator()) {
loop->Body()->Append(b_.Unreachable());
}
+
+ // Push id stack entry for the continuing block. We don't use EmitBlockParent to do this
+ // because we need the scope to exist until after we process any `continue_blk_phis_`.
+ id_stack_.emplace_back();
+
+ auto continue_id = loop_merge_inst->GetSingleWordInOperand(1);
+ if (continue_id != src.id()) {
+ const auto& bb_continue = current_spirv_function_->FindBlock(continue_id);
+
+ current_blocks_.insert(loop->Continuing());
+ // Emit the continuing block.
+ EmitBlock(loop->Continuing(), *bb_continue);
+
+ current_blocks_.erase(loop->Continuing());
+ }
+
if (!loop->Continuing()->Terminator()) {
loop->Continuing()->Append(b_.NextIteration(loop));
}
+ // If this continue block needs to pass any `phi` instructions back to
+ // the main loop body.
+ //
+ // We have to do this here because we need to have emitted the loop
+ // body before we can get the values used in the continue block.
+ auto phis = continue_blk_phis_.find(continue_id);
+ if (phis != continue_blk_phis_.end()) {
+ for (auto value_id : phis->second) {
+ auto* value = Value(value_id);
+
+ tint::Switch(
+ loop->Continuing()->Terminator(), //
+ [&](core::ir::NextIteration* ni) { ni->PushOperand(value); },
+ [&](core::ir::BreakIf* bi) {
+ // TODO(dsinclair): Need to change the break-if insertion of there
+ // happens to be exit values, but those are rare, so leave this for when
+ // we have test case.
+ TINT_ASSERT(bi->ExitValues().IsEmpty());
+
+ auto len = bi->NextIterValues().Length();
+ bi->PushOperand(value);
+ bi->SetNumNextIterValues(len + 1);
+ },
+ TINT_ICE_ON_NO_MATCH);
+ }
+ }
+
+ id_stack_.pop_back();
+ }
+
+ // For any `OpPhi` values we saw, insert their `Value` now. We do this
+ // at the end of the loop because a phi can refer to instructions
+ // defined after it in the block.
+ auto replace = values_to_replace_.back();
+ for (auto& val : replace) {
+ auto* value = ValueNoPropagate(val.value_id);
+ val.terminator->SetOperand(val.idx, value);
+ }
+ values_to_replace_.pop_back();
+
+ if (loop_merge_inst) {
+ auto* loop = StopWalkingAt(src.id())->As<core::ir::Loop>();
+ TINT_ASSERT(loop);
+
current_blocks_.erase(loop->Body());
id_stack_.pop_back();
+ // If we added phi's to the continuing block, we may have exits from the body which
+ // aren't valid.
+ auto continuing_param_count = loop->Continuing()->Params().Length();
+ if (continuing_param_count > 0) {
+ for (auto incoming : loop->Continuing()->InboundSiblingBranches()) {
+ TINT_ASSERT(incoming->Is<core::ir::Continue>());
+
+ // Check if the block this instruction exists in has default phi result that we
+ // can append.
+ auto inst_to_blk_iter = inst_to_spirv_block_.find(incoming);
+ if (inst_to_blk_iter != inst_to_spirv_block_.end()) {
+ uint32_t spirv_blk = inst_to_blk_iter->second;
+ auto phi_values_from_loop_header = block_phi_values_[spirv_blk];
+ // If there were phi values, push them to this instruction
+ for (auto value_id : phi_values_from_loop_header) {
+ auto* value = Value(value_id);
+ incoming->PushOperand(value);
+ }
+ }
+ }
+ }
+
+ // Emit the merge block
auto merge_id = loop_merge_inst->GetSingleWordInOperand(0);
const auto& merge_bb = current_spirv_function_->FindBlock(merge_id);
EmitBlock(dst, *merge_bb);
}
}
+ struct IfBranchValue {
+ core::ir::Value* value;
+ core::ir::If* if_;
+ };
+
+ void EmitPhi(spvtools::opt::Instruction& inst) {
+ auto num_ops = inst.NumInOperands();
+
+ // If there are only 2 arguments, that means we came directly from a block, so just emit the
+ // value directly.
+ if (num_ops == 2) {
+ AddValue(inst.result_id(), Value(inst.GetSingleWordInOperand(0)));
+ return;
+ }
+
+ std::unordered_map<core::ir::ControlInstruction*, const core::type::Type*>
+ ctrl_inst_result_types;
+ std::unordered_map<core::ir::MultiInBlock*, const core::type::Type*> blk_to_param_types;
+
+ std::optional<IfBranchValue> if_to_update_branch;
+
+ auto add_ctrl_inst = [&](core::ir::ControlInstruction* ctrl, const core::type::Type* type) {
+ auto iter = ctrl_inst_result_types.find(ctrl);
+ if (iter != ctrl_inst_result_types.end()) {
+ TINT_ASSERT(iter->second == type);
+ return;
+ }
+ ctrl_inst_result_types.insert({ctrl, type});
+ };
+
+ auto* type = Type(inst.type_id());
+ auto add_blk_inst = [&](core::ir::MultiInBlock* blk) {
+ auto iter = blk_to_param_types.find(blk);
+ if (iter != blk_to_param_types.end()) {
+ TINT_ASSERT(iter->second == type);
+ return;
+ }
+ blk_to_param_types.insert({blk, type});
+ };
+
+ auto* phi_spirv_block = spirv_context_->get_instr_block(&inst);
+ auto* phi_loop_merge_inst = phi_spirv_block->GetLoopMergeInst();
+
+ for (uint32_t i = 0; i < num_ops; i += 2) {
+ auto value_id = inst.GetSingleWordInOperand(i);
+ auto blk_id = inst.GetSingleWordInOperand(i + 1);
+
+ // Store this value away as a default phi value for this loop header.
+ block_phi_values_[blk_id].push_back(value_id);
+
+ auto value_blk_iter = spirv_id_to_block_.find(blk_id);
+
+ // The referenced block hasn't been emitted yet (continue blocks have this
+ // behaviour). So, store the fact that it needs to return a given value away for
+ // when we do emit the block.
+ if (value_blk_iter == spirv_id_to_block_.end()) {
+ auto continue_id = phi_loop_merge_inst->GetSingleWordInOperand(1);
+
+ // Note, we push it to the `continue_id` as the block and not
+ // `blk_id` so that we can emit them into the continuing block as
+ // a group.
+ continue_blk_phis_[continue_id].push_back(value_id);
+
+ // Add the phi to the current block set of input parameters
+ auto* mb = current_block_->As<core::ir::MultiInBlock>();
+ TINT_ASSERT(mb);
+ add_blk_inst(mb);
+ continue;
+ }
+
+ core::ir::Terminator* term = nullptr;
+
+ // The `OpPhi` is part of a loop header block, treat it special as we need to insert
+ // things into the phi's loop initializer/body/continuing block as needed.
+ if (phi_loop_merge_inst) {
+ auto* loop = StopWalkingAt(phi_spirv_block->id())->As<core::ir::Loop>();
+ TINT_ASSERT(loop);
+
+ // A phi from an explicit continue block is handled above as we haven't emitted the
+ // continue block so we wouldn't find it in the `spirv_id_to_block` list.
+
+ // If this loop header is also the continue block
+ if (blk_id == phi_spirv_block->id()) {
+ if (loop->Continuing()->IsEmpty()) {
+ b_.Append(loop->Continuing(), [&] { term = b_.NextIteration(loop); });
+ add_blk_inst(loop->Body());
+ } else {
+ // With multiple phis we my have already created the continuing
+ // block, so just get the terminator.
+ term = loop->Continuing()->Terminator();
+ TINT_ASSERT(term->Is<core::ir::NextIteration>());
+ }
+ } else {
+ // We know this isn't the continue as it hasn't emitted yet, so this has to be
+ // coming from the calling block. So, we need to add this item into the
+ // initializer `NextIteration` and as a parameter to the body.
+ if (loop->Initializer()->IsEmpty()) {
+ b_.Append(loop->Initializer(), [&] { term = b_.NextIteration(loop); });
+ add_blk_inst(loop->Body());
+ } else {
+ term = loop->Initializer()->Terminator();
+ TINT_ASSERT(term->Is<core::ir::NextIteration>());
+ }
+ }
+
+ } else {
+ auto* value_ir_blk = value_blk_iter->second;
+
+ // We know the phi isn't part of a loop. That means, all of the blocks making up the
+ // phi are known. The one trick is that an `if` may only have a single block (the
+ // true or false). In that case, we have to push the value into the other block
+ // as it has to return something.
+ //
+ // For a `Switch` we will have all the cases already so we can just get the
+ // terminator for the block.
+
+ if (!value_ir_blk->Terminator()) {
+ auto* if_ = value_ir_blk->Back()->As<core::ir::If>();
+ TINT_ASSERT(if_);
+ // No block terminator means the block that the `phi` is referencing
+ // isn't finished (in IR-land). This can only happen with an `if`
+ // instruction where you've branched directly to the `phi` as either the
+ // `then` or `else` clause.
+
+ TINT_ASSERT(!if_to_update_branch.has_value());
+
+ auto* value = ValueNoPropagate(value_id);
+ if_to_update_branch = IfBranchValue{
+ .value = value,
+ .if_ = if_,
+ };
+
+ continue;
+ }
+ term = value_ir_blk->Terminator();
+ }
+
+ // If we can't get to this part of the control flow, ignore the phi
+ if (term->Is<core::ir::Unreachable>()) {
+ continue;
+ }
+
+ // Push a placeholder for the operand value at this point. We'll
+ // store away the terminator/index pair along with the required
+ // value and then fill it in at the end of the block emission.
+ auto operand_idx = term->PushOperand(nullptr);
+ values_to_replace_.back().push_back(ReplacementValue{
+ .terminator = term,
+ .idx = operand_idx,
+ .value_id = value_id,
+ });
+
+ // For each incoming block to the phi, store either the control
+ // instruction to be updated, or the block to be updated and the
+ // type of result to return.
+ tint::Switch(
+ term, //
+ [&](core::ir::Exit* exit) { add_ctrl_inst(exit->ControlInstruction(), type); },
+ [&](core::ir::BreakIf* bi) { add_ctrl_inst(bi->Loop(), type); },
+ [&](core::ir::Continue* cont) { add_blk_inst(cont->Loop()->Continuing()); },
+ [&](core::ir::NextIteration* ni) { add_blk_inst(ni->Loop()->Body()); },
+ TINT_ICE_ON_NO_MATCH);
+ }
+
+ // We need to update one of the two `if` branches with the return value.
+ // Find the one where the terminator has less operands and update that
+ // one.
+ if (if_to_update_branch.has_value()) {
+ auto* value = if_to_update_branch->value;
+ auto* if_ = if_to_update_branch->if_;
+
+ core::ir::Terminator* term = nullptr;
+
+ if (if_->True()->Terminator()->Operands().Length() <
+ if_->False()->Terminator()->Operands().Length()) {
+ term = if_->True()->Terminator();
+ } else {
+ term = if_->False()->Terminator();
+ }
+
+ term->PushOperand(value);
+ add_ctrl_inst(if_, value->Type());
+ }
+
+ // Update control instruction results to contain the new type.
+ for (auto info : ctrl_inst_result_types) {
+ auto* ctrl = info.first;
+ auto* res_type = info.second;
+ auto* res = b_.InstructionResult(res_type);
+ ctrl->AddResult(res);
+ values_.Replace(inst.result_id(), res);
+ }
+
+ // Update block params to contain the new type.
+ for (auto info : blk_to_param_types) {
+ auto* blk = info.first;
+ auto* param_type = info.second;
+
+ auto* p = b_.BlockParam(param_type);
+ blk->AddParam(p);
+
+ TINT_ASSERT(blk == current_block_);
+ AddValue(inst.result_id(), p);
+ }
+ }
+
void EmitSampledImage(const spvtools::opt::Instruction& inst) {
auto* tex = Value(inst.GetSingleWordInOperand(0));
Emit(b_.CallExplicit<spirv::ir::BuiltinCall>(Type(inst.type_id()),
@@ -1810,6 +2195,7 @@
if (auto* ctrl_inst = StopWalkingAt(dest_id)) {
if (auto* loop = ctrl_inst->As<core::ir::Loop>()) {
// Going to the merge in a loop body has to be a break regardless of nesting level.
+
if (InBlock(loop->Body()) && !InBlock(loop->Continuing())) {
EmitWithoutResult(b_.Exit(ctrl_inst));
}
@@ -1907,27 +2293,30 @@
return ctrl;
}
- void EmitBranchStopBlock(core::ir::ControlInstruction* ctrl,
- core::ir::If* if_,
- core::ir::Block* blk,
- uint32_t target) {
+ core::ir::Instruction* EmitBranchStopBlock(core::ir::ControlInstruction* ctrl,
+ core::ir::If* if_,
+ core::ir::Block* blk,
+ uint32_t target) {
if (auto* loop = ContinueTarget(target)) {
- blk->Append(b_.Continue(loop));
- } else {
- auto iter = merge_to_premerge_.find(target);
- if (iter != merge_to_premerge_.end()) {
- // Branch to a merge block, but skipping over an expected premerge block
- // so we need a guard.
- if (!iter->second.condition) {
- b_.InsertBefore(iter->second.parent, [&] {
- iter->second.condition = b_.Var("execute_premerge", true);
- });
- }
- b_.Append(blk, [&] { b_.Store(iter->second.condition, false); });
- }
-
- blk->Append(b_.Exit(ExitFor(ctrl, if_)));
+ auto* cont = b_.Continue(loop);
+ blk->Append(cont);
+ return cont;
}
+
+ auto iter = merge_to_premerge_.find(target);
+ if (iter != merge_to_premerge_.end()) {
+ // Branch to a merge block, but skipping over an expected premerge block
+ // so we need a guard.
+ if (!iter->second.condition) {
+ b_.InsertBefore(iter->second.parent,
+ [&] { iter->second.condition = b_.Var("execute_premerge", true); });
+ }
+ b_.Append(blk, [&] { b_.Store(iter->second.condition, false); });
+ }
+
+ auto* exit = b_.Exit(ExitFor(ctrl, if_));
+ blk->Append(exit);
+ return exit;
}
bool ProcessBranchAsLoopHeader(core::ir::Value* cond, uint32_t true_id, uint32_t false_id) {
@@ -2062,7 +2451,8 @@
}
if (auto* ctrl = StopWalkingAt(true_id)) {
- EmitBranchStopBlock(ctrl, if_, if_->True(), true_id);
+ auto* new_inst = EmitBranchStopBlock(ctrl, if_, if_->True(), true_id);
+ inst_to_spirv_block_[new_inst] = bb.id();
} else {
EmitIfBranch(true_id, if_, if_->True());
}
@@ -2072,7 +2462,8 @@
if (false_id == true_id) {
if_->False()->Append(b_.Unreachable());
} else if (auto* ctrl = StopWalkingAt(false_id)) {
- EmitBranchStopBlock(ctrl, if_, if_->False(), false_id);
+ auto* new_inst = EmitBranchStopBlock(ctrl, if_, if_->False(), false_id);
+ inst_to_spirv_block_[new_inst] = bb.id();
} else {
EmitIfBranch(false_id, if_, if_->False());
}
@@ -2891,7 +3282,7 @@
std::unordered_map<uint32_t, core::ir::ControlInstruction*> walk_stop_blocks_;
// Map of continue target ID to the controlling IR loop.
std::unordered_map<uint32_t, core::ir::Loop*> continue_targets_;
- // Map of continue target ID to the controlling IR loop.
+ // Map of header target ID to the controlling IR loop.
std::unordered_map<uint32_t, core::ir::Loop*> loop_headers_;
struct PremergeInfo {
@@ -2902,11 +3293,31 @@
std::unordered_map<uint32_t, PremergeInfo> merge_to_premerge_;
std::unordered_set<core::ir::Block*> current_blocks_;
+
+ // For each block, we keep a set of SPIR-V `id`s which are known in that scope.
std::vector<std::unordered_set<uint32_t>> id_stack_;
// If we're in a switch, is populated with the IDs of the blocks for each of the switch
// selectors. This lets us watch for fallthrough when emitting branch instructions.
std::vector<std::unordered_set<uint32_t>> current_switch_blocks_;
+
+ /// Maps from a spirv-v block id to the corresponding block in the IR
+ std::unordered_map<uint32_t, core::ir::Block*> spirv_id_to_block_;
+
+ // Map of continue block id to the phi types which need to be returned by
+ // the continue target
+ std::unordered_map<uint32_t, std::vector<uint32_t>> continue_blk_phis_;
+
+ // A stack of values which need to be replaced as we finish processing a
+ // block. Used to store `phi` information so we can retrieve values which
+ // are defined after the `OpPhi` instruction.
+ std::vector<std::vector<ReplacementValue>> values_to_replace_;
+
+ // A map of loop header to phi values returned by that loop header
+ std::unordered_map<uint32_t, std::vector<uint32_t>> block_phi_values_;
+
+ // Map of certain instructions back to their originating spirv block
+ std::unordered_map<core::ir::Instruction*, uint32_t> inst_to_spirv_block_;
};
} // namespace
diff --git a/src/tint/lang/spirv/reader/parser/phi_test.cc b/src/tint/lang/spirv/reader/parser/phi_test.cc
new file mode 100644
index 0000000..22a6610
--- /dev/null
+++ b/src/tint/lang/spirv/reader/parser/phi_test.cc
@@ -0,0 +1,1538 @@
+// Copyright 2025 The Dawn & Tint Authors
+//
+// Redistribution and use in source and binary forms, with or without
+// modification, are permitted provided that the following conditions are met:
+//
+// 1. Redistributions of source code must retain the above copyright notice, this
+// list of conditions and the following disclaimer.
+//
+// 2. Redistributions in binary form must reproduce the above copyright notice,
+// this list of conditions and the following disclaimer in the documentation
+// and/or other materials provided with the distribution.
+//
+// 3. Neither the name of the copyright holder nor the names of its
+// contributors may be used to endorse or promote products derived from
+// this software without specific prior written permission.
+//
+// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+// DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE
+// FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL
+// DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR
+// SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER
+// CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY,
+// OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE
+// OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+
+#include "src/tint/lang/spirv/reader/parser/helper_test.h"
+
+namespace tint::spirv::reader {
+namespace {
+
+TEST_F(SpirvParserTest, Phi_FromBlock) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %int = OpTypeInt 32 1
+ %true = OpConstantTrue %bool
+ %int_2 = OpConstant %int 2
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ OpBranch %1
+ %1 = OpLabel
+ %2 = OpPhi %int %int_2 %main_start
+ %3 = OpIAdd %int %2 %2
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = spirv.add<i32> 2i, 2i
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Phi_If_Undef) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %int = OpTypeInt 32 1
+ %true = OpConstantTrue %bool
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ OpBranch %1
+ %1 = OpLabel
+ OpSelectionMerge %5 None
+ OpBranchConditional %true %3 %4
+ %3 = OpLabel
+ %6 = OpUndef %int
+ OpBranch %5
+ %4 = OpLabel
+ %7 = OpUndef %int
+ OpBranch %5
+ %5 = OpLabel
+ %8 = OpPhi %int %6 %3 %7 %4
+ %9 = OpIAdd %int %8 %8
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = if true [t: $B2, f: $B3] { # if_1
+ $B2: { # true
+ exit_if 0i # if_1
+ }
+ $B3: { # false
+ exit_if 0i # if_1
+ }
+ }
+ %3:i32 = spirv.add<i32> %2, %2
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Phi_If_ThenAndElse) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %int = OpTypeInt 32 1
+ %true = OpConstantTrue %bool
+ %int_1 = OpConstant %int 1
+ %int_2 = OpConstant %int 2
+ %int_3 = OpConstant %int 3
+ %int_4 = OpConstant %int 4
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ OpBranch %1
+ %1 = OpLabel
+ OpSelectionMerge %5 None
+ OpBranchConditional %true %3 %4
+ %3 = OpLabel
+ %6 = OpIAdd %int %int_1 %int_2
+ OpBranch %5
+ %4 = OpLabel
+ %7 = OpIAdd %int %int_3 %int_4
+ OpBranch %5
+ %5 = OpLabel
+ %8 = OpPhi %int %6 %3 %7 %4
+ %9 = OpIAdd %int %8 %8
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = if true [t: $B2, f: $B3] { # if_1
+ $B2: { # true
+ %3:i32 = spirv.add<i32> 1i, 2i
+ exit_if %3 # if_1
+ }
+ $B3: { # false
+ %4:i32 = spirv.add<i32> 3i, 4i
+ exit_if %4 # if_1
+ }
+ }
+ %5:i32 = spirv.add<i32> %2, %2
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Phi_If_ThenNoElse) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %int = OpTypeInt 32 1
+ %true = OpConstantTrue %bool
+ %int_1 = OpConstant %int 1
+ %int_2 = OpConstant %int 2
+ %int_3 = OpConstant %int 3
+ %int_4 = OpConstant %int 4
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ OpBranch %1
+ %1 = OpLabel
+ OpSelectionMerge %5 None
+ OpBranchConditional %true %3 %5
+ %3 = OpLabel
+ %6 = OpIAdd %int %int_1 %int_2
+ OpBranch %5
+ %5 = OpLabel
+ %8 = OpPhi %int %int_2 %1 %6 %3
+ %9 = OpIAdd %int %8 %8
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = if true [t: $B2, f: $B3] { # if_1
+ $B2: { # true
+ %3:i32 = spirv.add<i32> 1i, 2i
+ exit_if %3 # if_1
+ }
+ $B3: { # false
+ exit_if 2i # if_1
+ }
+ }
+ %4:i32 = spirv.add<i32> %2, %2
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Phi_If_NoThenElse) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %int = OpTypeInt 32 1
+ %true = OpConstantTrue %bool
+ %int_1 = OpConstant %int 1
+ %int_2 = OpConstant %int 2
+ %int_3 = OpConstant %int 3
+ %int_4 = OpConstant %int 4
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ OpBranch %1
+ %1 = OpLabel
+ OpSelectionMerge %5 None
+ OpBranchConditional %true %5 %3
+ %3 = OpLabel
+ %6 = OpIAdd %int %int_1 %int_2
+ OpBranch %5
+ %5 = OpLabel
+ %8 = OpPhi %int %6 %3 %int_2 %1
+ %9 = OpIAdd %int %8 %8
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = if true [t: $B2, f: $B3] { # if_1
+ $B2: { # true
+ exit_if 2i # if_1
+ }
+ $B3: { # false
+ %3:i32 = spirv.add<i32> 1i, 2i
+ exit_if %3 # if_1
+ }
+ }
+ %4:i32 = spirv.add<i32> %2, %2
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Phi_Switch) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %3 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %int_0 = OpConstant %int 0
+ %int_1 = OpConstant %int 1
+ %int_2 = OpConstant %int 2
+ %main = OpFunction %void None %3
+ %4 = OpLabel
+ OpSelectionMerge %13 None
+ OpSwitch %int_0 %10 0 %11 1 %12
+ %10 = OpLabel
+ OpBranch %13
+ %11 = OpLabel
+ OpBranch %13
+ %12 = OpLabel
+ OpBranch %13
+ %13 = OpLabel
+ %14 = OpPhi %int %int_0 %10 %int_1 %11 %int_2 %12
+ %15 = OpIAdd %int %14 %14
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = switch 0i [c: (default, $B2), c: (0i, $B3), c: (1i, $B4)] { # switch_1
+ $B2: { # case
+ exit_switch 0i # switch_1
+ }
+ $B3: { # case
+ exit_switch 1i # switch_1
+ }
+ $B4: { # case
+ exit_switch 2i # switch_1
+ }
+ }
+ %3:i32 = spirv.add<i32> %2, %2
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Phi_Loop_ContinueIsHeader) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %true = OpConstantTrue %bool
+ %false = OpConstantFalse %bool
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %10 = OpLabel
+ OpBranch %20
+ %20 = OpLabel
+ %101 = OpPhi %bool %true %10 %false %20
+ OpLoopMerge %99 %20 None
+ OpBranchConditional %101 %99 %20
+ %99 = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ loop [i: $B2, b: $B3, c: $B4] { # loop_1
+ $B2: { # initializer
+ next_iteration true # -> $B3
+ }
+ $B3 (%2:bool): { # body
+ if %2 [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_loop # loop_1
+ }
+ $B6: { # false
+ continue # -> $B4
+ }
+ }
+ unreachable
+ }
+ $B4: { # continuing
+ next_iteration false # -> $B3
+ }
+ }
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Phi_Loop_WithContinue) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %true = OpConstantTrue %bool
+ %false = OpConstantFalse %bool
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %10 = OpLabel
+ OpBranch %20
+ %20 = OpLabel
+ %101 = OpPhi %bool %true %10 %false %30
+ OpLoopMerge %99 %30 None
+ OpBranchConditional %101 %99 %25
+ %25 = OpLabel
+ OpBranch %30
+ %30 = OpLabel
+ %102 = OpCopyObject %bool %101
+ OpBranch %20
+ %99 = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ loop [i: $B2, b: $B3, c: $B4] { # loop_1
+ $B2: { # initializer
+ next_iteration true # -> $B3
+ }
+ $B3 (%2:bool): { # body
+ if %2 [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_loop # loop_1
+ }
+ $B6: { # false
+ continue # -> $B4
+ }
+ }
+ unreachable
+ }
+ $B4: { # continuing
+ %3:bool = let %2
+ next_iteration false # -> $B3
+ }
+ }
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Phi_Loop_WithContinue_PhiInContinue) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %true = OpConstantTrue %bool
+ %false = OpConstantFalse %bool
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %10 = OpLabel
+ OpBranch %20
+ %20 = OpLabel
+ OpLoopMerge %99 %30 None
+ OpBranchConditional %true %24 %30
+ %24 = OpLabel
+ OpBranch %30
+ %30 = OpLabel
+ %101 = OpPhi %bool %true %24 %false %20
+ %102 = OpCopyObject %bool %101
+ OpBranchConditional %true %99 %20
+ %99 = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ if true [t: $B4, f: $B5] { # if_1
+ $B4: { # true
+ continue true # -> $B3
+ }
+ $B5: { # false
+ continue false # -> $B3
+ }
+ }
+ unreachable
+ }
+ $B3 (%2:bool): { # continuing
+ %3:bool = let %2
+ break_if true # -> [t: exit_loop loop_1, f: $B2]
+ }
+ }
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Phi_Loop_WithMultiblockContinue) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %true = OpConstantTrue %bool
+ %false = OpConstantFalse %bool
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %10 = OpLabel
+ OpBranch %20
+ %20 = OpLabel
+ %101 = OpPhi %bool %true %10 %false %31
+ OpLoopMerge %99 %30 None
+ OpBranchConditional %101 %99 %25
+ %25 = OpLabel
+ OpBranch %30
+ %30 = OpLabel
+ %102 = OpCopyObject %bool %101
+ OpBranch %31
+ %31 = OpLabel
+ OpBranch %20
+ %99 = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ loop [i: $B2, b: $B3, c: $B4] { # loop_1
+ $B2: { # initializer
+ next_iteration true # -> $B3
+ }
+ $B3 (%2:bool): { # body
+ if %2 [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_loop # loop_1
+ }
+ $B6: { # false
+ continue # -> $B4
+ }
+ }
+ unreachable
+ }
+ $B4: { # continuing
+ %3:bool = let %2
+ next_iteration false # -> $B3
+ }
+ }
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Phi_Loop_BranchConditionalBreak) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %int = OpTypeInt 32 1
+ %true = OpConstantTrue %bool
+ %int_1 = OpConstant %int 1
+ %int_2 = OpConstant %int 2
+ %int_3 = OpConstant %int 3
+ %int_4 = OpConstant %int 4
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ OpBranch %1
+ %1 = OpLabel
+ OpLoopMerge %99 %20 None
+ OpBranchConditional %true %50 %20
+ %50 = OpLabel
+ %6 = OpIAdd %int %int_1 %int_2
+ OpBranch %99
+ %20 = OpLabel
+ OpBranch %1
+ %99 = OpLabel
+ %8 = OpPhi %int %6 %50
+ %9 = OpIAdd %int %8 %8
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:i32 = loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ if true [t: $B4, f: $B5] { # if_1
+ $B4: { # true
+ %3:i32 = spirv.add<i32> 1i, 2i
+ exit_loop %3 # loop_1
+ }
+ $B5: { # false
+ continue # -> $B3
+ }
+ }
+ unreachable
+ }
+ $B3: { # continuing
+ next_iteration # -> $B2
+ }
+ }
+ %4:i32 = spirv.add<i32> %2, %2
+ ret
+ }
+}
+)");
+}
+
+// Phis must act as if they are simultaneously assigned. %101 and %102 should exchange values on
+// each iteration, and never have the same value.
+TEST_F(SpirvParserTest, Phi_SimultaneousAssignment) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %101 "default_true"
+ OpName %102 "default_false"
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %true = OpConstantTrue %bool
+ %false = OpConstantFalse %bool
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %10 = OpLabel
+ OpBranch %20
+ %20 = OpLabel
+ %101 = OpPhi %bool %true %10 %102 %20
+ %102 = OpPhi %bool %false %10 %101 %20
+ OpLoopMerge %99 %20 None
+ OpBranchConditional %true %99 %20
+ %99 = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ loop [i: $B2, b: $B3, c: $B4] { # loop_1
+ $B2: { # initializer
+ next_iteration true, false # -> $B3
+ }
+ $B3 (%2:bool, %3:bool): { # body
+ if true [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_loop # loop_1
+ }
+ $B6: { # false
+ continue # -> $B4
+ }
+ }
+ unreachable
+ }
+ $B4: { # continuing
+ next_iteration %3, %2 # -> $B3
+ }
+ }
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Phi_SingleBlockLoopIndex) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %2 "computed"
+ OpName %3 "copied"
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %uint = OpTypeInt 32 0
+ %pty = OpTypePointer Private %uint
+ %1 = OpVariable %pty Private
+ %boolpty = OpTypePointer Private %bool
+ %7 = OpVariable %boolpty Private
+ %8 = OpVariable %boolpty Private
+ %uint_0 = OpConstant %uint 0
+ %uint_1 = OpConstant %uint 1
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ OpBranch %10
+; Use an outer loop to show we put the new variable in the
+; smallest enclosing scope.
+ %10 = OpLabel
+ %101 = OpLoad %bool %7
+ %102 = OpLoad %bool %8
+ %103 = OpIAdd %uint %uint_0 %uint_0
+ OpLoopMerge %99 %89 None
+ OpBranchConditional %101 %99 %20
+ %20 = OpLabel
+ %2 = OpPhi %uint %103 %10 %4 %20 ; gets computed value
+ %3 = OpPhi %uint %uint_1 %10 %3 %20 ; gets itself
+ %4 = OpIAdd %uint %2 %uint_1
+ OpLoopMerge %79 %20 None
+ OpBranchConditional %102 %79 %20
+ %79 = OpLabel
+ OpBranch %89
+ %89 = OpLabel
+ OpBranch %10
+ %99 = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:ptr<private, u32, read_write> = var undef
+ %2:ptr<private, bool, read_write> = var undef
+ %3:ptr<private, bool, read_write> = var undef
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ loop [b: $B3, c: $B4] { # loop_1
+ $B3: { # body
+ %5:bool = load %2
+ %6:bool = load %3
+ %7:u32 = spirv.add<u32> 0u, 0u
+ if %5 [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_loop # loop_1
+ }
+ $B6: { # false
+ loop [i: $B7, b: $B8, c: $B9] { # loop_2
+ $B7: { # initializer
+ next_iteration %7, 1u # -> $B8
+ }
+ $B8 (%8:u32, %9:u32): { # body
+ %10:u32 = spirv.add<u32> %8, 1u
+ if %6 [t: $B10, f: $B11] { # if_2
+ $B10: { # true
+ exit_loop # loop_2
+ }
+ $B11: { # false
+ continue # -> $B9
+ }
+ }
+ unreachable
+ }
+ $B9: { # continuing
+ next_iteration %10, %9 # -> $B8
+ }
+ }
+ continue # -> $B4
+ }
+ }
+ unreachable
+ }
+ $B4: { # continuing
+ next_iteration # -> $B3
+ }
+ }
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Phi_MultiBlockLoopIndex) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %uint = OpTypeInt 32 0
+ %pty = OpTypePointer Private %uint
+ %1 = OpVariable %pty Private
+ %boolpty = OpTypePointer Private %bool
+ %7 = OpVariable %boolpty Private
+ %8 = OpVariable %boolpty Private
+ %uint_0 = OpConstant %uint 0
+ %uint_1 = OpConstant %uint 1
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %5 = OpLabel
+ OpBranch %10
+; Use an outer loop to show we put the new variable in the
+; smallest enclosing scope.
+ %10 = OpLabel
+ %101 = OpLoad %bool %7
+ %102 = OpLoad %bool %8
+ OpLoopMerge %99 %89 None
+ OpBranchConditional %101 %99 %20
+ %20 = OpLabel
+ %2 = OpPhi %uint %uint_0 %10 %4 %30 ; gets computed value
+ %3 = OpPhi %uint %uint_1 %10 %3 %30 ; gets itself
+ OpLoopMerge %79 %30 None
+ OpBranchConditional %102 %79 %30
+ %30 = OpLabel ; continue target for inner loop
+ %4 = OpIAdd %uint %2 %uint_1
+ OpBranch %20
+ %79 = OpLabel ; merge for inner loop
+ OpBranch %89
+ %89 = OpLabel ; continue target for outer loop
+ OpBranch %10
+ %99 = OpLabel ; merge for outer loop
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:ptr<private, u32, read_write> = var undef
+ %2:ptr<private, bool, read_write> = var undef
+ %3:ptr<private, bool, read_write> = var undef
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ loop [b: $B3, c: $B4] { # loop_1
+ $B3: { # body
+ %5:bool = load %2
+ %6:bool = load %3
+ if %5 [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_loop # loop_1
+ }
+ $B6: { # false
+ loop [i: $B7, b: $B8, c: $B9] { # loop_2
+ $B7: { # initializer
+ next_iteration 0u, 1u # -> $B8
+ }
+ $B8 (%7:u32, %8:u32): { # body
+ if %6 [t: $B10, f: $B11] { # if_2
+ $B10: { # true
+ exit_loop # loop_2
+ }
+ $B11: { # false
+ continue # -> $B9
+ }
+ }
+ unreachable
+ }
+ $B9: { # continuing
+ %9:u32 = spirv.add<u32> %7, 1u
+ next_iteration %9, %8 # -> $B8
+ }
+ }
+ continue # -> $B4
+ }
+ }
+ unreachable
+ }
+ $B4: { # continuing
+ next_iteration # -> $B3
+ }
+ }
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Phi_ValueFromLoopBodyAndContinuing) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %uint = OpTypeInt 32 0
+ %pty = OpTypePointer Private %uint
+ %1 = OpVariable %pty Private
+ %boolpty = OpTypePointer Private %bool
+ %17 = OpVariable %boolpty Private
+ %uint_0 = OpConstant %uint 0
+ %uint_1 = OpConstant %uint 1
+ %true = OpConstantTrue %bool
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %9 = OpLabel
+ %101 = OpLoad %bool %17
+ OpBranch %10
+; Use an outer loop to show we put the new variable in the
+; smallest enclosing scope.
+ %10 = OpLabel
+ OpLoopMerge %99 %89 None
+ OpBranch %20
+ %20 = OpLabel
+ %2 = OpPhi %uint %uint_0 %10 %4 %30 ; gets computed value
+ %5 = OpPhi %uint %uint_1 %10 %7 %30
+ %4 = OpIAdd %uint %2 %uint_1 ; define %4
+ %6 = OpIAdd %uint %4 %uint_1 ; use %4
+ OpLoopMerge %79 %30 None
+ OpBranchConditional %101 %79 %30
+ %30 = OpLabel
+ %7 = OpIAdd %uint %4 %6 ; use %4 again
+ %8 = OpCopyObject %uint %5 ; use %5
+ OpBranchConditional %true %20 %79
+ %79 = OpLabel
+ OpBranch %89
+ %89 = OpLabel
+ OpBranchConditional %true %10 %99
+ %99 = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:ptr<private, u32, read_write> = var undef
+ %2:ptr<private, bool, read_write> = var undef
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %4:bool = load %2
+ loop [b: $B3, c: $B4] { # loop_1
+ $B3: { # body
+ loop [i: $B5, b: $B6, c: $B7] { # loop_2
+ $B5: { # initializer
+ next_iteration 0u, 1u # -> $B6
+ }
+ $B6 (%5:u32, %6:u32): { # body
+ %7:u32 = spirv.add<u32> %5, 1u
+ %8:u32 = spirv.add<u32> %7, 1u
+ if %4 [t: $B8, f: $B9] { # if_1
+ $B8: { # true
+ exit_loop # loop_2
+ }
+ $B9: { # false
+ continue # -> $B7
+ }
+ }
+ unreachable
+ }
+ $B7: { # continuing
+ %9:u32 = spirv.add<u32> %7, %8
+ %10:u32 = let %6
+ %11:bool = not true
+ break_if %11 next_iteration: [ %7, %9 ] # -> [t: exit_loop loop_2, f: $B6]
+ }
+ }
+ continue # -> $B4
+ }
+ $B4: { # continuing
+ %12:bool = not true
+ break_if %12 # -> [t: exit_loop loop_1, f: $B3]
+ }
+ }
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Phi_FromElseAndThen) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %uint = OpTypeInt 32 0
+ %pty = OpTypePointer Private %uint
+ %1 = OpVariable %pty Private
+ %boolpty = OpTypePointer Private %bool
+ %7 = OpVariable %boolpty Private
+ %8 = OpVariable %boolpty Private
+ %uint_0 = OpConstant %uint 0
+ %uint_1 = OpConstant %uint 1
+ %uint_3 = OpConstant %uint 3
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %5 = OpLabel
+ %101 = OpLoad %bool %7
+ %102 = OpLoad %bool %8
+ OpBranch %10
+; Use an outer loop to show we put the new variable in the
+; smallest enclosing scope.
+ %10 = OpLabel
+ OpLoopMerge %99 %89 None
+ OpBranchConditional %101 %99 %20
+ %20 = OpLabel ; if seleciton
+ OpSelectionMerge %79 None
+ OpBranchConditional %102 %30 %40
+ %30 = OpLabel
+ OpBranch %89
+ %40 = OpLabel
+ OpBranch %89
+ %79 = OpLabel ; disconnected selection merge node
+ OpBranch %89
+ %89 = OpLabel
+ %2 = OpPhi %uint %uint_0 %30 %uint_1 %40 %uint_3 %79
+ OpStore %1 %2
+ OpBranch %10
+ %99 = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:ptr<private, u32, read_write> = var undef
+ %2:ptr<private, bool, read_write> = var undef
+ %3:ptr<private, bool, read_write> = var undef
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %5:bool = load %2
+ %6:bool = load %3
+ loop [b: $B3, c: $B4] { # loop_1
+ $B3: { # body
+ if %5 [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_loop # loop_1
+ }
+ $B6: { # false
+ if %6 [t: $B7, f: $B8] { # if_2
+ $B7: { # true
+ continue 0u # -> $B4
+ }
+ $B8: { # false
+ continue 1u # -> $B4
+ }
+ }
+ continue 3u # -> $B4
+ }
+ }
+ unreachable
+ }
+ $B4 (%7:u32): { # continuing
+ store %1, %7
+ next_iteration # -> $B3
+ }
+ }
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, Phi_FromHeaderAndThen) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %uint = OpTypeInt 32 0
+ %pty = OpTypePointer Private %uint
+ %1 = OpVariable %pty Private
+ %boolpty = OpTypePointer Private %bool
+ %7 = OpVariable %boolpty Private
+ %8 = OpVariable %boolpty Private
+ %uint_0 = OpConstant %uint 0
+ %uint_1 = OpConstant %uint 1
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %5 = OpLabel
+ %101 = OpLoad %bool %7
+ %102 = OpLoad %bool %8
+ OpBranch %10
+; Use an outer loop to show we put the new variable in the
+; smallest enclosing scope.
+ %10 = OpLabel
+ OpLoopMerge %99 %89 None
+ OpBranchConditional %101 %99 %20
+ %20 = OpLabel ; if seleciton
+ OpSelectionMerge %79 None
+ OpBranchConditional %102 %30 %89
+ %30 = OpLabel
+ OpBranch %89
+ %79 = OpLabel ; disconnected selection merge node
+ OpUnreachable
+ %89 = OpLabel
+ %2 = OpPhi %uint %uint_0 %20 %uint_1 %30
+ %3 = OpPhi %uint %uint_1 %20 %uint_0 %30
+ %4 = OpIAdd %uint %2 %3
+ OpStore %1 %4
+ OpBranch %10
+ %99 = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+$B1: { # root
+ %1:ptr<private, u32, read_write> = var undef
+ %2:ptr<private, bool, read_write> = var undef
+ %3:ptr<private, bool, read_write> = var undef
+}
+
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B2: {
+ %5:bool = load %2
+ %6:bool = load %3
+ loop [b: $B3, c: $B4] { # loop_1
+ $B3: { # body
+ if %5 [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_loop # loop_1
+ }
+ $B6: { # false
+ if %6 [t: $B7, f: $B8] { # if_2
+ $B7: { # true
+ continue 1u, 0u # -> $B4
+ }
+ $B8: { # false
+ continue 0u, 1u # -> $B4
+ }
+ }
+ unreachable
+ }
+ }
+ unreachable
+ }
+ $B4 (%7:u32, %8:u32): { # continuing
+ %9:u32 = spirv.add<u32> %7, %8
+ store %1, %9
+ next_iteration # -> $B3
+ }
+ }
+ ret
+ }
+}
+)");
+}
+
+// If the only use of a combinatorially computed ID is as the value in an OpPhi, then we still have
+// to emit it. The algorithm fix is to always count uses in Phis. This is the reduced case from the
+// bug report.
+//
+// * The only use of %12 is in the phi.
+// * The only use of %11 is in %12.
+// * Both definitions need to be emitted to the output.
+//
+// https://crbug.com/215
+TEST_F(SpirvParserTest, Phi_UseInPhiCountsAsUse) {
+ EXPECT_IR(
+ R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %true = OpConstantTrue %bool
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %10 = OpLabel
+ %11 = OpLogicalAnd %bool %true %true
+ %12 = OpLogicalNot %bool %11 ;
+ OpSelectionMerge %99 None
+ OpBranchConditional %true %20 %99
+ %20 = OpLabel
+ OpBranch %99
+ %99 = OpLabel
+ %101 = OpPhi %bool %11 %10 %12 %20
+ %102 = OpCopyObject %bool %101 ;; ensure a use of %101
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:bool = and true, true
+ %3:bool = not %2
+ %4:bool = if true [t: $B2, f: $B3] { # if_1
+ $B2: { # true
+ exit_if %3 # if_1
+ }
+ $B3: { # false
+ exit_if %2 # if_1
+ }
+ }
+ %5:bool = let %4
+ ret
+ }
+}
+)");
+}
+
+// Value %999 is defined deep in control flow, then we arrange for it to dominate the backedge of
+// the outer loop. The %999 value is then fed back into the phi in the loop header. So %999 needs
+// to be hoisted out of the loop. The phi assignment needs to use the hoisted variable. The hoisted
+// variable needs to be placed such that its scope encloses that phi in the header of the outer
+// loop. The compiler needs to "see" that there is an implicit use of %999 in the backedge block of
+// that outer loop.
+//
+// https://crbug.com/1649
+TEST_F(SpirvParserTest, Phi_PhiInLoopHeader_FedByHoistedVar_PhiUnused) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %true = OpConstantTrue %bool
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %10 = OpLabel
+ OpBranch %20
+ %20 = OpLabel
+ %101 = OpPhi %bool %true %10 %999 %80
+ OpLoopMerge %99 %80 None
+ OpBranchConditional %true %30 %99
+ %30 = OpLabel
+ OpSelectionMerge %50 None
+ OpBranchConditional %true %40 %50
+ %40 = OpLabel
+ %999 = OpCopyObject %bool %true
+ OpBranch %60
+ %50 = OpLabel
+ OpReturn
+ %60 = OpLabel ; if merge
+ OpBranch %80
+ %80 = OpLabel ; continue target
+ OpBranch %20
+ %99 = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ loop [i: $B2, b: $B3, c: $B4] { # loop_1
+ $B2: { # initializer
+ next_iteration true # -> $B3
+ }
+ $B3 (%2:bool): { # body
+ if true [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ if true [t: $B7, f: $B8] { # if_2
+ $B7: { # true
+ %3:bool = let true
+ continue %3 # -> $B4
+ }
+ $B8: { # false
+ exit_if # if_2
+ }
+ }
+ ret
+ }
+ $B6: { # false
+ exit_loop # loop_1
+ }
+ }
+ unreachable
+ }
+ $B4 (%4:bool): { # continuing
+ next_iteration %4 # -> $B3
+ }
+ }
+ ret
+ }
+}
+)");
+}
+
+// Value %999 is defined deep in control flow, then we arrange for it to dominate the backedge of
+// the outer loop. The %999 value is then fed back into the phi in the loop header. So %999 needs
+// to be hoisted out of the loop. The phi assignment needs to use the hoisted variable. The hoisted
+// variable needs to be placed such that its scope encloses that phi in the header of the outer
+// loop. The compiler needs to "see" that there is an implicit use of %999 in the backedge block of
+// that outer loop.
+//
+// https://crbug.com/1649
+TEST_F(SpirvParserTest, Phi_PhiInLoopHeader_FedByHoistedVar_PhiUsed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %true = OpConstantTrue %bool
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %10 = OpLabel
+ OpBranch %20
+ %20 = OpLabel
+ %101 = OpPhi %bool %true %10 %999 %80
+ OpLoopMerge %99 %80 None
+ OpBranchConditional %true %30 %99
+ %30 = OpLabel
+ OpSelectionMerge %50 None
+ OpBranchConditional %true %40 %50
+ %40 = OpLabel
+ %999 = OpCopyObject %bool %true
+ OpBranch %60
+ %50 = OpLabel
+ OpReturn
+ %60 = OpLabel ; if merge
+ OpBranch %80
+ %80 = OpLabel ; continue target
+ OpBranch %20
+ %99 = OpLabel
+ %1000 = OpCopyObject %bool %101
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:bool = loop [i: $B2, b: $B3, c: $B4] { # loop_1
+ $B2: { # initializer
+ next_iteration true # -> $B3
+ }
+ $B3 (%3:bool): { # body
+ if true [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ if true [t: $B7, f: $B8] { # if_2
+ $B7: { # true
+ %4:bool = let true
+ continue %4 # -> $B4
+ }
+ $B8: { # false
+ exit_if # if_2
+ }
+ }
+ ret
+ }
+ $B6: { # false
+ exit_loop %3 # loop_1
+ }
+ }
+ unreachable
+ }
+ $B4 (%5:bool): { # continuing
+ next_iteration %5 # -> $B3
+ }
+ }
+ %6:bool = let %2
+ ret
+ }
+}
+)");
+}
+
+// This is a reduction of one of the hard parts of test case
+// vk-gl-cts/graphicsfuzz/stable-binarysearch-tree-false-if-discard-loop/1.spvasm
+// In particular, see the data flow around %114 in that case.
+//
+// Here value %999 is is a *phi* defined deep in control flow, then we arrange for it to dominate
+// the backedge of the outer loop. The %999 value is then fed back into the phi in the loop header.
+// The variable generated to hold the %999 value needs to be placed such that its scope encloses
+// that phi in the header of the outer loop. The compiler needs to "see" that there is an implicit
+// use of %999 in the backedge block of that outer loop.
+//
+// https://crbug.com/1649
+TEST_F(SpirvParserTest, Phi_PhiInLoopHeader_FedByPhi_PhiUnused) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %true = OpConstantTrue %bool
+ %false = OpConstantFalse %bool
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %10 = OpLabel
+ OpBranch %20
+ %20 = OpLabel
+ %101 = OpPhi %bool %true %10 %999 %80
+ OpLoopMerge %99 %80 None
+ OpBranchConditional %true %99 %30
+ %30 = OpLabel
+ OpLoopMerge %70 %60 None
+ OpBranch %40
+ %40 = OpLabel
+ OpBranchConditional %true %60 %50
+ %50 = OpLabel
+ OpBranch %60
+ %60 = OpLabel ; inner continue
+ %999 = OpPhi %bool %true %40 %false %50
+ OpBranchConditional %true %70 %30
+ %70 = OpLabel ; inner merge
+ OpBranch %80
+ %80 = OpLabel ; outer continue target
+ OpBranch %20
+ %99 = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ loop [i: $B2, b: $B3, c: $B4] { # loop_1
+ $B2: { # initializer
+ next_iteration true # -> $B3
+ }
+ $B3 (%2:bool): { # body
+ if true [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_loop # loop_1
+ }
+ $B6: { # false
+ %3:bool = loop [b: $B7, c: $B8] { # loop_2
+ $B7: { # body
+ if true [t: $B9, f: $B10] { # if_2
+ $B9: { # true
+ continue true # -> $B8
+ }
+ $B10: { # false
+ continue false # -> $B8
+ }
+ }
+ unreachable
+ }
+ $B8 (%4:bool): { # continuing
+ break_if true exit_loop: [ %4 ] # -> [t: exit_loop loop_2, f: $B7]
+ }
+ }
+ continue %3 # -> $B4
+ }
+ }
+ unreachable
+ }
+ $B4 (%5:bool): { # continuing
+ next_iteration %5 # -> $B3
+ }
+ }
+ ret
+ }
+}
+)");
+}
+
+// This is a reduction of one of the hard parts of test case
+// vk-gl-cts/graphicsfuzz/stable-binarysearch-tree-false-if-discard-loop/1.spvasm
+// In particular, see the data flow around %114 in that case.
+//
+// Here value %999 is is a *phi* defined deep in control flow, then we arrange for it to dominate
+// the backedge of the outer loop. The %999 value is then fed back into the phi in the loop header.
+// The variable generated to hold the %999 value needs to be placed such that its scope encloses
+// that phi in the header of the outer loop. The compiler needs to "see" that there is an implicit
+// use of %999 in the backedge block of that outer loop.
+//
+// https://crbug.com/1649
+TEST_F(SpirvParserTest, Phi_PhiInLoopHeader_FedByPhi_PhiUsed) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %bool = OpTypeBool
+ %true = OpConstantTrue %bool
+ %false = OpConstantFalse %bool
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %10 = OpLabel
+ OpBranch %20
+ %20 = OpLabel
+ %101 = OpPhi %bool %true %10 %999 %80
+ OpLoopMerge %99 %80 None
+ OpBranchConditional %true %99 %30
+ %30 = OpLabel
+ OpLoopMerge %70 %60 None
+ OpBranch %40
+ %40 = OpLabel
+ OpBranchConditional %true %60 %50
+ %50 = OpLabel
+ OpBranch %60
+ %60 = OpLabel ; inner continue
+ %999 = OpPhi %bool %true %40 %false %50
+ OpBranchConditional %true %70 %30
+ %70 = OpLabel ; inner merge
+ OpBranch %80
+ %80 = OpLabel ; outer continue target
+ OpBranch %20
+ %99 = OpLabel
+ %1000 = OpCopyObject %bool %101
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ %2:bool = loop [i: $B2, b: $B3, c: $B4] { # loop_1
+ $B2: { # initializer
+ next_iteration true # -> $B3
+ }
+ $B3 (%3:bool): { # body
+ if true [t: $B5, f: $B6] { # if_1
+ $B5: { # true
+ exit_loop %3 # loop_1
+ }
+ $B6: { # false
+ %4:bool = loop [b: $B7, c: $B8] { # loop_2
+ $B7: { # body
+ if true [t: $B9, f: $B10] { # if_2
+ $B9: { # true
+ continue true # -> $B8
+ }
+ $B10: { # false
+ continue false # -> $B8
+ }
+ }
+ unreachable
+ }
+ $B8 (%5:bool): { # continuing
+ break_if true exit_loop: [ %5 ] # -> [t: exit_loop loop_2, f: $B7]
+ }
+ }
+ continue %4 # -> $B4
+ }
+ }
+ unreachable
+ }
+ $B4 (%6:bool): { # continuing
+ next_iteration %6 # -> $B3
+ }
+ }
+ %7:bool = let %2
+ ret
+ }
+}
+)");
+}
+
+// A phi in an unreachable block may have no operands.
+TEST_F(SpirvParserTest, Phi_UnreachableLoopMerge) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ %void = OpTypeVoid
+ %uint = OpTypeInt 32 0
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %10 = OpLabel
+ OpBranch %99
+ %99 = OpLabel
+ OpLoopMerge %101 %99 None
+ OpBranch %99
+ %101 = OpLabel
+ %102 = OpPhi %uint
+ OpUnreachable
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(1u, 1u, 1u) func():void {
+ $B1: {
+ loop [b: $B2, c: $B3] { # loop_1
+ $B2: { # body
+ continue # -> $B3
+ }
+ $B3: { # continuing
+ next_iteration # -> $B2
+ }
+ }
+ unreachable
+ }
+}
+)");
+}
+
+} // namespace
+} // namespace tint::spirv::reader
