Remove fallthrough support from SPIRV-Reader.
This CL removes support for fallthrough from the SPIRV-Reader.
Bug: tint:1644
Change-Id: I80b63d627960a82ba90de83af407c539b0442080
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/109004
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
diff --git a/src/tint/reader/spirv/function.cc b/src/tint/reader/spirv/function.cc
index 07b5980..a71aa42 100644
--- a/src/tint/reader/spirv/function.cc
+++ b/src/tint/reader/spirv/function.cc
@@ -25,7 +25,6 @@
#include "src/tint/ast/call_statement.h"
#include "src/tint/ast/continue_statement.h"
#include "src/tint/ast/discard_statement.h"
-#include "src/tint/ast/fallthrough_statement.h"
#include "src/tint/ast/if_statement.h"
#include "src/tint/ast/loop_statement.h"
#include "src/tint/ast/return_statement.h"
@@ -2159,8 +2158,7 @@
// For each branch encountered, classify each edge (S,T) as:
// - a back-edge
// - a structured exit (specific ways of branching to enclosing construct)
- // - a normal (forward) edge, either natural control flow or a case
- // fallthrough
+ // - a normal (forward) edge, either natural control flow or a case fallthrough
//
// If more than one block is targeted by a normal edge, then S must be a
// structured header.
@@ -2194,11 +2192,10 @@
// There should only be one backedge per backedge block.
uint32_t num_backedges = 0;
- // Track destinations for normal forward edges, either kForward
- // or kCaseFallThrough. These count toward the need
- // to have a merge instruction. We also track kIfBreak edges
- // because when used with normal forward edges, we'll need
- // to generate a flow guard variable.
+ // Track destinations for normal forward edges, either kForward or kCaseFallThrough.
+ // These count toward the need to have a merge instruction. We also track kIfBreak edges
+ // because when used with normal forward edges, we'll need to generate a flow guard
+ // variable.
utils::Vector<uint32_t, 4> normal_forward_edges;
utils::Vector<uint32_t, 4> if_break_edges;
@@ -2376,6 +2373,12 @@
<< dest_construct.begin_id << " (dominance rule violated)";
}
}
+
+ // Error on the fallthrough at the end in order to allow the better error messages
+ // from the above checks to happen.
+ if (edge_kind == EdgeKind::kCaseFallThrough) {
+ return Fail() << "Fallthrough not permitted in WGSL";
+ }
} // end forward edge
} // end successor
@@ -3261,11 +3264,9 @@
// The fallthrough case is special because WGSL requires the fallthrough
// statement to be last in the case clause.
if (true_kind == EdgeKind::kCaseFallThrough) {
- return EmitConditionalCaseFallThrough(block_info, cond, false_kind, *false_info,
- true);
+ return Fail() << "Fallthrough not supported in WGSL";
} else if (false_kind == EdgeKind::kCaseFallThrough) {
- return EmitConditionalCaseFallThrough(block_info, cond, true_kind, *true_info,
- false);
+ return Fail() << "Fallthrough not supported in WGSL";
}
// At this point, at most one edge is kForward or kIfBreak.
@@ -3304,7 +3305,7 @@
const ast::Statement* FunctionEmitter::MakeBranchDetailed(const BlockInfo& src_info,
const BlockInfo& dest_info,
bool forced,
- std::string* flow_guard_name_ptr) const {
+ std::string* flow_guard_name_ptr) {
auto kind = src_info.succ_edge.find(dest_info.id)->second;
switch (kind) {
case EdgeKind::kBack:
@@ -3367,8 +3368,10 @@
// merge block is implicit.
break;
}
- case EdgeKind::kCaseFallThrough:
- return create<ast::FallthroughStatement>(Source{});
+ case EdgeKind::kCaseFallThrough: {
+ Fail() << "Fallthrough not supported in WGSL";
+ return nullptr;
+ }
case EdgeKind::kForward:
// Unconditional forward branch is implicit.
break;
@@ -3398,45 +3401,6 @@
return if_stmt;
}
-bool FunctionEmitter::EmitConditionalCaseFallThrough(const BlockInfo& src_info,
- const ast::Expression* cond,
- EdgeKind other_edge_kind,
- const BlockInfo& other_dest,
- bool fall_through_is_true_branch) {
- // In WGSL, the fallthrough statement must come last in the case clause.
- // So we'll emit an if statement for the other branch, and then emit
- // the fallthrough.
-
- // We have two distinct destinations. But we only get here if this
- // is a normal terminator; in particular the source block is *not* the
- // start of an if-selection. So at most one branch is a kForward or
- // kCaseFallThrough.
- if (other_edge_kind == EdgeKind::kForward) {
- return Fail() << "internal error: normal terminator OpBranchConditional has "
- "both forward and fallthrough edges";
- }
- if (other_edge_kind == EdgeKind::kIfBreak) {
- return Fail() << "internal error: normal terminator OpBranchConditional has "
- "both IfBreak and fallthrough edges. Violates nesting rule";
- }
- if (other_edge_kind == EdgeKind::kBack) {
- return Fail() << "internal error: normal terminator OpBranchConditional has "
- "both backedge and fallthrough edges. Violates nesting rule";
- }
- auto* other_branch = MakeForcedBranch(src_info, other_dest);
- if (other_branch == nullptr) {
- return Fail() << "internal error: expected a branch for edge-kind " << int(other_edge_kind);
- }
- if (fall_through_is_true_branch) {
- AddStatement(MakeSimpleIf(cond, nullptr, other_branch));
- } else {
- AddStatement(MakeSimpleIf(cond, other_branch, nullptr));
- }
- AddStatement(create<ast::FallthroughStatement>(Source{}));
-
- return success();
-}
-
bool FunctionEmitter::EmitStatementsInBasicBlock(const BlockInfo& block_info,
bool* already_emitted) {
if (*already_emitted) {
diff --git a/src/tint/reader/spirv/function.h b/src/tint/reader/spirv/function.h
index 8d06735..917b374 100644
--- a/src/tint/reader/spirv/function.h
+++ b/src/tint/reader/spirv/function.h
@@ -34,15 +34,13 @@
//
// The edge kinds are used in many ways.
//
-// For example, consider the edges leaving a basic block and going to distinct
-// targets. If the total number of kForward + kIfBreak + kCaseFallThrough edges
-// is more than 1, then the block must be a structured header, i.e. it needs
-// a merge instruction to declare the control flow divergence and associated
-// reconvergence point. Those those edge kinds count toward divergence
-// because SPIR-v is designed to easily map back to structured control flow
-// in GLSL (and C). In GLSL and C, those forward-flow edges don't have a
-// special statement to express them. The other forward edges: kSwitchBreak,
-// kLoopBreak, and kLoopContinue directly map to 'break', 'break', and
+// For example, consider the edges leaving a basic block and going to distinct targets. If the
+// total number of kForward + kIfBreak + kCaseFallThrough edges is more than 1, then the block must
+// be a structured header, i.e. it needs a merge instruction to declare the control flow divergence
+// and associated reconvergence point. Those those edge kinds count toward divergence because
+// SPIR-V is designed to easily map back to structured control flow in GLSL (and C). In GLSL and C,
+// those forward-flow edges don't have a special statement to express them. The other forward
+// edges: kSwitchBreak, kLoopBreak, and kLoopContinue directly map to 'break', 'break', and
// 'continue', respectively.
enum class EdgeKind {
// A back-edge: An edge from a node to one of its ancestors in a depth-first
@@ -64,7 +62,8 @@
// This can only occur for an "if" selection, i.e. where the selection
// header ends in OpBranchConditional.
kIfBreak,
- // An edge from one switch case to the next sibling switch case.
+ // An edge from one switch case to the next sibling switch case. Note, this is not valid in WGSL
+ // at the moment and will trigger an ICE if encountered. It is here for completeness.
kCaseFallThrough,
// None of the above.
kForward
@@ -708,8 +707,7 @@
/// Emits code for terminators, but that aren't part of entering or
/// resolving structured control flow. That is, if the basic block
- /// terminator calls for it, emit the fallthrough, break, continue, return,
- /// or kill commands.
+ /// terminator calls for it, emit the fallthrough break, continue, return, or kill commands.
/// @param block_info the block with the terminator to emit (if any)
/// @returns false if emission failed
bool EmitNormalTerminator(const BlockInfo& block_info);
@@ -722,7 +720,7 @@
/// @param src_info the source block
/// @param dest_info the destination block
/// @returns the new statement, or a null statement
- const ast::Statement* MakeBranch(const BlockInfo& src_info, const BlockInfo& dest_info) const {
+ const ast::Statement* MakeBranch(const BlockInfo& src_info, const BlockInfo& dest_info) {
return MakeBranchDetailed(src_info, dest_info, false, nullptr);
}
@@ -732,8 +730,7 @@
/// @param src_info the source block
/// @param dest_info the destination block
/// @returns the new statement, or a null statement
- const ast::Statement* MakeForcedBranch(const BlockInfo& src_info,
- const BlockInfo& dest_info) const {
+ const ast::Statement* MakeForcedBranch(const BlockInfo& src_info, const BlockInfo& dest_info) {
return MakeBranchDetailed(src_info, dest_info, true, nullptr);
}
@@ -754,7 +751,7 @@
const ast::Statement* MakeBranchDetailed(const BlockInfo& src_info,
const BlockInfo& dest_info,
bool forced,
- std::string* flow_guard_name_ptr) const;
+ std::string* flow_guard_name_ptr);
/// Returns a new if statement with the given statements as the then-clause
/// and the else-clause. Either or both clauses might be nullptr. If both
diff --git a/src/tint/reader/spirv/function_cfg_test.cc b/src/tint/reader/spirv/function_cfg_test.cc
index 538657b..88dc78e 100644
--- a/src/tint/reader/spirv/function_cfg_test.cc
+++ b/src/tint/reader/spirv/function_cfg_test.cc
@@ -1288,43 +1288,7 @@
EXPECT_THAT(fe.block_order(), ElementsAre(10, 40, 20, 30, 99));
}
-TEST_F(SpvParserCFGTest, ComputeBlockOrder_RespectSwitchCaseFallthrough) {
- auto assembly = CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpSelectionMerge %99 None
- ; SPIR-V validation requires a fallthrough destination to immediately
- ; follow the source. So %20 -> %40, %30 -> %50
- OpSwitch %selector %99 20 %20 40 %40 30 %30 50 %50
-
- %50 = OpLabel
- OpBranch %99
-
- %99 = OpLabel
- OpReturn
-
- %40 = OpLabel
- OpBranch %99
-
- %30 = OpLabel
- OpBranch %50 ; fallthrough
-
- %20 = OpLabel
- OpBranch %40 ; fallthrough
-
- OpFunctionEnd
- )";
- auto p = parser(test::Assemble(assembly));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- fe.RegisterBasicBlocks();
- fe.ComputeBlockOrderAndPositions();
-
- EXPECT_THAT(fe.block_order(), ElementsAre(10, 30, 50, 20, 40, 99)) << assembly;
-}
-
-TEST_F(SpvParserCFGTest, ComputeBlockOrder_RespectSwitchCaseFallthrough_FromDefault) {
+TEST_F(SpvParserCFGTest, ClassifyCFGEdges_Fallthrough_IsError) {
auto assembly = CommonTypes() + R"(
%100 = OpFunction %void None %voidfn
@@ -1352,122 +1316,9 @@
auto p = parser(test::Assemble(assembly));
ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
auto fe = p->function_emitter(100);
- fe.RegisterBasicBlocks();
- fe.ComputeBlockOrderAndPositions();
-
- EXPECT_THAT(fe.block_order(), ElementsAre(10, 20, 80, 30, 40, 99)) << assembly;
-}
-
-TEST_F(SpvParserCFGTest, ComputeBlockOrder_RespectSwitchCaseFallthrough_FromCaseToDefaultToCase) {
- auto assembly = CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpSelectionMerge %99 None
- OpSwitch %selector %80 20 %20 30 %30
-
- %20 = OpLabel
- OpBranch %80 ; fallthrough to default
-
- %80 = OpLabel ; the default case
- OpBranch %30 ; fallthrough to 30
-
- %30 = OpLabel
- OpBranch %99
-
- %99 = OpLabel ; dominated by %30, so follow %30
- OpReturn
-
- OpFunctionEnd
- )";
- auto p = parser(test::Assemble(assembly));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- fe.RegisterBasicBlocks();
- fe.ComputeBlockOrderAndPositions();
-
- EXPECT_THAT(fe.block_order(), ElementsAre(10, 20, 80, 30, 99)) << assembly;
-}
-
-TEST_F(SpvParserCFGTest, ComputeBlockOrder_SwitchCasesFallthrough_OppositeDirections) {
- auto assembly = CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpSelectionMerge %99 None
- OpSwitch %selector %99 20 %20 30 %30 40 %40 50 %50
-
- %99 = OpLabel
- OpReturn
-
- %20 = OpLabel
- OpBranch %30 ; forward
-
- %40 = OpLabel
- OpBranch %99
-
- %30 = OpLabel
- OpBranch %99
-
- ; SPIR-V doesn't actually allow a fall-through that goes backward in the
- ; module. But the block ordering algorithm tolerates it.
- %50 = OpLabel
- OpBranch %40 ; backward
-
- OpFunctionEnd
- )";
- auto p = parser(test::Assemble(assembly));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- fe.RegisterBasicBlocks();
- fe.ComputeBlockOrderAndPositions();
-
- EXPECT_THAT(fe.block_order(), ElementsAre(10, 50, 40, 20, 30, 99)) << assembly;
-
- // We're deliberately testing a case that SPIR-V doesn't allow.
- p->DeliberatelyInvalidSpirv();
-}
-
-TEST_F(SpvParserCFGTest, ComputeBlockOrder_RespectSwitchCaseFallthrough_Interleaved) {
- auto assembly = CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpSelectionMerge %99 None
- ; SPIR-V validation requires a fallthrough destination to immediately
- ; follow the source. So %20 -> %40
- OpSwitch %selector %99 20 %20 40 %40 30 %30 50 %50
-
- %99 = OpLabel
- OpReturn
-
- %20 = OpLabel
- OpBranch %40
-
- %30 = OpLabel
- OpBranch %50
-
- %40 = OpLabel
- OpBranch %60
-
- %50 = OpLabel
- OpBranch %70
-
- %60 = OpLabel
- OpBranch %99
-
- %70 = OpLabel
- OpBranch %99
-
- OpFunctionEnd
- )";
- auto p = parser(test::Assemble(assembly));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- fe.RegisterBasicBlocks();
- fe.ComputeBlockOrderAndPositions();
-
- EXPECT_THAT(fe.block_order(), ElementsAre(10, 30, 50, 70, 20, 40, 60, 99)) << assembly;
+ EXPECT_FALSE(FlowClassifyCFGEdges(&fe)) << p->error();
+ // Some further processing
+ EXPECT_THAT(p->error(), Eq("Fallthrough not permitted in WGSL"));
}
TEST_F(SpvParserCFGTest, ComputeBlockOrder_Nest_If_Contains_If) {
@@ -1566,54 +1417,6 @@
EXPECT_THAT(fe.block_order(), ElementsAre(10, 20, 30, 40, 49, 50, 60, 70, 79, 99)) << assembly;
}
-TEST_F(SpvParserCFGTest, ComputeBlockOrder_Nest_IfFallthrough_In_SwitchCase) {
- auto assembly = CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpSelectionMerge %99 None
- OpSwitch %selector %50 20 %20 50 %50
-
- %99 = OpLabel
- OpReturn
-
- %20 = OpLabel
- OpSelectionMerge %49 None
- OpBranchConditional %cond %30 %40
-
- %49 = OpLabel
- OpBranchConditional %cond %99 %50 ; fallthrough
-
- %30 = OpLabel
- OpBranch %49
-
- %40 = OpLabel
- OpBranch %49
-
- %50 = OpLabel
- OpSelectionMerge %79 None
- OpBranchConditional %cond %60 %70
-
- %79 = OpLabel
- OpBranch %99
-
- %60 = OpLabel
- OpBranch %79
-
- %70 = OpLabel
- OpBranch %79
-
- OpFunctionEnd
- )";
- auto p = parser(test::Assemble(assembly));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- fe.RegisterBasicBlocks();
- fe.ComputeBlockOrderAndPositions();
-
- EXPECT_THAT(fe.block_order(), ElementsAre(10, 20, 30, 40, 49, 50, 60, 70, 79, 99)) << assembly;
-}
-
TEST_F(SpvParserCFGTest, ComputeBlockOrder_Nest_IfBreak_In_SwitchCase) {
auto assembly = CommonTypes() + R"(
%100 = OpFunction %void None %voidfn
@@ -3904,7 +3707,7 @@
OpSwitch %selector %80 60 %60
%60 = OpLabel
- OpBranch %89 ; fallthrough
+ OpBranch %89
%80 = OpLabel ; default for both switches
OpBranch %89
@@ -4974,8 +4777,6 @@
}
TEST_F(SpvParserCFGTest, ClassifyCFGEdges_IfBreak_EscapeSwitchCase_IsError) {
- // Code generation assumes that you can't have kCaseFallThrough and kIfBreak
- // from the same OpBranchConditional.
// This checks one direction of that, where the IfBreak is shown it can't
// escape a switch case.
auto assembly = CommonTypes() + R"(
@@ -5936,440 +5737,6 @@
"starting at block 30; branch bypasses merge block 59"));
}
-TEST_F(SpvParserCFGTest, ClassifyCFGEdges_Fallthrough_CaseTailToCase) {
- auto assembly = CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpSelectionMerge %99 None
- OpSwitch %selector %99 20 %20 40 %40
-
- %20 = OpLabel ; case 20
- OpBranch %30
-
- %30 = OpLabel
- OpBranch %40 ; fallthrough
-
- %40 = OpLabel ; case 40
- OpBranch %99
-
- %99 = OpLabel
- OpReturn
-
- OpFunctionEnd
-)";
- auto p = parser(test::Assemble(assembly));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_TRUE(FlowClassifyCFGEdges(&fe));
-
- auto* bi = fe.GetBlockInfo(30);
- ASSERT_NE(bi, nullptr);
- EXPECT_EQ(bi->succ_edge.count(40), 1u);
- EXPECT_EQ(bi->succ_edge[40], EdgeKind::kCaseFallThrough);
-}
-
-TEST_F(SpvParserCFGTest, ClassifyCFGEdges_Fallthrough_CaseTailToDefaultNotMerge) {
- auto assembly = CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpSelectionMerge %99 None
- OpSwitch %selector %40 20 %20
-
- %20 = OpLabel ; case 20
- OpBranch %30
-
- %30 = OpLabel
- OpBranch %40 ; fallthrough
-
- %40 = OpLabel ; case 40
- OpBranch %99
-
- %99 = OpLabel
- OpReturn
-
- OpFunctionEnd
-)";
- auto p = parser(test::Assemble(assembly));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_TRUE(FlowClassifyCFGEdges(&fe));
-
- auto* bi = fe.GetBlockInfo(30);
- ASSERT_NE(bi, nullptr);
- EXPECT_EQ(bi->succ_edge.count(40), 1u);
- EXPECT_EQ(bi->succ_edge[40], EdgeKind::kCaseFallThrough);
-}
-
-TEST_F(SpvParserCFGTest, ClassifyCFGEdges_Fallthrough_DefaultToCase) {
- auto assembly = CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpSelectionMerge %99 None
- OpSwitch %selector %20 40 %40
-
- %20 = OpLabel ; default
- OpBranch %30
-
- %30 = OpLabel
- OpBranch %40 ; fallthrough
-
- %40 = OpLabel ; case 40
- OpBranch %99
-
- %99 = OpLabel
- OpReturn
-
- OpFunctionEnd
-)";
- auto p = parser(test::Assemble(assembly));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_TRUE(FlowClassifyCFGEdges(&fe));
-
- auto* bi = fe.GetBlockInfo(30);
- ASSERT_NE(bi, nullptr);
- EXPECT_EQ(bi->succ_edge.count(40), 1u);
- EXPECT_EQ(bi->succ_edge[40], EdgeKind::kCaseFallThrough);
-}
-
-TEST_F(SpvParserCFGTest, ClassifyCFGEdges_Fallthrough_BranchConditionalWith_IfBreak_IsError) {
- // Code generation assumes OpBranchConditional can't have kCaseFallThrough
- // with kIfBreak.
- auto assembly = CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpSelectionMerge %99 None ; Set up if-break to %99
- OpBranchConditional %cond %20 %99
-
- %20 = OpLabel
- OpSelectionMerge %80 None ; switch-selection
- OpSwitch %selector %80 30 %30 40 %40
-
- %30 = OpLabel ; first case
- ; branch to %99 would be an if-break, but it bypasess the switch merge
- ; Also has case fall-through
- OpBranchConditional %cond2 %99 %40
-
- %40 = OpLabel ; second case
- OpBranch %80
-
- %80 = OpLabel ; switch-selection's merge
- OpBranch %99
-
- %99 = OpLabel ; if-selection's merge
- OpReturn
-
- OpFunctionEnd
-)";
- auto p = parser(test::Assemble(assembly));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_FALSE(FlowClassifyCFGEdges(&fe));
- EXPECT_THAT(p->error(), Eq("Branch from block 30 to block 99 is an invalid exit from "
- "construct starting at block 20; branch bypasses merge block 80"));
-}
-
-TEST_F(SpvParserCFGTest, ClassifyCFGEdges_Fallthrough_BranchConditionalWith_Forward_IsError) {
- // Code generation assumes OpBranchConditional can't have kCaseFallThrough
- // with kForward.
- auto assembly = CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpSelectionMerge %99 None ; switch-selection
- OpSwitch %selector %99 20 %20 30 %30
-
- ; Try to make branch to 35 a kForward branch
- %20 = OpLabel ; first case
- OpBranchConditional %cond2 %25 %30
-
- %25 = OpLabel
- OpBranch %99
-
- %30 = OpLabel ; second case
- OpBranch %99
-
- %99 = OpLabel ; if-selection's merge
- OpReturn
-
- OpFunctionEnd
-)";
- auto p = parser(test::Assemble(assembly));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_FALSE(FlowClassifyCFGEdges(&fe));
- EXPECT_THAT(p->error(), Eq("Control flow diverges at block 20 (to 25, 30) but it is not "
- "a structured header (it has no merge instruction)"));
-}
-
-TEST_F(SpvParserCFGTest,
- ClassifyCFGEdges_Fallthrough_BranchConditionalWith_Back_LoopOnOutside_IsError) { // NOLINT
- // Code generation assumes OpBranchConditional can't have kCaseFallThrough
- // with kBack.
- //
- // This test has the loop on the outside. The backedge coming from a case
- // clause means the switch is inside the continue construct, and the nesting
- // of the switch's merge means the backedge is coming from a block that is not
- // at the end of the continue construct.
- auto assembly = CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpBranch %20
-
- %20 = OpLabel
- OpLoopMerge %99 %30 None
- OpBranch %30
-
- %30 = OpLabel ; continue target and
- OpSelectionMerge %80 None ; switch-selection
- OpSwitch %selector %80 40 %40 50 %50
-
- ; try to make a back edge with a fallthrough
- %40 = OpLabel ; first case
- OpBranchConditional %cond2 %20 %50
-
- %50 = OpLabel ; second case
- OpBranch %80
-
- %80 = OpLabel ; switch merge
- OpBranch %20 ; also backedge
-
- %99 = OpLabel ; loop merge
- OpReturn
-
- OpFunctionEnd
-)";
- auto p = parser(test::Assemble(assembly));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_FALSE(FlowClassifyCFGEdges(&fe));
- EXPECT_THAT(p->error(), Eq("Invalid exit (40->20) from continue construct: 40 is not the "
- "last block in the continue construct starting at 30 "
- "(violates post-dominance rule)"));
-}
-
-TEST_F(
- SpvParserCFGTest,
- FindSwitchCaseSelectionHeaders_Fallthrough_BranchConditionalWith_Back_LoopOnInside_FallthroughIsMerge_IsError) { // NOLINT
- // Code generation assumes OpBranchConditional can't have kCaseFallThrough
- // with kBack.
- //
- // This test has the loop on the inside. The merge block is also the
- // fallthrough target.
- auto assembly = CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel ; continue target and
- OpSelectionMerge %99 None ; switch-selection
- OpSwitch %selector %99 20 %20 50 %50
-
- %20 = OpLabel ; first case, and loop header
- OpLoopMerge %50 %40 None
- OpBranch %40
-
- ; try to make a back edge with a fallthrough
- %40 = OpLabel
- OpBranchConditional %cond2 %20 %50
-
- %50 = OpLabel ; second case. also the loop merge ; header must dominate its merge block !
- OpBranch %99
-
- %99 = OpLabel ; switch merge
- OpReturn
-
- OpFunctionEnd
-)";
- auto p = parser(test::Assemble(assembly));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_FALSE(FlowFindSwitchCaseHeaders(&fe));
- EXPECT_THAT(fe.block_order(), ElementsAre(10, 20, 40, 50, 99));
- EXPECT_THAT(p->error(), Eq("Block 50 is a case block for switch-selection header 10 and "
- "also the merge block for 20 (violates dominance rule)"));
-}
-
-TEST_F(
- SpvParserCFGTest,
- ClassifyCFGEdges_Fallthrough_BranchConditionalWith_Back_LoopOnInside_FallthroughIsNotMerge_IsError) { // NOLINT
- // Code generation assumes OpBranchConditional can't have kCaseFallThrough
- // with kBack.
- //
- // This test has the loop on the inside. The merge block is not the merge
- // target But the block order gets messed up because of the weird
- // connectivity.
- auto assembly = CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel ; continue target and
- OpSelectionMerge %99 None ; switch-selection
- OpSwitch %selector %99 20 %20 50 %50
-
- %20 = OpLabel ; first case, and loop header
- OpLoopMerge %45 %40 None ; move the merge to an unreachable block
- OpBranch %40
-
- ; try to make a back edge with a fallthrough
- %40 = OpLabel
- OpBranchConditional %cond2 %20 %50
-
- %45 = OpLabel ; merge for the loop
- OpUnreachable
-
- %50 = OpLabel ; second case. target of fallthrough
- OpBranch %99
-
- %99 = OpLabel ; switch merge
- OpReturn
-
- OpFunctionEnd
-)";
- auto p = parser(test::Assemble(assembly));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_FALSE(FlowClassifyCFGEdges(&fe));
- EXPECT_THAT(p->error(), Eq("Branch from 10 to 50 bypasses continue target 40 "
- "(dominance rule violated)"));
-}
-
-TEST_F(
- SpvParserCFGTest,
- ClassifyCFGEdges_Fallthrough_BranchConditionalWith_Back_LoopOnInside_NestedMerge_IsError) { // NOLINT
- // Code generation assumes OpBranchConditional can't have kCaseFallThrough
- // with kBack.
- //
- // This test has the loop on the inside. The fallthrough is an invalid exit
- // from the loop. However, the block order gets all messed up because going
- // from 40 to 50 ends up pulling in 99
- auto assembly = CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel ; continue target and
- OpSelectionMerge %99 None ; switch-selection
- OpSwitch %selector %99 20 %20 50 %50
-
- %20 = OpLabel ; first case, and loop header
- OpLoopMerge %49 %40 None
- OpBranch %40
-
- ; try to make a back edge with a fallthrough
- %40 = OpLabel
- OpBranchConditional %cond2 %20 %50
-
- %49 = OpLabel ; loop merge
- OpBranch %99
-
- %50 = OpLabel ; second case
- OpBranch %99
-
- %99 = OpLabel ; switch merge
- OpReturn
-
- OpFunctionEnd
-)";
- auto p = parser(test::Assemble(assembly));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_FALSE(FlowClassifyCFGEdges(&fe));
- EXPECT_THAT(fe.block_order(), ElementsAre(10, 20, 40, 50, 49, 99));
- EXPECT_THAT(p->error(), Eq("Branch from 10 to 50 bypasses continue target 40 "
- "(dominance rule violated)"));
-}
-
-TEST_F(SpvParserCFGTest, ClassifyCFGEdges_Fallthrough_CaseNonTailToCase_TrueBranch) {
- // This is an unusual one, and is an error. Structurally it looks like this:
- // switch (val) {
- // case 0: {
- // if (cond) {
- // fallthrough;
- // }
- // something = 1;
- // }
- // case 1: { }
- // }
- auto assembly = CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpSelectionMerge %99 None
- OpSwitch %selector %99 20 %20 50 %50
-
- %20 = OpLabel
- OpSelectionMerge %49 None
- OpBranchConditional %cond %30 %49
-
- %30 = OpLabel
- OpBranch %50 ; attempt to fallthrough
-
- %49 = OpLabel
- OpBranch %99
-
- %50 = OpLabel
- OpBranch %99
-
- %99 = OpLabel
- OpReturn
-
- OpFunctionEnd
-)";
- auto p = parser(test::Assemble(assembly));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_FALSE(FlowClassifyCFGEdges(&fe));
- EXPECT_THAT(p->error(),
- Eq("Branch from 10 to 50 bypasses header 20 (dominance rule violated)"));
-}
-
-TEST_F(SpvParserCFGTest, ClassifyCFGEdges_Fallthrough_CaseNonTailToCase_FalseBranch) {
- // Like previous test, but taking the false branch.
-
- // This is an unusual one, and is an error. Structurally it looks like this:
- // switch (val) {
- // case 0: {
- // if (cond) {
- // fallthrough;
- // }
- // something = 1;
- // }
- // case 1: { }
- // }
- auto assembly = CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpSelectionMerge %99 None
- OpSwitch %selector %99 20 %20 50 %50
-
- %20 = OpLabel
- OpSelectionMerge %49 None
- OpBranchConditional %cond %49 %30 ;; this is the difference
-
- %30 = OpLabel
- OpBranch %50 ; attempt to fallthrough
-
- %49 = OpLabel
- OpBranch %99
-
- %50 = OpLabel
- OpBranch %99
-
- %99 = OpLabel
- OpReturn
-
- OpFunctionEnd
-)";
- auto p = parser(test::Assemble(assembly));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_FALSE(FlowClassifyCFGEdges(&fe));
- EXPECT_THAT(p->error(),
- Eq("Branch from 10 to 50 bypasses header 20 (dominance rule violated)"));
-}
-
TEST_F(SpvParserCFGTest, ClassifyCFGEdges_Forward_IfToThen) {
auto assembly = CommonTypes() + R"(
%100 = OpFunction %void None %voidfn
@@ -9120,7 +8487,7 @@
ASSERT_EQ(expect, got);
}
-// First do no special control flow: no fallthroughs, breaks, continues.
+// First do no special control flow: no breaks, continues.
TEST_F(SpvParserCFGTest, EmitBody_Switch_DefaultIsMerge_OneCase) {
auto p = parser(test::Assemble(CommonTypes() + R"(
%100 = OpFunction %void None %voidfn
@@ -9307,8 +8674,7 @@
TEST_F(SpvParserCFGTest, EmitBody_Switch_DefaultIsCase_WithDupCase) {
// The default block is not the merge block and is the same as a case.
- // We emit the default case separately, but just before the labeled
- // case, and with a fallthrough.
+ // We emit the default case as part of the labeled case.
auto p = parser(test::Assemble(CommonTypes() + R"(
%100 = OpFunction %void None %voidfn
@@ -10423,53 +9789,6 @@
ASSERT_EQ(expect, got);
}
-TEST_F(SpvParserCFGTest, EmitBody_Branch_Fallthrough) {
- auto p = parser(test::Assemble(CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpStore %var %uint_1
- OpSelectionMerge %99 None
- OpSwitch %selector %99 20 %20 30 %30
-
- %20 = OpLabel
- OpStore %var %uint_20
- OpBranch %30 ; uncondtional fallthrough
-
- %30 = OpLabel
- OpStore %var %uint_30
- OpBranch %99
-
- %99 = OpLabel
- OpStore %var %uint_7
- OpReturn
-
- OpFunctionEnd
- )"));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_TRUE(fe.EmitBody()) << p->error();
-
- auto ast_body = fe.ast_body();
- auto got = test::ToString(p->program(), ast_body);
- auto* expect = R"(var_1 = 1u;
-switch(42u) {
- case 20u: {
- var_1 = 20u;
- fallthrough;
- }
- case 30u: {
- var_1 = 30u;
- }
- default: {
- }
-}
-var_1 = 7u;
-return;
-)";
- ASSERT_EQ(expect, got);
-}
-
TEST_F(SpvParserCFGTest, EmitBody_Branch_Forward) {
auto p = parser(test::Assemble(CommonTypes() + R"(
%100 = OpFunction %void None %voidfn
@@ -10507,7 +9826,6 @@
// If continue is forward, then it's a continue from a
// continue which is also invalid.
// kIfBreak: invalid: loop and if must have distinct merge blocks
-// kCaseFallThrough: invalid: loop header must dominate its merge
// kForward: impossible; would be a loop break
//
// kSwitchBreak with:
@@ -10516,7 +9834,6 @@
// kLoopBreak: invalid; only one kind of break allowed
// kLoopContinue: TESTED
// kIfBreak: invalid: switch and if must have distinct merge blocks
-// kCaseFallThrough: TESTED
// kForward: TESTED
//
// kLoopBreak with:
@@ -10525,7 +9842,6 @@
// kLoopBreak: dup general case
// kLoopContinue: TESTED
// kIfBreak: invalid: switch and if must have distinct merge blocks
-// kCaseFallThrough: not possible, because switch break conflicts with loop
// break kForward: TESTED
//
// kLoopContinue with:
@@ -10534,7 +9850,6 @@
// kLoopBreak: symmetry
// kLoopContinue: dup general case
// kIfBreak: TESTED
-// kCaseFallThrough: TESTED
// kForward: TESTED
//
// kIfBreak with:
@@ -10543,25 +9858,14 @@
// kLoopBreak: symmetry
// kLoopContinue: symmetry
// kIfBreak: dup general case
-// kCaseFallThrough: invalid; violates nesting or unique merges
// kForward: invalid: needs a merge instruction
//
-// kCaseFallThrough with:
-// kBack : symmetry
-// kSwitchBreak: symmetry
-// kLoopBreak: symmetry
-// kLoopContinue: symmetry
-// kIfBreak: symmetry
-// kCaseFallThrough: dup general case
-// kForward: invalid (tested)
-//
// kForward with:
// kBack : symmetry
// kSwitchBreak: symmetry
// kLoopBreak: symmetry
// kLoopContinue: symmetry
// kIfBreak: symmetry
-// kCaseFallThrough: symmetry
// kForward: dup general case
TEST_F(SpvParserCFGTest, EmitBody_BranchConditional_Back_SingleBlock_Back) {
@@ -11085,107 +10389,6 @@
ASSERT_EQ(expect, got);
}
-TEST_F(SpvParserCFGTest, EmitBody_BranchConditional_SwitchBreak_Fallthrough_OnTrue) {
- auto p = parser(test::Assemble(CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpStore %var %uint_1
- OpSelectionMerge %99 None
- OpSwitch %selector %99 20 %20 30 %30
-
- %20 = OpLabel
- OpStore %var %uint_20
- OpBranchConditional %cond %30 %99; fallthrough on true
-
- %30 = OpLabel
- OpStore %var %uint_30
- OpBranch %99
-
- %99 = OpLabel
- OpStore %var %uint_7
- OpReturn
-
- OpFunctionEnd
- )"));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_TRUE(fe.EmitBody()) << p->error();
-
- auto ast_body = fe.ast_body();
- auto got = test::ToString(p->program(), ast_body);
- auto* expect = R"(var_1 = 1u;
-switch(42u) {
- case 20u: {
- var_1 = 20u;
- if (false) {
- } else {
- break;
- }
- fallthrough;
- }
- case 30u: {
- var_1 = 30u;
- }
- default: {
- }
-}
-var_1 = 7u;
-return;
-)";
- ASSERT_EQ(expect, got);
-}
-
-TEST_F(SpvParserCFGTest, EmitBody_BranchConditional_SwitchBreak_Fallthrough_OnFalse) {
- auto p = parser(test::Assemble(CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpStore %var %uint_1
- OpSelectionMerge %99 None
- OpSwitch %selector %99 20 %20 30 %30
-
- %20 = OpLabel
- OpStore %var %uint_20
- OpBranchConditional %cond %99 %30; fallthrough on false
-
- %30 = OpLabel
- OpStore %var %uint_30
- OpBranch %99
-
- %99 = OpLabel
- OpStore %var %uint_7
- OpReturn
-
- OpFunctionEnd
- )"));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_TRUE(fe.EmitBody()) << p->error();
-
- auto ast_body = fe.ast_body();
- auto got = test::ToString(p->program(), ast_body);
- auto* expect = R"(var_1 = 1u;
-switch(42u) {
- case 20u: {
- var_1 = 20u;
- if (false) {
- break;
- }
- fallthrough;
- }
- case 30u: {
- var_1 = 30u;
- }
- default: {
- }
-}
-var_1 = 7u;
-return;
-)";
- ASSERT_EQ(expect, got);
-}
-
TEST_F(SpvParserCFGTest, EmitBody_BranchConditional_LoopBreak_SingleBlock_LoopBreak) {
auto p = parser(test::Assemble(CommonTypes() + R"(
%100 = OpFunction %void None %voidfn
@@ -11411,53 +10614,6 @@
ASSERT_EQ(expect, got);
}
-TEST_F(SpvParserCFGTest, EmitBody_BranchConditional_LoopBreak_Fallthrough_IsError) {
- // It's an error because switch break conflicts with loop break.
- auto p = parser(test::Assemble(CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpStore %var %uint_0
- OpBranch %20
-
- %20 = OpLabel
- OpStore %var %uint_1
- OpLoopMerge %99 %80 None
- OpBranch %30
-
- %30 = OpLabel
- OpSelectionMerge %79 None
- OpSwitch %selector %79 40 %40 50 %50
-
- %40 = OpLabel
- OpStore %var %uint_40
- ; error: branch to 99 bypasses switch's merge
- OpBranchConditional %cond %99 %50 ; loop break; fall through
-
- %50 = OpLabel
- OpStore %var %uint_50
- OpBranch %79
-
- %79 = OpLabel ; switch merge
- OpBranch %80
-
- %80 = OpLabel ; continue target
- OpStore %var %uint_4
- OpBranch %20
-
- %99 = OpLabel
- OpStore %var %uint_5
- OpReturn
-
- OpFunctionEnd
- )"));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_FALSE(fe.EmitBody()) << p->error();
- EXPECT_THAT(p->error(), Eq("Branch from block 40 to block 99 is an invalid exit from construct "
- "starting at block 30; branch bypasses merge block 79"));
-}
-
TEST_F(SpvParserCFGTest, EmitBody_BranchConditional_LoopBreak_Forward_OnTrue) {
auto p = parser(test::Assemble(CommonTypes() + R"(
%100 = OpFunction %void None %voidfn
@@ -11662,7 +10818,6 @@
TEST_F(SpvParserCFGTest, EmitBody_BranchConditional_Continue_Continue_AfterHeader_Conditional) {
// Create an intervening block so we actually require a "continue" statement
- // instead of just an adjacent fallthrough to the continue target.
auto p = parser(test::Assemble(CommonTypes() + R"(
%100 = OpFunction %void None %voidfn
@@ -11977,157 +11132,6 @@
ASSERT_EQ(expect, got);
}
-TEST_F(SpvParserCFGTest, EmitBody_BranchConditional_Continue_Fallthrough_OnTrue) {
- auto p = parser(test::Assemble(CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpStore %var %uint_0
- OpBranch %20
-
- %20 = OpLabel
- OpStore %var %uint_1
- OpLoopMerge %99 %80 None
- OpBranch %30
-
- %30 = OpLabel
- OpStore %var %uint_2
- OpSelectionMerge %79 None
- OpSwitch %selector %79 40 %40 50 %50
-
- %40 = OpLabel
- OpStore %var %uint_40
- OpBranchConditional %cond %50 %80 ; loop continue; fall through on true
-
- %50 = OpLabel
- OpStore %var %uint_50
- OpBranch %79
-
- %79 = OpLabel ; switch merge
- OpStore %var %uint_3
- OpBranch %80
-
- %80 = OpLabel ; continue target
- OpStore %var %uint_4
- OpBranch %20
-
- %99 = OpLabel
- OpStore %var %uint_5
- OpReturn
-
- OpFunctionEnd
- )"));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_TRUE(fe.EmitBody()) << p->error();
- auto ast_body = fe.ast_body();
- auto got = test::ToString(p->program(), ast_body);
- auto* expect = R"(var_1 = 0u;
-loop {
- var_1 = 1u;
- var_1 = 2u;
- switch(42u) {
- case 40u: {
- var_1 = 40u;
- if (false) {
- } else {
- continue;
- }
- fallthrough;
- }
- case 50u: {
- var_1 = 50u;
- }
- default: {
- }
- }
- var_1 = 3u;
-
- continuing {
- var_1 = 4u;
- }
-}
-var_1 = 5u;
-return;
-)";
- ASSERT_EQ(expect, got);
-}
-
-TEST_F(SpvParserCFGTest, EmitBody_BranchConditional_Continue_Fallthrough_OnFalse) {
- auto p = parser(test::Assemble(CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpStore %var %uint_0
- OpBranch %20
-
- %20 = OpLabel
- OpStore %var %uint_1
- OpLoopMerge %99 %80 None
- OpBranch %30
-
- %30 = OpLabel
- OpStore %var %uint_2
- OpSelectionMerge %79 None
- OpSwitch %selector %79 40 %40 50 %50
-
- %40 = OpLabel
- OpStore %var %uint_40
- OpBranchConditional %cond %80 %50 ; loop continue; fall through on false
-
- %50 = OpLabel
- OpStore %var %uint_50
- OpBranch %79
-
- %79 = OpLabel ; switch merge
- OpStore %var %uint_3
- OpBranch %80
-
- %80 = OpLabel ; continue target
- OpStore %var %uint_4
- OpBranch %20
-
- %99 = OpLabel
- OpStore %var %uint_5
- OpReturn
-
- OpFunctionEnd
- )"));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_TRUE(fe.EmitBody()) << p->error();
- auto ast_body = fe.ast_body();
- auto got = test::ToString(p->program(), ast_body);
- auto* expect = R"(var_1 = 0u;
-loop {
- var_1 = 1u;
- var_1 = 2u;
- switch(42u) {
- case 40u: {
- var_1 = 40u;
- if (false) {
- continue;
- }
- fallthrough;
- }
- case 50u: {
- var_1 = 50u;
- }
- default: {
- }
- }
- var_1 = 3u;
-
- continuing {
- var_1 = 4u;
- }
-}
-var_1 = 5u;
-return;
-)";
- ASSERT_EQ(expect, got);
-}
-
TEST_F(SpvParserCFGTest, EmitBody_BranchConditional_Continue_Forward_OnTrue) {
auto p = parser(test::Assemble(CommonTypes() + R"(
%100 = OpFunction %void None %voidfn
@@ -12308,91 +11312,6 @@
"starting at block 20; branch bypasses merge block 89"));
}
-TEST_F(SpvParserCFGTest, EmitBody_BranchConditional_Fallthrough_Fallthrough_Same) {
- auto p = parser(test::Assemble(CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpStore %var %uint_1
- OpSelectionMerge %99 None
- OpSwitch %selector %99 20 %20 30 %30
-
- %20 = OpLabel
- OpStore %var %uint_20
- OpBranchConditional %cond %30 %30 ; fallthrough fallthrough
-
- %30 = OpLabel
- OpStore %var %uint_30
- OpBranch %99
-
- %99 = OpLabel
- OpStore %var %uint_7
- OpReturn
-
- OpFunctionEnd
- )"));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_TRUE(fe.EmitBody()) << p->error();
-
- auto ast_body = fe.ast_body();
- auto got = test::ToString(p->program(), ast_body);
- auto* expect = R"(var_1 = 1u;
-switch(42u) {
- case 20u: {
- var_1 = 20u;
- fallthrough;
- }
- case 30u: {
- var_1 = 30u;
- }
- default: {
- }
-}
-var_1 = 7u;
-return;
-)";
- ASSERT_EQ(expect, got);
-}
-
-TEST_F(SpvParserCFGTest, EmitBody_BranchConditional_Fallthrough_NotLastInCase_IsError) {
- // See also
- // ClassifyCFGEdges_Fallthrough_BranchConditionalWith_Forward_IsError.
- auto p = parser(test::Assemble(CommonTypes() + R"(
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpSelectionMerge %99 None
- OpSwitch %selector %99 20 %20 40 %40
-
- %20 = OpLabel ; case 30
- OpSelectionMerge %39 None
- OpBranchConditional %cond %40 %30 ; fallthrough and forward
-
- %30 = OpLabel
- OpBranch %39
-
- %39 = OpLabel
- OpBranch %99
-
- %40 = OpLabel ; case 40
- OpBranch %99
-
- %99 = OpLabel
- OpReturn
-
- OpFunctionEnd
- )"));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << p->error();
- auto fe = p->function_emitter(100);
- EXPECT_FALSE(fe.EmitBody());
- // The weird forward branch pulls in 40 as part of the selection rather than
- // as a case.
- EXPECT_THAT(fe.block_order(), ElementsAre(10, 20, 40, 30, 39, 99));
- EXPECT_THAT(p->error(),
- Eq("Branch from 10 to 40 bypasses header 20 (dominance rule violated)"));
-}
-
TEST_F(SpvParserCFGTest, EmitBody_BranchConditional_Forward_Forward_Same) {
auto p = parser(test::Assemble(CommonTypes() + R"(
%100 = OpFunction %void None %voidfn
diff --git a/src/tint/reader/spirv/function_var_test.cc b/src/tint/reader/spirv/function_var_test.cc
index e3d2acf..5b83105 100644
--- a/src/tint/reader/spirv/function_var_test.cc
+++ b/src/tint/reader/spirv/function_var_test.cc
@@ -1344,73 +1344,6 @@
EXPECT_EQ(expect, got) << got;
}
-TEST_F(SpvParserFunctionVarTest, EmitStatement_Phi_InMerge_PredecessorsDominatdByNestedSwitchCase) {
- // This is the essence of the bug report from crbug.com/tint/495
- auto assembly = Preamble() + R"(
- %cond = OpConstantTrue %bool
- %pty = OpTypePointer Private %uint
- %1 = OpVariable %pty Private
- %boolpty = OpTypePointer Private %bool
- %7 = OpVariable %boolpty Private
- %8 = OpVariable %boolpty Private
-
- %100 = OpFunction %void None %voidfn
-
- %10 = OpLabel
- OpSelectionMerge %99 None
- OpSwitch %uint_1 %20 0 %20 1 %30
-
- %20 = OpLabel ; case 0
- OpBranch %30 ;; fall through
-
- %30 = OpLabel ; case 1
- OpSelectionMerge %50 None
- OpBranchConditional %true %40 %45
-
- %40 = OpLabel
- OpBranch %50
-
- %45 = OpLabel
- OpBranch %99 ; break
-
- %50 = OpLabel ; end the case
- OpBranch %99
-
- %99 = OpLabel
- ; predecessors are all dominated by case construct head at %30
- %41 = OpPhi %uint %uint_0 %45 %uint_1 %50
- %101 = OpCopyObject %uint %41 ; give it a use so it's emitted
- OpReturn
-
- OpFunctionEnd
- )";
- auto p = parser(test::Assemble(assembly));
- ASSERT_TRUE(p->BuildAndParseInternalModuleExceptFunctions()) << assembly;
- auto fe = p->function_emitter(100);
- EXPECT_TRUE(fe.EmitBody()) << p->error();
-
- auto ast_body = fe.ast_body();
- auto got = test::ToString(p->program(), ast_body);
- auto* expect = R"(var x_41 : u32;
-switch(1u) {
- case 0u, default: {
- fallthrough;
- }
- case 1u: {
- if (true) {
- } else {
- x_41 = 0u;
- break;
- }
- x_41 = 1u;
- }
-}
-let x_101 : u32 = x_41;
-return;
-)";
- EXPECT_EQ(expect, got) << got << assembly;
-}
-
TEST_F(SpvParserFunctionVarTest, EmitStatement_Phi_UseInPhiCountsAsUse) {
// From crbug.com/215
// If the only use of a combinatorially computed ID is as the value
diff --git a/test/tint/statements/switch/fallthrough.wgsl b/test/tint/statements/switch/fallthrough.wgsl
deleted file mode 100644
index 667a8f2..0000000
--- a/test/tint/statements/switch/fallthrough.wgsl
+++ /dev/null
@@ -1,12 +0,0 @@
-@compute @workgroup_size(1)
-fn f() {
- var i : i32;
- switch(i) {
- case 0: {
- fallthrough;
- }
- default: {
- break;
- }
- }
-}
diff --git a/test/tint/statements/switch/fallthrough.wgsl.expected.dxc.hlsl b/test/tint/statements/switch/fallthrough.wgsl.expected.dxc.hlsl
deleted file mode 100644
index de066b7..0000000
--- a/test/tint/statements/switch/fallthrough.wgsl.expected.dxc.hlsl
+++ /dev/null
@@ -1,20 +0,0 @@
-statements/switch/fallthrough.wgsl:6:13 warning: use of deprecated language feature: fallthrough is set to be removed from WGSL. Case can accept multiple selectors if the existing case bodies are empty. (e.g. `case 1, 2, 3:`) `default` is a valid case selector value. (e.g. `case 1, default:`)
- fallthrough;
- ^^^^^^^^^^^
-
-[numthreads(1, 1, 1)]
-void f() {
- int i = 0;
- switch(i) {
- case 0: {
- /* fallthrough */
- {
- break;
- }
- }
- default: {
- break;
- }
- }
- return;
-}
diff --git a/test/tint/statements/switch/fallthrough.wgsl.expected.fxc.hlsl b/test/tint/statements/switch/fallthrough.wgsl.expected.fxc.hlsl
deleted file mode 100644
index de066b7..0000000
--- a/test/tint/statements/switch/fallthrough.wgsl.expected.fxc.hlsl
+++ /dev/null
@@ -1,20 +0,0 @@
-statements/switch/fallthrough.wgsl:6:13 warning: use of deprecated language feature: fallthrough is set to be removed from WGSL. Case can accept multiple selectors if the existing case bodies are empty. (e.g. `case 1, 2, 3:`) `default` is a valid case selector value. (e.g. `case 1, default:`)
- fallthrough;
- ^^^^^^^^^^^
-
-[numthreads(1, 1, 1)]
-void f() {
- int i = 0;
- switch(i) {
- case 0: {
- /* fallthrough */
- {
- break;
- }
- }
- default: {
- break;
- }
- }
- return;
-}
diff --git a/test/tint/statements/switch/fallthrough.wgsl.expected.glsl b/test/tint/statements/switch/fallthrough.wgsl.expected.glsl
deleted file mode 100644
index 6ca3be4..0000000
--- a/test/tint/statements/switch/fallthrough.wgsl.expected.glsl
+++ /dev/null
@@ -1,23 +0,0 @@
-statements/switch/fallthrough.wgsl:6:13 warning: use of deprecated language feature: fallthrough is set to be removed from WGSL. Case can accept multiple selectors if the existing case bodies are empty. (e.g. `case 1, 2, 3:`) `default` is a valid case selector value. (e.g. `case 1, default:`)
- fallthrough;
- ^^^^^^^^^^^
-
-#version 310 es
-
-void f() {
- int i = 0;
- switch(i) {
- case 0: {
- /* fallthrough */
- }
- default: {
- break;
- }
- }
-}
-
-layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
-void main() {
- f();
- return;
-}
diff --git a/test/tint/statements/switch/fallthrough.wgsl.expected.msl b/test/tint/statements/switch/fallthrough.wgsl.expected.msl
deleted file mode 100644
index d2cd679..0000000
--- a/test/tint/statements/switch/fallthrough.wgsl.expected.msl
+++ /dev/null
@@ -1,20 +0,0 @@
-statements/switch/fallthrough.wgsl:6:13 warning: use of deprecated language feature: fallthrough is set to be removed from WGSL. Case can accept multiple selectors if the existing case bodies are empty. (e.g. `case 1, 2, 3:`) `default` is a valid case selector value. (e.g. `case 1, default:`)
- fallthrough;
- ^^^^^^^^^^^
-
-#include <metal_stdlib>
-
-using namespace metal;
-kernel void f() {
- int i = 0;
- switch(i) {
- case 0: {
- /* fallthrough */
- }
- default: {
- break;
- }
- }
- return;
-}
-
diff --git a/test/tint/statements/switch/fallthrough.wgsl.expected.spvasm b/test/tint/statements/switch/fallthrough.wgsl.expected.spvasm
deleted file mode 100644
index dc8c063..0000000
--- a/test/tint/statements/switch/fallthrough.wgsl.expected.spvasm
+++ /dev/null
@@ -1,33 +0,0 @@
-statements/switch/fallthrough.wgsl:6:13 warning: use of deprecated language feature: fallthrough is set to be removed from WGSL. Case can accept multiple selectors if the existing case bodies are empty. (e.g. `case 1, 2, 3:`) `default` is a valid case selector value. (e.g. `case 1, default:`)
- fallthrough;
- ^^^^^^^^^^^
-
-; SPIR-V
-; Version: 1.3
-; Generator: Google Tint Compiler; 0
-; Bound: 13
-; Schema: 0
- OpCapability Shader
- OpMemoryModel Logical GLSL450
- OpEntryPoint GLCompute %f "f"
- OpExecutionMode %f LocalSize 1 1 1
- OpName %f "f"
- OpName %i "i"
- %void = OpTypeVoid
- %1 = OpTypeFunction %void
- %int = OpTypeInt 32 1
-%_ptr_Function_int = OpTypePointer Function %int
- %8 = OpConstantNull %int
- %f = OpFunction %void None %1
- %4 = OpLabel
- %i = OpVariable %_ptr_Function_int Function %8
- %10 = OpLoad %int %i
- OpSelectionMerge %9 None
- OpSwitch %10 %11 0 %12
- %12 = OpLabel
- OpBranch %11
- %11 = OpLabel
- OpBranch %9
- %9 = OpLabel
- OpReturn
- OpFunctionEnd
diff --git a/test/tint/statements/switch/fallthrough.wgsl.expected.wgsl b/test/tint/statements/switch/fallthrough.wgsl.expected.wgsl
deleted file mode 100644
index 50254ab..0000000
--- a/test/tint/statements/switch/fallthrough.wgsl.expected.wgsl
+++ /dev/null
@@ -1,16 +0,0 @@
-statements/switch/fallthrough.wgsl:6:13 warning: use of deprecated language feature: fallthrough is set to be removed from WGSL. Case can accept multiple selectors if the existing case bodies are empty. (e.g. `case 1, 2, 3:`) `default` is a valid case selector value. (e.g. `case 1, default:`)
- fallthrough;
- ^^^^^^^^^^^
-
-@compute @workgroup_size(1)
-fn f() {
- var i : i32;
- switch(i) {
- case 0: {
- fallthrough;
- }
- default: {
- break;
- }
- }
-}
diff --git a/test/tint/unittest/reader/spirv/SpvParserCFGTest_ClassifyCFGEdges_Fallthrough_CaseTailToCase.spvasm b/test/tint/unittest/reader/spirv/SpvParserCFGTest_ClassifyCFGEdges_Fallthrough_CaseTailToCase.spvasm
deleted file mode 100644
index bf5da38..0000000
--- a/test/tint/unittest/reader/spirv/SpvParserCFGTest_ClassifyCFGEdges_Fallthrough_CaseTailToCase.spvasm
+++ /dev/null
@@ -1,55 +0,0 @@
-; Test: SpvParserCFGTest_ClassifyCFGEdges_Fallthrough_CaseTailToCase.spvasm
-; SPIR-V
-; Version: 1.0
-; Generator: Khronos SPIR-V Tools Assembler; 0
-; Bound: 1000
-; Schema: 0
-OpCapability Shader
-OpMemoryModel Logical Simple
-OpEntryPoint Fragment %100 "main"
-OpExecutionMode %100 OriginUpperLeft
-OpName %var "var"
-%void = OpTypeVoid
-%3 = OpTypeFunction %void
-%bool = OpTypeBool
-%5 = OpConstantNull %bool
-%true = OpConstantTrue %bool
-%false = OpConstantFalse %bool
-%uint = OpTypeInt 32 0
-%int = OpTypeInt 32 1
-%uint_42 = OpConstant %uint 42
-%int_42 = OpConstant %int 42
-%13 = OpTypeFunction %uint
-%uint_0 = OpConstant %uint 0
-%uint_1 = OpConstant %uint 1
-%uint_2 = OpConstant %uint 2
-%uint_3 = OpConstant %uint 3
-%uint_4 = OpConstant %uint 4
-%uint_5 = OpConstant %uint 5
-%uint_6 = OpConstant %uint 6
-%uint_7 = OpConstant %uint 7
-%uint_8 = OpConstant %uint 8
-%uint_10 = OpConstant %uint 10
-%uint_20 = OpConstant %uint 20
-%uint_30 = OpConstant %uint 30
-%uint_40 = OpConstant %uint 40
-%uint_50 = OpConstant %uint 50
-%uint_90 = OpConstant %uint 90
-%uint_99 = OpConstant %uint 99
-%_ptr_Private_uint = OpTypePointer Private %uint
-%var = OpVariable %_ptr_Private_uint Private
-%uint_999 = OpConstant %uint 999
-%100 = OpFunction %void None %3
-%10 = OpLabel
-OpSelectionMerge %99 None
-OpSwitch %uint_42 %99 20 %20 40 %40
-%20 = OpLabel
-OpBranch %30
-%30 = OpLabel
-OpBranch %40
-%40 = OpLabel
-OpBranch %99
-%99 = OpLabel
-OpReturn
-OpFunctionEnd
-
diff --git a/test/tint/unittest/reader/spirv/SpvParserCFGTest_ClassifyCFGEdges_Fallthrough_CaseTailToDefaultNotMerge.spvasm b/test/tint/unittest/reader/spirv/SpvParserCFGTest_ClassifyCFGEdges_Fallthrough_CaseTailToDefaultNotMerge.spvasm
deleted file mode 100644
index e3f3d8d..0000000
--- a/test/tint/unittest/reader/spirv/SpvParserCFGTest_ClassifyCFGEdges_Fallthrough_CaseTailToDefaultNotMerge.spvasm
+++ /dev/null
@@ -1,55 +0,0 @@
-; Test: SpvParserCFGTest_ClassifyCFGEdges_Fallthrough_CaseTailToDefaultNotMerge.spvasm
-; SPIR-V
-; Version: 1.0
-; Generator: Khronos SPIR-V Tools Assembler; 0
-; Bound: 1000
-; Schema: 0
-OpCapability Shader
-OpMemoryModel Logical Simple
-OpEntryPoint Fragment %100 "main"
-OpExecutionMode %100 OriginUpperLeft
-OpName %var "var"
-%void = OpTypeVoid
-%3 = OpTypeFunction %void
-%bool = OpTypeBool
-%5 = OpConstantNull %bool
-%true = OpConstantTrue %bool
-%false = OpConstantFalse %bool
-%uint = OpTypeInt 32 0
-%int = OpTypeInt 32 1
-%uint_42 = OpConstant %uint 42
-%int_42 = OpConstant %int 42
-%13 = OpTypeFunction %uint
-%uint_0 = OpConstant %uint 0
-%uint_1 = OpConstant %uint 1
-%uint_2 = OpConstant %uint 2
-%uint_3 = OpConstant %uint 3
-%uint_4 = OpConstant %uint 4
-%uint_5 = OpConstant %uint 5
-%uint_6 = OpConstant %uint 6
-%uint_7 = OpConstant %uint 7
-%uint_8 = OpConstant %uint 8
-%uint_10 = OpConstant %uint 10
-%uint_20 = OpConstant %uint 20
-%uint_30 = OpConstant %uint 30
-%uint_40 = OpConstant %uint 40
-%uint_50 = OpConstant %uint 50
-%uint_90 = OpConstant %uint 90
-%uint_99 = OpConstant %uint 99
-%_ptr_Private_uint = OpTypePointer Private %uint
-%var = OpVariable %_ptr_Private_uint Private
-%uint_999 = OpConstant %uint 999
-%100 = OpFunction %void None %3
-%10 = OpLabel
-OpSelectionMerge %99 None
-OpSwitch %uint_42 %40 20 %20
-%20 = OpLabel
-OpBranch %30
-%30 = OpLabel
-OpBranch %40
-%40 = OpLabel
-OpBranch %99
-%99 = OpLabel
-OpReturn
-OpFunctionEnd
-
diff --git a/test/tint/unittest/reader/spirv/SpvParserCFGTest_ClassifyCFGEdges_Fallthrough_DefaultToCase.spvasm b/test/tint/unittest/reader/spirv/SpvParserCFGTest_ClassifyCFGEdges_Fallthrough_DefaultToCase.spvasm
deleted file mode 100644
index 5ab7265..0000000
--- a/test/tint/unittest/reader/spirv/SpvParserCFGTest_ClassifyCFGEdges_Fallthrough_DefaultToCase.spvasm
+++ /dev/null
@@ -1,55 +0,0 @@
-; Test: SpvParserCFGTest_ClassifyCFGEdges_Fallthrough_DefaultToCase.spvasm
-; SPIR-V
-; Version: 1.0
-; Generator: Khronos SPIR-V Tools Assembler; 0
-; Bound: 1000
-; Schema: 0
-OpCapability Shader
-OpMemoryModel Logical Simple
-OpEntryPoint Fragment %100 "main"
-OpExecutionMode %100 OriginUpperLeft
-OpName %var "var"
-%void = OpTypeVoid
-%3 = OpTypeFunction %void
-%bool = OpTypeBool
-%5 = OpConstantNull %bool
-%true = OpConstantTrue %bool
-%false = OpConstantFalse %bool
-%uint = OpTypeInt 32 0
-%int = OpTypeInt 32 1
-%uint_42 = OpConstant %uint 42
-%int_42 = OpConstant %int 42
-%13 = OpTypeFunction %uint
-%uint_0 = OpConstant %uint 0
-%uint_1 = OpConstant %uint 1
-%uint_2 = OpConstant %uint 2
-%uint_3 = OpConstant %uint 3
-%uint_4 = OpConstant %uint 4
-%uint_5 = OpConstant %uint 5
-%uint_6 = OpConstant %uint 6
-%uint_7 = OpConstant %uint 7
-%uint_8 = OpConstant %uint 8
-%uint_10 = OpConstant %uint 10
-%uint_20 = OpConstant %uint 20
-%uint_30 = OpConstant %uint 30
-%uint_40 = OpConstant %uint 40
-%uint_50 = OpConstant %uint 50
-%uint_90 = OpConstant %uint 90
-%uint_99 = OpConstant %uint 99
-%_ptr_Private_uint = OpTypePointer Private %uint
-%var = OpVariable %_ptr_Private_uint Private
-%uint_999 = OpConstant %uint 999
-%100 = OpFunction %void None %3
-%10 = OpLabel
-OpSelectionMerge %99 None
-OpSwitch %uint_42 %20 40 %40
-%20 = OpLabel
-OpBranch %30
-%30 = OpLabel
-OpBranch %40
-%40 = OpLabel
-OpBranch %99
-%99 = OpLabel
-OpReturn
-OpFunctionEnd
-
diff --git a/test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_Nest_IfFallthrough_In_SwitchCase.spvasm b/test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_Nest_IfFallthrough_In_SwitchCase.spvasm
deleted file mode 100644
index 87bdbc8..0000000
--- a/test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_Nest_IfFallthrough_In_SwitchCase.spvasm
+++ /dev/null
@@ -1,67 +0,0 @@
-; Test: SpvParserCFGTest_ComputeBlockOrder_Nest_IfFallthrough_In_SwitchCase.spvasm
-; SPIR-V
-; Version: 1.0
-; Generator: Khronos SPIR-V Tools Assembler; 0
-; Bound: 1000
-; Schema: 0
-OpCapability Shader
-OpMemoryModel Logical Simple
-OpEntryPoint Fragment %100 "main"
-OpExecutionMode %100 OriginUpperLeft
-OpName %var "var"
-%void = OpTypeVoid
-%3 = OpTypeFunction %void
-%bool = OpTypeBool
-%5 = OpConstantNull %bool
-%true = OpConstantTrue %bool
-%false = OpConstantFalse %bool
-%uint = OpTypeInt 32 0
-%int = OpTypeInt 32 1
-%uint_42 = OpConstant %uint 42
-%int_42 = OpConstant %int 42
-%13 = OpTypeFunction %uint
-%uint_0 = OpConstant %uint 0
-%uint_1 = OpConstant %uint 1
-%uint_2 = OpConstant %uint 2
-%uint_3 = OpConstant %uint 3
-%uint_4 = OpConstant %uint 4
-%uint_5 = OpConstant %uint 5
-%uint_6 = OpConstant %uint 6
-%uint_7 = OpConstant %uint 7
-%uint_8 = OpConstant %uint 8
-%uint_10 = OpConstant %uint 10
-%uint_20 = OpConstant %uint 20
-%uint_30 = OpConstant %uint 30
-%uint_40 = OpConstant %uint 40
-%uint_50 = OpConstant %uint 50
-%uint_90 = OpConstant %uint 90
-%uint_99 = OpConstant %uint 99
-%_ptr_Private_uint = OpTypePointer Private %uint
-%var = OpVariable %_ptr_Private_uint Private
-%uint_999 = OpConstant %uint 999
-%100 = OpFunction %void None %3
-%10 = OpLabel
-OpSelectionMerge %99 None
-OpSwitch %uint_42 %50 20 %20 50 %50
-%99 = OpLabel
-OpReturn
-%20 = OpLabel
-OpSelectionMerge %49 None
-OpBranchConditional %5 %30 %40
-%49 = OpLabel
-OpBranchConditional %5 %99 %50
-%30 = OpLabel
-OpBranch %49
-%40 = OpLabel
-OpBranch %49
-%50 = OpLabel
-OpSelectionMerge %79 None
-OpBranchConditional %5 %60 %70
-%79 = OpLabel
-OpBranch %99
-%60 = OpLabel
-OpBranch %79
-%70 = OpLabel
-OpBranch %79
-OpFunctionEnd
-
diff --git a/test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough.spvasm b/test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough.spvasm
deleted file mode 100644
index 92feaa3..0000000
--- a/test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough.spvasm
+++ /dev/null
@@ -1,57 +0,0 @@
-; Test: SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough.spvasm
-; SPIR-V
-; Version: 1.0
-; Generator: Khronos SPIR-V Tools Assembler; 0
-; Bound: 1000
-; Schema: 0
-OpCapability Shader
-OpMemoryModel Logical Simple
-OpEntryPoint Fragment %100 "main"
-OpExecutionMode %100 OriginUpperLeft
-OpName %var "var"
-%void = OpTypeVoid
-%3 = OpTypeFunction %void
-%bool = OpTypeBool
-%5 = OpConstantNull %bool
-%true = OpConstantTrue %bool
-%false = OpConstantFalse %bool
-%uint = OpTypeInt 32 0
-%int = OpTypeInt 32 1
-%uint_42 = OpConstant %uint 42
-%int_42 = OpConstant %int 42
-%13 = OpTypeFunction %uint
-%uint_0 = OpConstant %uint 0
-%uint_1 = OpConstant %uint 1
-%uint_2 = OpConstant %uint 2
-%uint_3 = OpConstant %uint 3
-%uint_4 = OpConstant %uint 4
-%uint_5 = OpConstant %uint 5
-%uint_6 = OpConstant %uint 6
-%uint_7 = OpConstant %uint 7
-%uint_8 = OpConstant %uint 8
-%uint_10 = OpConstant %uint 10
-%uint_20 = OpConstant %uint 20
-%uint_30 = OpConstant %uint 30
-%uint_40 = OpConstant %uint 40
-%uint_50 = OpConstant %uint 50
-%uint_90 = OpConstant %uint 90
-%uint_99 = OpConstant %uint 99
-%_ptr_Private_uint = OpTypePointer Private %uint
-%var = OpVariable %_ptr_Private_uint Private
-%uint_999 = OpConstant %uint 999
-%100 = OpFunction %void None %3
-%10 = OpLabel
-OpSelectionMerge %99 None
-OpSwitch %uint_42 %99 20 %20 40 %40 30 %30 50 %50
-%50 = OpLabel
-OpBranch %99
-%99 = OpLabel
-OpReturn
-%40 = OpLabel
-OpBranch %99
-%30 = OpLabel
-OpBranch %50
-%20 = OpLabel
-OpBranch %40
-OpFunctionEnd
-
diff --git a/test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough_FromCaseToDefaultToCase.spvasm b/test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough_FromCaseToDefaultToCase.spvasm
deleted file mode 100644
index a74bcc1..0000000
--- a/test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough_FromCaseToDefaultToCase.spvasm
+++ /dev/null
@@ -1,55 +0,0 @@
-; Test: SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough_FromCaseToDefaultToCase.spvasm
-; SPIR-V
-; Version: 1.0
-; Generator: Khronos SPIR-V Tools Assembler; 0
-; Bound: 1000
-; Schema: 0
-OpCapability Shader
-OpMemoryModel Logical Simple
-OpEntryPoint Fragment %100 "main"
-OpExecutionMode %100 OriginUpperLeft
-OpName %var "var"
-%void = OpTypeVoid
-%3 = OpTypeFunction %void
-%bool = OpTypeBool
-%5 = OpConstantNull %bool
-%true = OpConstantTrue %bool
-%false = OpConstantFalse %bool
-%uint = OpTypeInt 32 0
-%int = OpTypeInt 32 1
-%uint_42 = OpConstant %uint 42
-%int_42 = OpConstant %int 42
-%13 = OpTypeFunction %uint
-%uint_0 = OpConstant %uint 0
-%uint_1 = OpConstant %uint 1
-%uint_2 = OpConstant %uint 2
-%uint_3 = OpConstant %uint 3
-%uint_4 = OpConstant %uint 4
-%uint_5 = OpConstant %uint 5
-%uint_6 = OpConstant %uint 6
-%uint_7 = OpConstant %uint 7
-%uint_8 = OpConstant %uint 8
-%uint_10 = OpConstant %uint 10
-%uint_20 = OpConstant %uint 20
-%uint_30 = OpConstant %uint 30
-%uint_40 = OpConstant %uint 40
-%uint_50 = OpConstant %uint 50
-%uint_90 = OpConstant %uint 90
-%uint_99 = OpConstant %uint 99
-%_ptr_Private_uint = OpTypePointer Private %uint
-%var = OpVariable %_ptr_Private_uint Private
-%uint_999 = OpConstant %uint 999
-%100 = OpFunction %void None %3
-%10 = OpLabel
-OpSelectionMerge %99 None
-OpSwitch %uint_42 %80 20 %20 30 %30
-%20 = OpLabel
-OpBranch %80
-%80 = OpLabel
-OpBranch %30
-%30 = OpLabel
-OpBranch %99
-%99 = OpLabel
-OpReturn
-OpFunctionEnd
-
diff --git a/test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough_FromDefault.spvasm b/test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough_FromDefault.spvasm
deleted file mode 100644
index d6e1cb3..0000000
--- a/test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough_FromDefault.spvasm
+++ /dev/null
@@ -1,57 +0,0 @@
-; Test: SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough_FromDefault.spvasm
-; SPIR-V
-; Version: 1.0
-; Generator: Khronos SPIR-V Tools Assembler; 0
-; Bound: 1000
-; Schema: 0
-OpCapability Shader
-OpMemoryModel Logical Simple
-OpEntryPoint Fragment %100 "main"
-OpExecutionMode %100 OriginUpperLeft
-OpName %var "var"
-%void = OpTypeVoid
-%3 = OpTypeFunction %void
-%bool = OpTypeBool
-%5 = OpConstantNull %bool
-%true = OpConstantTrue %bool
-%false = OpConstantFalse %bool
-%uint = OpTypeInt 32 0
-%int = OpTypeInt 32 1
-%uint_42 = OpConstant %uint 42
-%int_42 = OpConstant %int 42
-%13 = OpTypeFunction %uint
-%uint_0 = OpConstant %uint 0
-%uint_1 = OpConstant %uint 1
-%uint_2 = OpConstant %uint 2
-%uint_3 = OpConstant %uint 3
-%uint_4 = OpConstant %uint 4
-%uint_5 = OpConstant %uint 5
-%uint_6 = OpConstant %uint 6
-%uint_7 = OpConstant %uint 7
-%uint_8 = OpConstant %uint 8
-%uint_10 = OpConstant %uint 10
-%uint_20 = OpConstant %uint 20
-%uint_30 = OpConstant %uint 30
-%uint_40 = OpConstant %uint 40
-%uint_50 = OpConstant %uint 50
-%uint_90 = OpConstant %uint 90
-%uint_99 = OpConstant %uint 99
-%_ptr_Private_uint = OpTypePointer Private %uint
-%var = OpVariable %_ptr_Private_uint Private
-%uint_999 = OpConstant %uint 999
-%100 = OpFunction %void None %3
-%10 = OpLabel
-OpSelectionMerge %99 None
-OpSwitch %uint_42 %80 20 %20 30 %30 40 %40
-%80 = OpLabel
-OpBranch %30
-%99 = OpLabel
-OpReturn
-%40 = OpLabel
-OpBranch %99
-%30 = OpLabel
-OpBranch %40
-%20 = OpLabel
-OpBranch %99
-OpFunctionEnd
-
diff --git a/test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough_Interleaved.spvasm b/test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough_Interleaved.spvasm
deleted file mode 100644
index 3f26450..0000000
--- a/test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough_Interleaved.spvasm
+++ /dev/null
@@ -1,61 +0,0 @@
-; Test: SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough_Interleaved.spvasm
-; SPIR-V
-; Version: 1.0
-; Generator: Khronos SPIR-V Tools Assembler; 0
-; Bound: 1000
-; Schema: 0
-OpCapability Shader
-OpMemoryModel Logical Simple
-OpEntryPoint Fragment %100 "main"
-OpExecutionMode %100 OriginUpperLeft
-OpName %var "var"
-%void = OpTypeVoid
-%3 = OpTypeFunction %void
-%bool = OpTypeBool
-%5 = OpConstantNull %bool
-%true = OpConstantTrue %bool
-%false = OpConstantFalse %bool
-%uint = OpTypeInt 32 0
-%int = OpTypeInt 32 1
-%uint_42 = OpConstant %uint 42
-%int_42 = OpConstant %int 42
-%13 = OpTypeFunction %uint
-%uint_0 = OpConstant %uint 0
-%uint_1 = OpConstant %uint 1
-%uint_2 = OpConstant %uint 2
-%uint_3 = OpConstant %uint 3
-%uint_4 = OpConstant %uint 4
-%uint_5 = OpConstant %uint 5
-%uint_6 = OpConstant %uint 6
-%uint_7 = OpConstant %uint 7
-%uint_8 = OpConstant %uint 8
-%uint_10 = OpConstant %uint 10
-%uint_20 = OpConstant %uint 20
-%uint_30 = OpConstant %uint 30
-%uint_40 = OpConstant %uint 40
-%uint_50 = OpConstant %uint 50
-%uint_90 = OpConstant %uint 90
-%uint_99 = OpConstant %uint 99
-%_ptr_Private_uint = OpTypePointer Private %uint
-%var = OpVariable %_ptr_Private_uint Private
-%uint_999 = OpConstant %uint 999
-%100 = OpFunction %void None %3
-%10 = OpLabel
-OpSelectionMerge %99 None
-OpSwitch %uint_42 %99 20 %20 40 %40 30 %30 50 %50
-%99 = OpLabel
-OpReturn
-%20 = OpLabel
-OpBranch %40
-%30 = OpLabel
-OpBranch %50
-%40 = OpLabel
-OpBranch %60
-%50 = OpLabel
-OpBranch %70
-%60 = OpLabel
-OpBranch %99
-%70 = OpLabel
-OpBranch %99
-OpFunctionEnd
-
diff --git a/test/tint/unittest/reader/spirv/SpvParserCFGTest_EmitBody_BranchConditional_Fallthrough_Fallthrough_Same.spvasm b/test/tint/unittest/reader/spirv/SpvParserCFGTest_EmitBody_BranchConditional_Fallthrough_Fallthrough_Same.spvasm
deleted file mode 100644
index 0560ea2..0000000
--- a/test/tint/unittest/reader/spirv/SpvParserCFGTest_EmitBody_BranchConditional_Fallthrough_Fallthrough_Same.spvasm
+++ /dev/null
@@ -1,57 +0,0 @@
-; Test: SpvParserCFGTest_EmitBody_BranchConditional_Fallthrough_Fallthrough_Same.spvasm
-; SPIR-V
-; Version: 1.0
-; Generator: Khronos SPIR-V Tools Assembler; 0
-; Bound: 1000
-; Schema: 0
-OpCapability Shader
-OpMemoryModel Logical Simple
-OpEntryPoint Fragment %100 "main"
-OpExecutionMode %100 OriginUpperLeft
-OpName %var "var"
-%void = OpTypeVoid
-%3 = OpTypeFunction %void
-%bool = OpTypeBool
-%5 = OpConstantNull %bool
-%true = OpConstantTrue %bool
-%false = OpConstantFalse %bool
-%uint = OpTypeInt 32 0
-%int = OpTypeInt 32 1
-%uint_42 = OpConstant %uint 42
-%int_42 = OpConstant %int 42
-%13 = OpTypeFunction %uint
-%uint_0 = OpConstant %uint 0
-%uint_1 = OpConstant %uint 1
-%uint_2 = OpConstant %uint 2
-%uint_3 = OpConstant %uint 3
-%uint_4 = OpConstant %uint 4
-%uint_5 = OpConstant %uint 5
-%uint_6 = OpConstant %uint 6
-%uint_7 = OpConstant %uint 7
-%uint_8 = OpConstant %uint 8
-%uint_10 = OpConstant %uint 10
-%uint_20 = OpConstant %uint 20
-%uint_30 = OpConstant %uint 30
-%uint_40 = OpConstant %uint 40
-%uint_50 = OpConstant %uint 50
-%uint_90 = OpConstant %uint 90
-%uint_99 = OpConstant %uint 99
-%_ptr_Private_uint = OpTypePointer Private %uint
-%var = OpVariable %_ptr_Private_uint Private
-%uint_999 = OpConstant %uint 999
-%100 = OpFunction %void None %3
-%10 = OpLabel
-OpStore %var %uint_1
-OpSelectionMerge %99 None
-OpSwitch %uint_42 %99 20 %20 30 %30
-%20 = OpLabel
-OpStore %var %uint_20
-OpBranchConditional %5 %30 %30
-%30 = OpLabel
-OpStore %var %uint_30
-OpBranch %99
-%99 = OpLabel
-OpStore %var %uint_7
-OpReturn
-OpFunctionEnd
-
diff --git a/test/tint/unittest/reader/spirv/SpvParserCFGTest_EmitBody_BranchConditional_SwitchBreak_Fallthrough_OnFalse.spvasm b/test/tint/unittest/reader/spirv/SpvParserCFGTest_EmitBody_BranchConditional_SwitchBreak_Fallthrough_OnFalse.spvasm
deleted file mode 100644
index 7f1d386..0000000
--- a/test/tint/unittest/reader/spirv/SpvParserCFGTest_EmitBody_BranchConditional_SwitchBreak_Fallthrough_OnFalse.spvasm
+++ /dev/null
@@ -1,57 +0,0 @@
-; Test: SpvParserCFGTest_EmitBody_BranchConditional_SwitchBreak_Fallthrough_OnFalse.spvasm
-; SPIR-V
-; Version: 1.0
-; Generator: Khronos SPIR-V Tools Assembler; 0
-; Bound: 1000
-; Schema: 0
-OpCapability Shader
-OpMemoryModel Logical Simple
-OpEntryPoint Fragment %100 "main"
-OpExecutionMode %100 OriginUpperLeft
-OpName %var "var"
-%void = OpTypeVoid
-%3 = OpTypeFunction %void
-%bool = OpTypeBool
-%5 = OpConstantNull %bool
-%true = OpConstantTrue %bool
-%false = OpConstantFalse %bool
-%uint = OpTypeInt 32 0
-%int = OpTypeInt 32 1
-%uint_42 = OpConstant %uint 42
-%int_42 = OpConstant %int 42
-%13 = OpTypeFunction %uint
-%uint_0 = OpConstant %uint 0
-%uint_1 = OpConstant %uint 1
-%uint_2 = OpConstant %uint 2
-%uint_3 = OpConstant %uint 3
-%uint_4 = OpConstant %uint 4
-%uint_5 = OpConstant %uint 5
-%uint_6 = OpConstant %uint 6
-%uint_7 = OpConstant %uint 7
-%uint_8 = OpConstant %uint 8
-%uint_10 = OpConstant %uint 10
-%uint_20 = OpConstant %uint 20
-%uint_30 = OpConstant %uint 30
-%uint_40 = OpConstant %uint 40
-%uint_50 = OpConstant %uint 50
-%uint_90 = OpConstant %uint 90
-%uint_99 = OpConstant %uint 99
-%_ptr_Private_uint = OpTypePointer Private %uint
-%var = OpVariable %_ptr_Private_uint Private
-%uint_999 = OpConstant %uint 999
-%100 = OpFunction %void None %3
-%10 = OpLabel
-OpStore %var %uint_1
-OpSelectionMerge %99 None
-OpSwitch %uint_42 %99 20 %20 30 %30
-%20 = OpLabel
-OpStore %var %uint_20
-OpBranchConditional %5 %99 %30
-%30 = OpLabel
-OpStore %var %uint_30
-OpBranch %99
-%99 = OpLabel
-OpStore %var %uint_7
-OpReturn
-OpFunctionEnd
-
diff --git a/test/tint/unittest/reader/spirv/SpvParserCFGTest_EmitBody_BranchConditional_SwitchBreak_Fallthrough_OnTrue.spvasm b/test/tint/unittest/reader/spirv/SpvParserCFGTest_EmitBody_BranchConditional_SwitchBreak_Fallthrough_OnTrue.spvasm
deleted file mode 100644
index 5991963..0000000
--- a/test/tint/unittest/reader/spirv/SpvParserCFGTest_EmitBody_BranchConditional_SwitchBreak_Fallthrough_OnTrue.spvasm
+++ /dev/null
@@ -1,57 +0,0 @@
-; Test: SpvParserCFGTest_EmitBody_BranchConditional_SwitchBreak_Fallthrough_OnTrue.spvasm
-; SPIR-V
-; Version: 1.0
-; Generator: Khronos SPIR-V Tools Assembler; 0
-; Bound: 1000
-; Schema: 0
-OpCapability Shader
-OpMemoryModel Logical Simple
-OpEntryPoint Fragment %100 "main"
-OpExecutionMode %100 OriginUpperLeft
-OpName %var "var"
-%void = OpTypeVoid
-%3 = OpTypeFunction %void
-%bool = OpTypeBool
-%5 = OpConstantNull %bool
-%true = OpConstantTrue %bool
-%false = OpConstantFalse %bool
-%uint = OpTypeInt 32 0
-%int = OpTypeInt 32 1
-%uint_42 = OpConstant %uint 42
-%int_42 = OpConstant %int 42
-%13 = OpTypeFunction %uint
-%uint_0 = OpConstant %uint 0
-%uint_1 = OpConstant %uint 1
-%uint_2 = OpConstant %uint 2
-%uint_3 = OpConstant %uint 3
-%uint_4 = OpConstant %uint 4
-%uint_5 = OpConstant %uint 5
-%uint_6 = OpConstant %uint 6
-%uint_7 = OpConstant %uint 7
-%uint_8 = OpConstant %uint 8
-%uint_10 = OpConstant %uint 10
-%uint_20 = OpConstant %uint 20
-%uint_30 = OpConstant %uint 30
-%uint_40 = OpConstant %uint 40
-%uint_50 = OpConstant %uint 50
-%uint_90 = OpConstant %uint 90
-%uint_99 = OpConstant %uint 99
-%_ptr_Private_uint = OpTypePointer Private %uint
-%var = OpVariable %_ptr_Private_uint Private
-%uint_999 = OpConstant %uint 999
-%100 = OpFunction %void None %3
-%10 = OpLabel
-OpStore %var %uint_1
-OpSelectionMerge %99 None
-OpSwitch %uint_42 %99 20 %20 30 %30
-%20 = OpLabel
-OpStore %var %uint_20
-OpBranchConditional %5 %30 %99
-%30 = OpLabel
-OpStore %var %uint_30
-OpBranch %99
-%99 = OpLabel
-OpStore %var %uint_7
-OpReturn
-OpFunctionEnd
-
diff --git a/test/tint/unittest/reader/spirv/SpvParserCFGTest_EmitBody_Branch_Fallthrough.spvasm b/test/tint/unittest/reader/spirv/SpvParserCFGTest_EmitBody_Branch_Fallthrough.spvasm
deleted file mode 100644
index 6e9d1eb..0000000
--- a/test/tint/unittest/reader/spirv/SpvParserCFGTest_EmitBody_Branch_Fallthrough.spvasm
+++ /dev/null
@@ -1,57 +0,0 @@
-; Test: SpvParserCFGTest_EmitBody_Branch_Fallthrough.spvasm
-; SPIR-V
-; Version: 1.0
-; Generator: Khronos SPIR-V Tools Assembler; 0
-; Bound: 1000
-; Schema: 0
-OpCapability Shader
-OpMemoryModel Logical Simple
-OpEntryPoint Fragment %100 "main"
-OpExecutionMode %100 OriginUpperLeft
-OpName %var "var"
-%void = OpTypeVoid
-%3 = OpTypeFunction %void
-%bool = OpTypeBool
-%5 = OpConstantNull %bool
-%true = OpConstantTrue %bool
-%false = OpConstantFalse %bool
-%uint = OpTypeInt 32 0
-%int = OpTypeInt 32 1
-%uint_42 = OpConstant %uint 42
-%int_42 = OpConstant %int 42
-%13 = OpTypeFunction %uint
-%uint_0 = OpConstant %uint 0
-%uint_1 = OpConstant %uint 1
-%uint_2 = OpConstant %uint 2
-%uint_3 = OpConstant %uint 3
-%uint_4 = OpConstant %uint 4
-%uint_5 = OpConstant %uint 5
-%uint_6 = OpConstant %uint 6
-%uint_7 = OpConstant %uint 7
-%uint_8 = OpConstant %uint 8
-%uint_10 = OpConstant %uint 10
-%uint_20 = OpConstant %uint 20
-%uint_30 = OpConstant %uint 30
-%uint_40 = OpConstant %uint 40
-%uint_50 = OpConstant %uint 50
-%uint_90 = OpConstant %uint 90
-%uint_99 = OpConstant %uint 99
-%_ptr_Private_uint = OpTypePointer Private %uint
-%var = OpVariable %_ptr_Private_uint Private
-%uint_999 = OpConstant %uint 999
-%100 = OpFunction %void None %3
-%10 = OpLabel
-OpStore %var %uint_1
-OpSelectionMerge %99 None
-OpSwitch %uint_42 %99 20 %20 30 %30
-%20 = OpLabel
-OpStore %var %uint_20
-OpBranch %30
-%30 = OpLabel
-OpStore %var %uint_30
-OpBranch %99
-%99 = OpLabel
-OpStore %var %uint_7
-OpReturn
-OpFunctionEnd
-
diff --git a/test/tint/unittest/reader/spirv/SpvParserCFGTest_TerminatorsAreValid_Switch.spvasm b/test/tint/unittest/reader/spirv/SpvParserCFGTest_TerminatorsAreValid_Switch.spvasm
deleted file mode 100644
index 4269988..0000000
--- a/test/tint/unittest/reader/spirv/SpvParserCFGTest_TerminatorsAreValid_Switch.spvasm
+++ /dev/null
@@ -1,55 +0,0 @@
-; Test: SpvParserCFGTest_TerminatorsAreValid_Switch.spvasm
-; SPIR-V
-; Version: 1.0
-; Generator: Khronos SPIR-V Tools Assembler; 0
-; Bound: 1000
-; Schema: 0
-OpCapability Shader
-OpMemoryModel Logical Simple
-OpEntryPoint Fragment %100 "main"
-OpExecutionMode %100 OriginUpperLeft
-OpName %var "var"
-%void = OpTypeVoid
-%3 = OpTypeFunction %void
-%bool = OpTypeBool
-%5 = OpConstantNull %bool
-%true = OpConstantTrue %bool
-%false = OpConstantFalse %bool
-%uint = OpTypeInt 32 0
-%int = OpTypeInt 32 1
-%uint_42 = OpConstant %uint 42
-%int_42 = OpConstant %int 42
-%13 = OpTypeFunction %uint
-%uint_0 = OpConstant %uint 0
-%uint_1 = OpConstant %uint 1
-%uint_2 = OpConstant %uint 2
-%uint_3 = OpConstant %uint 3
-%uint_4 = OpConstant %uint 4
-%uint_5 = OpConstant %uint 5
-%uint_6 = OpConstant %uint 6
-%uint_7 = OpConstant %uint 7
-%uint_8 = OpConstant %uint 8
-%uint_10 = OpConstant %uint 10
-%uint_20 = OpConstant %uint 20
-%uint_30 = OpConstant %uint 30
-%uint_40 = OpConstant %uint 40
-%uint_50 = OpConstant %uint 50
-%uint_90 = OpConstant %uint 90
-%uint_99 = OpConstant %uint 99
-%_ptr_Private_uint = OpTypePointer Private %uint
-%var = OpVariable %_ptr_Private_uint Private
-%uint_999 = OpConstant %uint 999
-%100 = OpFunction %void None %3
-%10 = OpLabel
-OpSelectionMerge %99 None
-OpSwitch %uint_42 %80 20 %20 30 %30
-%20 = OpLabel
-OpBranch %30
-%30 = OpLabel
-OpBranch %99
-%80 = OpLabel
-OpBranch %99
-%99 = OpLabel
-OpReturn
-OpFunctionEnd
-
diff --git a/test/tint/unittest/reader/spirv/SpvParserFunctionVarTest_EmitStatement_Phi_InMerge_PredecessorsDominatdByNestedSwitchCase.spvasm b/test/tint/unittest/reader/spirv/SpvParserFunctionVarTest_EmitStatement_Phi_InMerge_PredecessorsDominatdByNestedSwitchCase.spvasm
deleted file mode 100644
index 83d008f..0000000
--- a/test/tint/unittest/reader/spirv/SpvParserFunctionVarTest_EmitStatement_Phi_InMerge_PredecessorsDominatdByNestedSwitchCase.spvasm
+++ /dev/null
@@ -1,66 +0,0 @@
-; Test: SpvParserFunctionVarTest_EmitStatement_Phi_InMerge_PredecessorsDominatdByNestedSwitchCase.spvasm
-; SPIR-V
-; Version: 1.0
-; Generator: Khronos SPIR-V Tools Assembler; 0
-; Bound: 101
-; Schema: 0
-OpCapability Shader
-OpMemoryModel Logical Simple
-OpEntryPoint Fragment %100 "main"
-OpExecutionMode %100 OriginUpperLeft
-%void = OpTypeVoid
-%3 = OpTypeFunction %void
-%bool = OpTypeBool
-%float = OpTypeFloat 32
-%uint = OpTypeInt 32 0
-%int = OpTypeInt 32 1
-%_ptr_Function_bool = OpTypePointer Function %bool
-%_ptr_Function_float = OpTypePointer Function %float
-%_ptr_Function_uint = OpTypePointer Function %uint
-%_ptr_Function_int = OpTypePointer Function %int
-%true = OpConstantTrue %bool
-%false = OpConstantFalse %bool
-%float_0 = OpConstant %float 0
-%float_1_5 = OpConstant %float 1.5
-%uint_0 = OpConstant %uint 0
-%uint_1 = OpConstant %uint 1
-%int_n1 = OpConstant %int -1
-%int_0 = OpConstant %int 0
-%int_1 = OpConstant %int 1
-%int_3 = OpConstant %int 3
-%uint_2 = OpConstant %uint 2
-%uint_3 = OpConstant %uint 3
-%uint_4 = OpConstant %uint 4
-%uint_5 = OpConstant %uint 5
-%v2int = OpTypeVector %int 2
-%v2float = OpTypeVector %float 2
-%mat3v2float = OpTypeMatrix %v2float 3
-%34 = OpConstantNull %v2int
-%_arr_uint_uint_2 = OpTypeArray %uint %uint_2
-%_struct_36 = OpTypeStruct %uint %float %_arr_uint_uint_2
-%true_0 = OpConstantTrue %bool
-%_ptr_Private_uint = OpTypePointer Private %uint
-%1 = OpVariable %_ptr_Private_uint Private
-%_ptr_Private_bool = OpTypePointer Private %bool
-%7 = OpVariable %_ptr_Private_bool Private
-%8 = OpVariable %_ptr_Private_bool Private
-%100 = OpFunction %void None %3
-%10 = OpLabel
-OpSelectionMerge %99 None
-OpSwitch %uint_1 %20 0 %20 1 %30
-%20 = OpLabel
-OpBranch %30
-%30 = OpLabel
-OpSelectionMerge %50 None
-OpBranchConditional %true %40 %45
-%40 = OpLabel
-OpBranch %50
-%45 = OpLabel
-OpBranch %99
-%50 = OpLabel
-OpBranch %99
-%99 = OpLabel
-%41 = OpPhi %uint %uint_0 %45 %uint_1 %50
-OpReturn
-OpFunctionEnd
-
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/call-if-while-switch/0-opt.spvasm b/test/tint/vk-gl-cts/graphicsfuzz/call-if-while-switch/0-opt.spvasm
deleted file mode 100644
index 07f5e01..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/call-if-while-switch/0-opt.spvasm
+++ /dev/null
@@ -1,93 +0,0 @@
- OpCapability Shader
- %1 = OpExtInstImport "GLSL.std.450"
- OpMemoryModel Logical GLSL450
- OpEntryPoint Fragment %main "main" %_GLF_color
- OpExecutionMode %main OriginUpperLeft
- OpSource ESSL 310
- OpName %main "main"
- OpName %data "data"
- OpName %buf0 "buf0"
- OpMemberName %buf0 0 "injectionSwitch"
- OpName %_ ""
- OpName %_GLF_color "_GLF_color"
- OpDecorate %data RelaxedPrecision
- OpMemberDecorate %buf0 0 Offset 0
- OpDecorate %buf0 Block
- OpDecorate %_ DescriptorSet 0
- OpDecorate %_ Binding 0
- OpDecorate %_GLF_color Location 0
- OpDecorate %7 RelaxedPrecision
- OpDecorate %8 RelaxedPrecision
- OpDecorate %9 RelaxedPrecision
- OpDecorate %10 RelaxedPrecision
- OpDecorate %11 RelaxedPrecision
- %void = OpTypeVoid
- %13 = OpTypeFunction %void
- %int = OpTypeInt 32 1
-%_ptr_Function_int = OpTypePointer Function %int
- %int_1 = OpConstant %int 1
- %uint = OpTypeInt 32 0
- %uint_10 = OpConstant %uint 10
-%_arr_int_uint_10 = OpTypeArray %int %uint_10
- %bool = OpTypeBool
- %int_3 = OpConstant %int 3
- %float = OpTypeFloat 32
- %v2float = OpTypeVector %float 2
- %buf0 = OpTypeStruct %v2float
-%_ptr_Uniform_buf0 = OpTypePointer Uniform %buf0
- %_ = OpVariable %_ptr_Uniform_buf0 Uniform
- %int_0 = OpConstant %int 0
- %uint_0 = OpConstant %uint 0
-%_ptr_Uniform_float = OpTypePointer Uniform %float
- %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
- %_GLF_color = OpVariable %_ptr_Output_v4float Output
- %float_1 = OpConstant %float 1
- %31 = OpConstantComposite %v4float %float_1 %float_1 %float_1 %float_1
- %float_0 = OpConstant %float 0
- %33 = OpConstantComposite %v4float %float_1 %float_0 %float_0 %float_1
-%_ptr_Function__arr_int_uint_10 = OpTypePointer Function %_arr_int_uint_10
- %int_2 = OpConstant %int 2
- %main = OpFunction %void None %13
- %36 = OpLabel
- %data = OpVariable %_ptr_Function__arr_int_uint_10 Function
- %37 = OpAccessChain %_ptr_Function_int %data %int_1
- %7 = OpLoad %int %37
- %38 = OpSLessThan %bool %int_1 %7
- %10 = OpSelect %int %38 %int_2 %int_1
- OpBranch %39
- %39 = OpLabel
- %40 = OpPhi %int %int_1 %36 %41 %42
- %11 = OpPhi %int %10 %36 %8 %42
- %43 = OpSLessThan %bool %11 %int_3
- OpLoopMerge %44 %42 None
- OpBranchConditional %43 %45 %44
- %45 = OpLabel
- %8 = OpIAdd %int %11 %int_1
- %46 = OpAccessChain %_ptr_Uniform_float %_ %int_0 %uint_0
- %47 = OpLoad %float %46
- %48 = OpConvertFToS %int %47
- OpSelectionMerge %49 None
- OpSwitch %48 %50 19 %51 38 %52 23 %52 78 %53
- %50 = OpLabel
- OpBranch %42
- %51 = OpLabel
- %9 = OpIAdd %int %40 %int_1
- OpBranch %52
- %52 = OpLabel
- %54 = OpPhi %int %40 %45 %9 %51
- OpBranch %42
- %53 = OpLabel
- OpStore %_GLF_color %31
- OpBranch %50
- %49 = OpLabel
- OpBranch %42
- %42 = OpLabel
- %41 = OpPhi %int %40 %50 %54 %52 %int_0 %49
- OpBranch %39
- %44 = OpLabel
- %55 = OpAccessChain %_ptr_Function_int %data %40
- OpStore %55 %int_1
- OpStore %_GLF_color %33
- OpReturn
- OpFunctionEnd
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/call-if-while-switch/0-opt.wgsl b/test/tint/vk-gl-cts/graphicsfuzz/call-if-while-switch/0-opt.wgsl
deleted file mode 100644
index 00cdb63..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/call-if-while-switch/0-opt.wgsl
+++ /dev/null
@@ -1,72 +0,0 @@
-struct buf0 {
- injectionSwitch : vec2<f32>,
-}
-
-@group(0) @binding(0) var<uniform> x_6 : buf0;
-
-var<private> x_GLF_color : vec4<f32>;
-
-fn main_1() {
- var data : array<i32, 10u>;
- var x_40 : i32;
- var x_40_phi : i32;
- var x_11_phi : i32;
- let x_7 : i32 = data[1];
- let x_10 : i32 = select(1, 2, (1 < x_7));
- x_40_phi = 1;
- x_11_phi = x_10;
- loop {
- var x_54 : i32;
- var x_41 : i32;
- var x_41_phi : i32;
- x_40 = x_40_phi;
- let x_11 : i32 = x_11_phi;
- if ((x_11 < 3)) {
- } else {
- break;
- }
- var x_54_phi : i32;
- let x_8 : i32 = (x_11 + 1);
- let x_47 : f32 = x_6.injectionSwitch.x;
- x_54_phi = x_40;
- switch(i32(x_47)) {
- case 78: {
- x_GLF_color = vec4<f32>(1.0, 1.0, 1.0, 1.0);
- fallthrough;
- }
- case 19: {
- x_54_phi = bitcast<i32>((x_40 + bitcast<i32>(1)));
- fallthrough;
- }
- case 23, 38: {
- x_54 = x_54_phi;
- x_41_phi = x_54;
- continue;
- }
- default: {
- x_41_phi = x_40;
- continue;
- }
- }
-
- continuing {
- x_41 = x_41_phi;
- x_40_phi = x_41;
- x_11_phi = x_8;
- }
- }
- data[x_40] = 1;
- x_GLF_color = vec4<f32>(1.0, 0.0, 0.0, 1.0);
- return;
-}
-
-struct main_out {
- @location(0)
- x_GLF_color_1 : vec4<f32>,
-}
-
-@fragment
-fn main() -> main_out {
- main_1();
- return main_out(x_GLF_color);
-}
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-dead-code-unreachable-merge/0-opt.spvasm b/test/tint/vk-gl-cts/graphicsfuzz/cov-dead-code-unreachable-merge/0-opt.spvasm
deleted file mode 100644
index 3ccb8a9..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-dead-code-unreachable-merge/0-opt.spvasm
+++ /dev/null
@@ -1,154 +0,0 @@
- OpCapability Shader
- %1 = OpExtInstImport "GLSL.std.450"
- OpMemoryModel Logical GLSL450
- OpEntryPoint Fragment %main "main" %gl_FragCoord %_GLF_color
- OpExecutionMode %main OriginUpperLeft
- OpSource ESSL 310
- OpName %main "main"
- OpName %q "q"
- OpName %i "i"
- OpName %gl_FragCoord "gl_FragCoord"
- OpName %c "c"
- OpName %array0 "array0"
- OpName %array1 "array1"
- OpName %buf0 "buf0"
- OpMemberName %buf0 0 "injectionSwitch"
- OpName %_ ""
- OpName %_GLF_color "_GLF_color"
- OpDecorate %q RelaxedPrecision
- OpDecorate %i RelaxedPrecision
- OpDecorate %gl_FragCoord BuiltIn FragCoord
- OpDecorate %12 RelaxedPrecision
- OpDecorate %13 RelaxedPrecision
- OpDecorate %c RelaxedPrecision
- OpDecorate %14 RelaxedPrecision
- OpDecorate %15 RelaxedPrecision
- OpDecorate %16 RelaxedPrecision
- OpMemberDecorate %buf0 0 Offset 0
- OpDecorate %buf0 Block
- OpDecorate %_ DescriptorSet 0
- OpDecorate %_ Binding 0
- OpDecorate %17 RelaxedPrecision
- OpDecorate %18 RelaxedPrecision
- OpDecorate %19 RelaxedPrecision
- OpDecorate %20 RelaxedPrecision
- OpDecorate %21 RelaxedPrecision
- OpDecorate %22 RelaxedPrecision
- OpDecorate %23 RelaxedPrecision
- OpDecorate %_GLF_color Location 0
- OpDecorate %24 RelaxedPrecision
- OpDecorate %25 RelaxedPrecision
- OpDecorate %26 RelaxedPrecision
- %void = OpTypeVoid
- %28 = OpTypeFunction %void
- %int = OpTypeInt 32 1
-%_ptr_Function_int = OpTypePointer Function %int
- %int_0 = OpConstant %int 0
- %float = OpTypeFloat 32
- %v4float = OpTypeVector %float 4
-%_ptr_Input_v4float = OpTypePointer Input %v4float
-%gl_FragCoord = OpVariable %_ptr_Input_v4float Input
- %uint = OpTypeInt 32 0
- %uint_0 = OpConstant %uint 0
-%_ptr_Input_float = OpTypePointer Input %float
- %int_3 = OpConstant %int 3
- %bool = OpTypeBool
- %uint_3 = OpConstant %uint 3
-%_arr_float_uint_3 = OpTypeArray %float %uint_3
-%_ptr_Private__arr_float_uint_3 = OpTypePointer Private %_arr_float_uint_3
- %array0 = OpVariable %_ptr_Private__arr_float_uint_3 Private
- %float_0 = OpConstant %float 0
-%_ptr_Private_float = OpTypePointer Private %float
- %array1 = OpVariable %_ptr_Private__arr_float_uint_3 Private
- %v2float = OpTypeVector %float 2
- %buf0 = OpTypeStruct %v2float
-%_ptr_Uniform_buf0 = OpTypePointer Uniform %buf0
- %_ = OpVariable %_ptr_Uniform_buf0 Uniform
-%_ptr_Uniform_float = OpTypePointer Uniform %float
- %int_61 = OpConstant %int 61
- %true = OpConstantTrue %bool
- %float_1 = OpConstant %float 1
- %int_1 = OpConstant %int 1
-%_ptr_Output_v4float = OpTypePointer Output %v4float
- %_GLF_color = OpVariable %_ptr_Output_v4float Output
- %main = OpFunction %void None %28
- %53 = OpLabel
- %q = OpVariable %_ptr_Function_int Function
- %i = OpVariable %_ptr_Function_int Function
- %c = OpVariable %_ptr_Function_int Function
- OpStore %q %int_0
- %54 = OpAccessChain %_ptr_Input_float %gl_FragCoord %uint_0
- %55 = OpLoad %float %54
- %12 = OpConvertFToS %int %55
- %13 = OpSMod %int %12 %int_3
- OpStore %i %13
- OpStore %c %int_0
- OpBranch %56
- %56 = OpLabel
- OpLoopMerge %57 %58 None
- OpBranch %59
- %59 = OpLabel
- %14 = OpLoad %int %c
- %60 = OpSLessThan %bool %14 %int_3
- OpBranchConditional %60 %61 %57
- %61 = OpLabel
- %15 = OpLoad %int %c
- %62 = OpAccessChain %_ptr_Private_float %array0 %15
- OpStore %62 %float_0
- %16 = OpLoad %int %c
- %63 = OpAccessChain %_ptr_Private_float %array1 %16
- OpStore %63 %float_0
- %64 = OpAccessChain %_ptr_Uniform_float %_ %int_0 %uint_0
- %65 = OpLoad %float %64
- %17 = OpConvertFToS %int %65
- %18 = OpLoad %int %q
- %19 = OpIAdd %int %17 %18
- OpSelectionMerge %66 None
- OpSwitch %19 %66 0 %67 51 %68 61 %69
- %67 = OpLabel
- OpStore %q %int_61
- OpBranch %66
- %68 = OpLabel
- OpBranch %70
- %70 = OpLabel
- OpLoopMerge %71 %72 None
- OpBranch %73
- %73 = OpLabel
- OpBranchConditional %true %74 %71
- %74 = OpLabel
- OpBranch %72
- %72 = OpLabel
- OpBranch %70
- %71 = OpLabel
- %20 = OpLoad %int %c
- %75 = OpAccessChain %_ptr_Private_float %array0 %20
- OpStore %75 %float_1
- OpBranch %69
- %69 = OpLabel
- %76 = OpAccessChain %_ptr_Private_float %array1 %int_0
- OpStore %76 %float_1
- %21 = OpLoad %int %c
- %77 = OpAccessChain %_ptr_Private_float %array1 %21
- OpStore %77 %float_1
- OpBranch %66
- %66 = OpLabel
- OpBranch %58
- %58 = OpLabel
- %22 = OpLoad %int %c
- %23 = OpIAdd %int %22 %int_1
- OpStore %c %23
- OpBranch %56
- %57 = OpLabel
- %24 = OpLoad %int %i
- %78 = OpAccessChain %_ptr_Private_float %array1 %24
- %79 = OpLoad %float %78
- %25 = OpLoad %int %i
- %80 = OpAccessChain %_ptr_Private_float %array0 %25
- %81 = OpLoad %float %80
- %26 = OpLoad %int %i
- %82 = OpAccessChain %_ptr_Private_float %array0 %26
- %83 = OpLoad %float %82
- %84 = OpCompositeConstruct %v4float %79 %81 %83 %float_1
- OpStore %_GLF_color %84
- OpReturn
- OpFunctionEnd
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-dead-code-unreachable-merge/0-opt.spvasm.expected.dxc.hlsl b/test/tint/vk-gl-cts/graphicsfuzz/cov-dead-code-unreachable-merge/0-opt.spvasm.expected.dxc.hlsl
deleted file mode 100755
index 2e7ba76..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-dead-code-unreachable-merge/0-opt.spvasm.expected.dxc.hlsl
+++ /dev/null
@@ -1,93 +0,0 @@
-SKIP: FAILED
-
-static float4 gl_FragCoord = float4(0.0f, 0.0f, 0.0f, 0.0f);
-static float array0[3] = (float[3])0;
-static float array1[3] = (float[3])0;
-cbuffer cbuffer_x_11 : register(b0, space0) {
- uint4 x_11[1];
-};
-static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f);
-
-void main_1() {
- int q = 0;
- int i = 0;
- int c = 0;
- q = 0;
- const float x_55 = gl_FragCoord.x;
- i = (int(x_55) % 3);
- c = 0;
- {
- for(; (c < 3); c = (c + 1)) {
- array0[c] = 0.0f;
- array1[c] = 0.0f;
- const float x_65 = asfloat(x_11[0].x);
- switch((int(x_65) + q)) {
- case 51: {
- while (true) {
- if (true) {
- } else {
- break;
- }
- }
- array0[c] = 1.0f;
- /* fallthrough */
- {
- array1[0] = 1.0f;
- array1[c] = 1.0f;
- }
- break;
- }
- case 61: {
- array1[0] = 1.0f;
- array1[c] = 1.0f;
- break;
- }
- case 0: {
- q = 61;
- break;
- }
- default: {
- break;
- }
- }
- }
- }
- const float x_79 = array1[i];
- const float x_81 = array0[i];
- const float x_83 = array0[i];
- x_GLF_color = float4(x_79, x_81, x_83, 1.0f);
- return;
-}
-
-struct main_out {
- float4 x_GLF_color_1;
-};
-struct tint_symbol_1 {
- float4 gl_FragCoord_param : SV_Position;
-};
-struct tint_symbol_2 {
- float4 x_GLF_color_1 : SV_Target0;
-};
-
-main_out main_inner(float4 gl_FragCoord_param) {
- gl_FragCoord = gl_FragCoord_param;
- main_1();
- const main_out tint_symbol_4 = {x_GLF_color};
- return tint_symbol_4;
-}
-
-tint_symbol_2 main(tint_symbol_1 tint_symbol) {
- const main_out inner_result = main_inner(tint_symbol.gl_FragCoord_param);
- tint_symbol_2 wrapper_result = (tint_symbol_2)0;
- wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1;
- return wrapper_result;
-}
-DXC validation failure:
-warning: DXIL.dll not found. Resulting DXIL will not be signed for use in release environments.
-
-error: validation errors
-shader.hlsl:77: error: Loop must have break.
-Validation failed.
-
-
-
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-dead-code-unreachable-merge/0-opt.wgsl b/test/tint/vk-gl-cts/graphicsfuzz/cov-dead-code-unreachable-merge/0-opt.wgsl
deleted file mode 100644
index 6042980..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-dead-code-unreachable-merge/0-opt.wgsl
+++ /dev/null
@@ -1,84 +0,0 @@
-struct buf0 {
- injectionSwitch : vec2<f32>,
-}
-
-var<private> gl_FragCoord : vec4<f32>;
-
-var<private> array0 : array<f32, 3u>;
-
-var<private> array1 : array<f32, 3u>;
-
-@group(0) @binding(0) var<uniform> x_11 : buf0;
-
-var<private> x_GLF_color : vec4<f32>;
-
-fn main_1() {
- var q : i32;
- var i : i32;
- var c : i32;
- q = 0;
- let x_55 : f32 = gl_FragCoord.x;
- i = (i32(x_55) % 3);
- c = 0;
- loop {
- let x_14 : i32 = c;
- if ((x_14 < 3)) {
- } else {
- break;
- }
- let x_15 : i32 = c;
- array0[x_15] = 0.0;
- let x_16 : i32 = c;
- array1[x_16] = 0.0;
- let x_65 : f32 = x_11.injectionSwitch.x;
- let x_18 : i32 = q;
- switch((i32(x_65) + x_18)) {
- case 51: {
- loop {
- if (true) {
- } else {
- break;
- }
- }
- let x_20 : i32 = c;
- array0[x_20] = 1.0;
- fallthrough;
- }
- case 61: {
- array1[0] = 1.0;
- let x_21 : i32 = c;
- array1[x_21] = 1.0;
- }
- case 0: {
- q = 61;
- }
- default: {
- }
- }
-
- continuing {
- let x_22 : i32 = c;
- c = (x_22 + 1);
- }
- }
- let x_24 : i32 = i;
- let x_79 : f32 = array1[x_24];
- let x_25 : i32 = i;
- let x_81 : f32 = array0[x_25];
- let x_26 : i32 = i;
- let x_83 : f32 = array0[x_26];
- x_GLF_color = vec4<f32>(x_79, x_81, x_83, 1.0);
- return;
-}
-
-struct main_out {
- @location(0)
- x_GLF_color_1 : vec4<f32>,
-}
-
-@fragment
-fn main(@builtin(position) gl_FragCoord_param : vec4<f32>) -> main_out {
- gl_FragCoord = gl_FragCoord_param;
- main_1();
- return main_out(x_GLF_color);
-}
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-dead-code-unreachable-merge/0-opt.wgsl.expected.dxc.hlsl b/test/tint/vk-gl-cts/graphicsfuzz/cov-dead-code-unreachable-merge/0-opt.wgsl.expected.dxc.hlsl
deleted file mode 100755
index 2e7ba76..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-dead-code-unreachable-merge/0-opt.wgsl.expected.dxc.hlsl
+++ /dev/null
@@ -1,93 +0,0 @@
-SKIP: FAILED
-
-static float4 gl_FragCoord = float4(0.0f, 0.0f, 0.0f, 0.0f);
-static float array0[3] = (float[3])0;
-static float array1[3] = (float[3])0;
-cbuffer cbuffer_x_11 : register(b0, space0) {
- uint4 x_11[1];
-};
-static float4 x_GLF_color = float4(0.0f, 0.0f, 0.0f, 0.0f);
-
-void main_1() {
- int q = 0;
- int i = 0;
- int c = 0;
- q = 0;
- const float x_55 = gl_FragCoord.x;
- i = (int(x_55) % 3);
- c = 0;
- {
- for(; (c < 3); c = (c + 1)) {
- array0[c] = 0.0f;
- array1[c] = 0.0f;
- const float x_65 = asfloat(x_11[0].x);
- switch((int(x_65) + q)) {
- case 51: {
- while (true) {
- if (true) {
- } else {
- break;
- }
- }
- array0[c] = 1.0f;
- /* fallthrough */
- {
- array1[0] = 1.0f;
- array1[c] = 1.0f;
- }
- break;
- }
- case 61: {
- array1[0] = 1.0f;
- array1[c] = 1.0f;
- break;
- }
- case 0: {
- q = 61;
- break;
- }
- default: {
- break;
- }
- }
- }
- }
- const float x_79 = array1[i];
- const float x_81 = array0[i];
- const float x_83 = array0[i];
- x_GLF_color = float4(x_79, x_81, x_83, 1.0f);
- return;
-}
-
-struct main_out {
- float4 x_GLF_color_1;
-};
-struct tint_symbol_1 {
- float4 gl_FragCoord_param : SV_Position;
-};
-struct tint_symbol_2 {
- float4 x_GLF_color_1 : SV_Target0;
-};
-
-main_out main_inner(float4 gl_FragCoord_param) {
- gl_FragCoord = gl_FragCoord_param;
- main_1();
- const main_out tint_symbol_4 = {x_GLF_color};
- return tint_symbol_4;
-}
-
-tint_symbol_2 main(tint_symbol_1 tint_symbol) {
- const main_out inner_result = main_inner(tint_symbol.gl_FragCoord_param);
- tint_symbol_2 wrapper_result = (tint_symbol_2)0;
- wrapper_result.x_GLF_color_1 = inner_result.x_GLF_color_1;
- return wrapper_result;
-}
-DXC validation failure:
-warning: DXIL.dll not found. Resulting DXIL will not be signed for use in release environments.
-
-error: validation errors
-shader.hlsl:77: error: Loop must have break.
-Validation failed.
-
-
-
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-inst-combine-and-or-xor-switch/0-opt.spvasm b/test/tint/vk-gl-cts/graphicsfuzz/cov-inst-combine-and-or-xor-switch/0-opt.spvasm
deleted file mode 100644
index 07cab15..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-inst-combine-and-or-xor-switch/0-opt.spvasm
+++ /dev/null
@@ -1,133 +0,0 @@
- OpCapability Shader
- %1 = OpExtInstImport "GLSL.std.450"
- OpMemoryModel Logical GLSL450
- OpEntryPoint Fragment %main "main" %_GLF_color
- OpExecutionMode %main OriginUpperLeft
- OpSource ESSL 310
- OpName %main "main"
- OpName %count0 "count0"
- OpName %buf0 "buf0"
- OpMemberName %buf0 0 "_GLF_uniform_int_values"
- OpName %_ ""
- OpName %count1 "count1"
- OpName %i "i"
- OpName %_GLF_color "_GLF_color"
- OpDecorate %_arr_int_uint_5 ArrayStride 16
- OpMemberDecorate %buf0 0 Offset 0
- OpDecorate %buf0 Block
- OpDecorate %_ DescriptorSet 0
- OpDecorate %_ Binding 0
- OpDecorate %_GLF_color Location 0
- %void = OpTypeVoid
- %11 = OpTypeFunction %void
- %int = OpTypeInt 32 1
-%_ptr_Function_int = OpTypePointer Function %int
- %uint = OpTypeInt 32 0
- %uint_5 = OpConstant %uint 5
-%_arr_int_uint_5 = OpTypeArray %int %uint_5
- %buf0 = OpTypeStruct %_arr_int_uint_5
-%_ptr_Uniform_buf0 = OpTypePointer Uniform %buf0
- %_ = OpVariable %_ptr_Uniform_buf0 Uniform
- %int_0 = OpConstant %int 0
- %int_2 = OpConstant %int 2
-%_ptr_Uniform_int = OpTypePointer Uniform %int
- %int_4 = OpConstant %int 4
- %bool = OpTypeBool
- %int_1 = OpConstant %int 1
- %float = OpTypeFloat 32
- %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
- %_GLF_color = OpVariable %_ptr_Output_v4float Output
- %int_3 = OpConstant %int 3
- %main = OpFunction %void None %11
- %27 = OpLabel
- %count0 = OpVariable %_ptr_Function_int Function
- %count1 = OpVariable %_ptr_Function_int Function
- %i = OpVariable %_ptr_Function_int Function
- %28 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_2
- %29 = OpLoad %int %28
- OpStore %count0 %29
- %30 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_2
- %31 = OpLoad %int %30
- OpStore %count1 %31
- %32 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_2
- %33 = OpLoad %int %32
- OpStore %i %33
- OpBranch %34
- %34 = OpLabel
- OpLoopMerge %35 %36 None
- OpBranch %37
- %37 = OpLabel
- %38 = OpLoad %int %i
- %39 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_4
- %40 = OpLoad %int %39
- %41 = OpSLessThan %bool %38 %40
- OpBranchConditional %41 %42 %35
- %42 = OpLabel
- %43 = OpLoad %int %i
- OpSelectionMerge %44 None
- OpSwitch %43 %44 0 %45 1 %45 3 %46 2 %46
- %45 = OpLabel
- %47 = OpLoad %int %count0
- %48 = OpIAdd %int %47 %int_1
- OpStore %count0 %48
- OpBranch %46
- %46 = OpLabel
- %49 = OpLoad %int %count1
- %50 = OpIAdd %int %49 %int_1
- OpStore %count1 %50
- OpBranch %44
- %44 = OpLabel
- OpBranch %36
- %36 = OpLabel
- %51 = OpLoad %int %i
- %52 = OpIAdd %int %51 %int_1
- OpStore %i %52
- OpBranch %34
- %35 = OpLabel
- %53 = OpLoad %int %count1
- %54 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_0
- %55 = OpLoad %int %54
- %56 = OpIEqual %bool %53 %55
- OpSelectionMerge %57 None
- OpBranchConditional %56 %58 %59
- %58 = OpLabel
- %60 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_3
- %61 = OpLoad %int %60
- %62 = OpConvertSToF %float %61
- %63 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_2
- %64 = OpLoad %int %63
- %65 = OpConvertSToF %float %64
- %66 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_2
- %67 = OpLoad %int %66
- %68 = OpConvertSToF %float %67
- %69 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_3
- %70 = OpLoad %int %69
- %71 = OpConvertSToF %float %70
- %72 = OpCompositeConstruct %v4float %62 %65 %68 %71
- OpStore %_GLF_color %72
- OpBranch %57
- %59 = OpLabel
- %73 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_2
- %74 = OpLoad %int %73
- %75 = OpConvertSToF %float %74
- %76 = OpCompositeConstruct %v4float %75 %75 %75 %75
- OpStore %_GLF_color %76
- OpBranch %57
- %57 = OpLabel
- %77 = OpLoad %int %count0
- %78 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_1
- %79 = OpLoad %int %78
- %80 = OpINotEqual %bool %77 %79
- OpSelectionMerge %81 None
- OpBranchConditional %80 %82 %81
- %82 = OpLabel
- %83 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_2
- %84 = OpLoad %int %83
- %85 = OpConvertSToF %float %84
- %86 = OpCompositeConstruct %v4float %85 %85 %85 %85
- OpStore %_GLF_color %86
- OpBranch %81
- %81 = OpLabel
- OpReturn
- OpFunctionEnd
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-inst-combine-and-or-xor-switch/0-opt.wgsl b/test/tint/vk-gl-cts/graphicsfuzz/cov-inst-combine-and-or-xor-switch/0-opt.wgsl
deleted file mode 100644
index 5256543..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-inst-combine-and-or-xor-switch/0-opt.wgsl
+++ /dev/null
@@ -1,85 +0,0 @@
-struct strided_arr {
- @size(16)
- el : i32,
-}
-
-type Arr = array<strided_arr, 5u>;
-
-struct buf0 {
- x_GLF_uniform_int_values : Arr,
-}
-
-@group(0) @binding(0) var<uniform> x_6 : buf0;
-
-var<private> x_GLF_color : vec4<f32>;
-
-fn main_1() {
- var count0 : i32;
- var count1 : i32;
- var i : i32;
- let x_29 : i32 = x_6.x_GLF_uniform_int_values[2].el;
- count0 = x_29;
- let x_31 : i32 = x_6.x_GLF_uniform_int_values[2].el;
- count1 = x_31;
- let x_33 : i32 = x_6.x_GLF_uniform_int_values[2].el;
- i = x_33;
- loop {
- let x_38 : i32 = i;
- let x_40 : i32 = x_6.x_GLF_uniform_int_values[4].el;
- if ((x_38 < x_40)) {
- } else {
- break;
- }
- let x_43 : i32 = i;
- switch(x_43) {
- case 0, 1: {
- let x_47 : i32 = count0;
- count0 = (x_47 + 1);
- fallthrough;
- }
- case 2, 3: {
- let x_49 : i32 = count1;
- count1 = (x_49 + 1);
- }
- default: {
- }
- }
-
- continuing {
- let x_51 : i32 = i;
- i = (x_51 + 1);
- }
- }
- let x_53 : i32 = count1;
- let x_55 : i32 = x_6.x_GLF_uniform_int_values[0].el;
- if ((x_53 == x_55)) {
- let x_61 : i32 = x_6.x_GLF_uniform_int_values[3].el;
- let x_64 : i32 = x_6.x_GLF_uniform_int_values[2].el;
- let x_67 : i32 = x_6.x_GLF_uniform_int_values[2].el;
- let x_70 : i32 = x_6.x_GLF_uniform_int_values[3].el;
- x_GLF_color = vec4<f32>(f32(x_61), f32(x_64), f32(x_67), f32(x_70));
- } else {
- let x_74 : i32 = x_6.x_GLF_uniform_int_values[2].el;
- let x_75 : f32 = f32(x_74);
- x_GLF_color = vec4<f32>(x_75, x_75, x_75, x_75);
- }
- let x_77 : i32 = count0;
- let x_79 : i32 = x_6.x_GLF_uniform_int_values[1].el;
- if ((x_77 != x_79)) {
- let x_84 : i32 = x_6.x_GLF_uniform_int_values[2].el;
- let x_85 : f32 = f32(x_84);
- x_GLF_color = vec4<f32>(x_85, x_85, x_85, x_85);
- }
- return;
-}
-
-struct main_out {
- @location(0)
- x_GLF_color_1 : vec4<f32>,
-}
-
-@fragment
-fn main() -> main_out {
- main_1();
- return main_out(x_GLF_color);
-}
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-loop-returns-behind-true-and-false/0-opt.spvasm b/test/tint/vk-gl-cts/graphicsfuzz/cov-loop-returns-behind-true-and-false/0-opt.spvasm
deleted file mode 100644
index 97d3744..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-loop-returns-behind-true-and-false/0-opt.spvasm
+++ /dev/null
@@ -1,117 +0,0 @@
- OpCapability Shader
- %1 = OpExtInstImport "GLSL.std.450"
- OpMemoryModel Logical GLSL450
- OpEntryPoint Fragment %main "main" %_GLF_color
- OpExecutionMode %main OriginUpperLeft
- OpSource ESSL 320
- OpName %main "main"
- OpName %_GLF_global_loop_count "_GLF_global_loop_count"
- OpName %buf0 "buf0"
- OpMemberName %buf0 0 "_GLF_uniform_int_values"
- OpName %_ ""
- OpName %_GLF_color "_GLF_color"
- OpDecorate %_arr_int_uint_2 ArrayStride 16
- OpMemberDecorate %buf0 0 Offset 0
- OpDecorate %buf0 Block
- OpDecorate %_ DescriptorSet 0
- OpDecorate %_ Binding 0
- OpDecorate %_GLF_color Location 0
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
- %int = OpTypeInt 32 1
-%_ptr_Private_int = OpTypePointer Private %int
-%_GLF_global_loop_count = OpVariable %_ptr_Private_int Private
- %int_0 = OpConstant %int 0
- %uint = OpTypeInt 32 0
- %uint_2 = OpConstant %uint 2
-%_arr_int_uint_2 = OpTypeArray %int %uint_2
- %buf0 = OpTypeStruct %_arr_int_uint_2
-%_ptr_Uniform_buf0 = OpTypePointer Uniform %buf0
- %_ = OpVariable %_ptr_Uniform_buf0 Uniform
-%_ptr_Uniform_int = OpTypePointer Uniform %int
- %bool = OpTypeBool
- %true = OpConstantTrue %bool
- %float = OpTypeFloat 32
- %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
- %_GLF_color = OpVariable %_ptr_Output_v4float Output
- %int_1 = OpConstant %int 1
- %false = OpConstantFalse %bool
- %int_100 = OpConstant %int 100
- %main = OpFunction %void None %9
- %25 = OpLabel
- OpStore %_GLF_global_loop_count %int_0
- %26 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_0
- %27 = OpLoad %int %26
- OpSelectionMerge %28 None
- OpSwitch %27 %28 0 %29 1 %30
- %29 = OpLabel
- OpSelectionMerge %31 None
- OpBranchConditional %true %32 %31
- %32 = OpLabel
- %33 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_1
- %34 = OpLoad %int %33
- %35 = OpConvertSToF %float %34
- %36 = OpCompositeConstruct %v4float %35 %35 %35 %35
- OpStore %_GLF_color %36
- OpReturn
- %31 = OpLabel
- OpBranch %30
- %30 = OpLabel
- OpSelectionMerge %37 None
- OpBranchConditional %true %38 %37
- %38 = OpLabel
- %39 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_0
- %40 = OpLoad %int %39
- %41 = OpConvertSToF %float %40
- %42 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_1
- %43 = OpLoad %int %42
- %44 = OpConvertSToF %float %43
- %45 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_1
- %46 = OpLoad %int %45
- %47 = OpConvertSToF %float %46
- %48 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_0
- %49 = OpLoad %int %48
- %50 = OpConvertSToF %float %49
- %51 = OpCompositeConstruct %v4float %41 %44 %47 %50
- OpStore %_GLF_color %51
- OpSelectionMerge %52 None
- OpBranchConditional %false %53 %52
- %53 = OpLabel
- %54 = OpAccessChain %_ptr_Uniform_int %_ %int_0 %int_0
- %55 = OpLoad %int %54
- %56 = OpConvertSToF %float %55
- %57 = OpCompositeConstruct %v4float %56 %56 %56 %56
- OpStore %_GLF_color %57
- OpBranch %58
- %58 = OpLabel
- OpLoopMerge %59 %60 None
- OpBranch %61
- %61 = OpLabel
- %62 = OpLoad %int %_GLF_global_loop_count
- %63 = OpIAdd %int %62 %int_1
- OpStore %_GLF_global_loop_count %63
- OpSelectionMerge %64 None
- OpBranchConditional %false %65 %64
- %65 = OpLabel
- OpReturn
- %64 = OpLabel
- OpSelectionMerge %66 None
- OpBranchConditional %true %67 %66
- %67 = OpLabel
- OpReturn
- %66 = OpLabel
- OpBranch %60
- %60 = OpLabel
- %68 = OpLoad %int %_GLF_global_loop_count
- %69 = OpSLessThan %bool %68 %int_100
- OpBranchConditional %69 %58 %59
- %59 = OpLabel
- OpBranch %52
- %52 = OpLabel
- OpBranch %37
- %37 = OpLabel
- OpBranch %28
- %28 = OpLabel
- OpReturn
- OpFunctionEnd
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-loop-returns-behind-true-and-false/0-opt.wgsl b/test/tint/vk-gl-cts/graphicsfuzz/cov-loop-returns-behind-true-and-false/0-opt.wgsl
deleted file mode 100644
index 7254319..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-loop-returns-behind-true-and-false/0-opt.wgsl
+++ /dev/null
@@ -1,78 +0,0 @@
-struct strided_arr {
- @size(16)
- el : i32,
-}
-
-type Arr = array<strided_arr, 2u>;
-
-struct buf0 {
- x_GLF_uniform_int_values : Arr,
-}
-
-var<private> x_GLF_global_loop_count : i32;
-
-@group(0) @binding(0) var<uniform> x_6 : buf0;
-
-var<private> x_GLF_color : vec4<f32>;
-
-fn main_1() {
- x_GLF_global_loop_count = 0;
- let x_27 : i32 = x_6.x_GLF_uniform_int_values[0].el;
- switch(x_27) {
- case 0: {
- if (true) {
- let x_34 : i32 = x_6.x_GLF_uniform_int_values[1].el;
- let x_35 : f32 = f32(x_34);
- x_GLF_color = vec4<f32>(x_35, x_35, x_35, x_35);
- return;
- }
- fallthrough;
- }
- case 1: {
- if (true) {
- let x_40 : i32 = x_6.x_GLF_uniform_int_values[0].el;
- let x_43 : i32 = x_6.x_GLF_uniform_int_values[1].el;
- let x_46 : i32 = x_6.x_GLF_uniform_int_values[1].el;
- let x_49 : i32 = x_6.x_GLF_uniform_int_values[0].el;
- x_GLF_color = vec4<f32>(f32(x_40), f32(x_43), f32(x_46), f32(x_49));
- if (false) {
- let x_55 : i32 = x_6.x_GLF_uniform_int_values[0].el;
- let x_56 : f32 = f32(x_55);
- x_GLF_color = vec4<f32>(x_56, x_56, x_56, x_56);
- loop {
- let x_62 : i32 = x_GLF_global_loop_count;
- x_GLF_global_loop_count = (x_62 + 1);
- if (false) {
- return;
- }
- if (true) {
- return;
- }
-
- continuing {
- let x_68 : i32 = x_GLF_global_loop_count;
- if ((x_68 < 100)) {
- } else {
- break;
- }
- }
- }
- }
- }
- }
- default: {
- }
- }
- return;
-}
-
-struct main_out {
- @location(0)
- x_GLF_color_1 : vec4<f32>,
-}
-
-@fragment
-fn main() -> main_out {
- main_1();
- return main_out(x_GLF_color);
-}
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-ssa-rewrite-case-with-default/0-opt.spvasm b/test/tint/vk-gl-cts/graphicsfuzz/cov-ssa-rewrite-case-with-default/0-opt.spvasm
deleted file mode 100644
index 5ef5478..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-ssa-rewrite-case-with-default/0-opt.spvasm
+++ /dev/null
@@ -1,92 +0,0 @@
- OpCapability Shader
- %1 = OpExtInstImport "GLSL.std.450"
- OpMemoryModel Logical GLSL450
- OpEntryPoint Fragment %main "main" %_GLF_color
- OpExecutionMode %main OriginUpperLeft
- OpSource ESSL 310
- OpName %main "main"
- OpName %func_ "func("
- OpName %buf0 "buf0"
- OpMemberName %buf0 0 "one"
- OpName %_ ""
- OpName %_GLF_color "_GLF_color"
- OpName %i "i"
- OpMemberDecorate %buf0 0 Offset 0
- OpDecorate %buf0 Block
- OpDecorate %_ DescriptorSet 0
- OpDecorate %_ Binding 0
- OpDecorate %_GLF_color Location 0
- %void = OpTypeVoid
- %9 = OpTypeFunction %void
- %float = OpTypeFloat 32
- %v4float = OpTypeVector %float 4
- %12 = OpTypeFunction %v4float
- %int = OpTypeInt 32 1
- %buf0 = OpTypeStruct %int
-%_ptr_Uniform_buf0 = OpTypePointer Uniform %buf0
- %_ = OpVariable %_ptr_Uniform_buf0 Uniform
- %int_0 = OpConstant %int 0
-%_ptr_Uniform_int = OpTypePointer Uniform %int
- %int_1 = OpConstant %int 1
- %bool = OpTypeBool
- %float_1 = OpConstant %float 1
- %float_0 = OpConstant %float 0
- %21 = OpConstantComposite %v4float %float_1 %float_0 %float_0 %float_1
- %22 = OpConstantComposite %v4float %float_0 %float_0 %float_0 %float_0
-%_ptr_Output_v4float = OpTypePointer Output %v4float
- %_GLF_color = OpVariable %_ptr_Output_v4float Output
-%_ptr_Function_int = OpTypePointer Function %int
- %uint = OpTypeInt 32 0
- %uint_1 = OpConstant %uint 1
-%_ptr_Output_float = OpTypePointer Output %float
- %main = OpFunction %void None %9
- %28 = OpLabel
- %i = OpVariable %_ptr_Function_int Function
- OpStore %_GLF_color %22
- OpStore %i %int_0
- OpBranch %29
- %29 = OpLabel
- OpLoopMerge %30 %31 None
- OpBranch %32
- %32 = OpLabel
- %33 = OpLoad %int %i
- %34 = OpAccessChain %_ptr_Uniform_int %_ %int_0
- %35 = OpLoad %int %34
- %36 = OpSLessThanEqual %bool %33 %35
- OpBranchConditional %36 %37 %30
- %37 = OpLabel
- %38 = OpLoad %int %i
- OpSelectionMerge %39 None
- OpSwitch %38 %40 1 %41 0 %40
- %40 = OpLabel
- %42 = OpAccessChain %_ptr_Output_float %_GLF_color %uint_1
- OpStore %42 %float_0
- OpBranch %39
- %41 = OpLabel
- %43 = OpFunctionCall %v4float %func_
- OpStore %_GLF_color %43
- OpBranch %40
- %39 = OpLabel
- OpBranch %31
- %31 = OpLabel
- %44 = OpLoad %int %i
- %45 = OpIAdd %int %44 %int_1
- OpStore %i %45
- OpBranch %29
- %30 = OpLabel
- OpReturn
- OpFunctionEnd
- %func_ = OpFunction %v4float None %12
- %46 = OpLabel
- %47 = OpAccessChain %_ptr_Uniform_int %_ %int_0
- %48 = OpLoad %int %47
- %49 = OpIEqual %bool %48 %int_1
- OpSelectionMerge %50 None
- OpBranchConditional %49 %51 %52
- %51 = OpLabel
- OpReturnValue %21
- %52 = OpLabel
- OpReturnValue %22
- %50 = OpLabel
- OpUnreachable
- OpFunctionEnd
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-ssa-rewrite-case-with-default/0-opt.wgsl b/test/tint/vk-gl-cts/graphicsfuzz/cov-ssa-rewrite-case-with-default/0-opt.wgsl
deleted file mode 100644
index 0cbcc97..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-ssa-rewrite-case-with-default/0-opt.wgsl
+++ /dev/null
@@ -1,61 +0,0 @@
-struct buf0 {
- one : i32,
-}
-
-@group(0) @binding(0) var<uniform> x_6 : buf0;
-
-var<private> x_GLF_color : vec4<f32>;
-
-fn func_() -> vec4<f32> {
- let x_48 : i32 = x_6.one;
- if ((x_48 == 1)) {
- return vec4<f32>(1.0, 0.0, 0.0, 1.0);
- } else {
- return vec4<f32>(0.0, 0.0, 0.0, 0.0);
- }
-}
-
-fn main_1() {
- var i : i32;
- x_GLF_color = vec4<f32>(0.0, 0.0, 0.0, 0.0);
- i = 0;
- loop {
- let x_33 : i32 = i;
- let x_35 : i32 = x_6.one;
- if ((x_33 <= x_35)) {
- } else {
- break;
- }
- let x_38 : i32 = i;
- switch(x_38) {
- case 1: {
- let x_43 : vec4<f32> = func_();
- x_GLF_color = x_43;
- fallthrough;
- }
- default: {
- fallthrough;
- }
- case 0: {
- x_GLF_color.y = 0.0;
- }
- }
-
- continuing {
- let x_44 : i32 = i;
- i = (x_44 + 1);
- }
- }
- return;
-}
-
-struct main_out {
- @location(0)
- x_GLF_color_1 : vec4<f32>,
-}
-
-@fragment
-fn main() -> main_out {
- main_1();
- return main_out(x_GLF_color);
-}
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-val-cfg-case-fallthrough/0-opt.spvasm b/test/tint/vk-gl-cts/graphicsfuzz/cov-val-cfg-case-fallthrough/0-opt.spvasm
deleted file mode 100644
index 5591664..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-val-cfg-case-fallthrough/0-opt.spvasm
+++ /dev/null
@@ -1,67 +0,0 @@
- OpCapability Shader
- %1 = OpExtInstImport "GLSL.std.450"
- OpMemoryModel Logical GLSL450
- OpEntryPoint Fragment %main "main" %_GLF_color
- OpExecutionMode %main OriginUpperLeft
- OpSource ESSL 310
- OpName %main "main"
- OpName %a "a"
- OpName %buf0 "buf0"
- OpMemberName %buf0 0 "one"
- OpName %_ ""
- OpName %_GLF_color "_GLF_color"
- OpMemberDecorate %buf0 0 Offset 0
- OpDecorate %buf0 Block
- OpDecorate %_ DescriptorSet 0
- OpDecorate %_ Binding 0
- OpDecorate %_GLF_color Location 0
- %void = OpTypeVoid
- %8 = OpTypeFunction %void
- %int = OpTypeInt 32 1
-%_ptr_Function_int = OpTypePointer Function %int
- %int_0 = OpConstant %int 0
- %buf0 = OpTypeStruct %int
-%_ptr_Uniform_buf0 = OpTypePointer Uniform %buf0
- %_ = OpVariable %_ptr_Uniform_buf0 Uniform
-%_ptr_Uniform_int = OpTypePointer Uniform %int
- %int_1 = OpConstant %int 1
- %int_2 = OpConstant %int 2
- %bool = OpTypeBool
- %float = OpTypeFloat 32
- %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
- %_GLF_color = OpVariable %_ptr_Output_v4float Output
- %float_1 = OpConstant %float 1
- %float_0 = OpConstant %float 0
- %22 = OpConstantComposite %v4float %float_1 %float_0 %float_0 %float_1
- %23 = OpConstantComposite %v4float %float_0 %float_0 %float_0 %float_0
- %main = OpFunction %void None %8
- %24 = OpLabel
- %a = OpVariable %_ptr_Function_int Function
- OpStore %a %int_0
- %25 = OpAccessChain %_ptr_Uniform_int %_ %int_0
- %26 = OpLoad %int %25
- OpSelectionMerge %27 None
- OpSwitch %26 %28 2 %29 3 %29 4 %30
- %28 = OpLabel
- OpStore %a %int_2
- OpBranch %27
- %29 = OpLabel
- OpStore %a %int_1
- OpBranch %30
- %30 = OpLabel
- OpBranch %27
- %27 = OpLabel
- %31 = OpLoad %int %a
- %32 = OpIEqual %bool %31 %int_2
- OpSelectionMerge %33 None
- OpBranchConditional %32 %34 %35
- %34 = OpLabel
- OpStore %_GLF_color %22
- OpBranch %33
- %35 = OpLabel
- OpStore %_GLF_color %23
- OpBranch %33
- %33 = OpLabel
- OpReturn
- OpFunctionEnd
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/cov-val-cfg-case-fallthrough/0-opt.wgsl b/test/tint/vk-gl-cts/graphicsfuzz/cov-val-cfg-case-fallthrough/0-opt.wgsl
deleted file mode 100644
index e58d17e..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/cov-val-cfg-case-fallthrough/0-opt.wgsl
+++ /dev/null
@@ -1,42 +0,0 @@
-struct buf0 {
- one : i32,
-}
-
-@group(0) @binding(0) var<uniform> x_6 : buf0;
-
-var<private> x_GLF_color : vec4<f32>;
-
-fn main_1() {
- var a : i32;
- a = 0;
- let x_26 : i32 = x_6.one;
- switch(x_26) {
- case 2, 3: {
- a = 1;
- fallthrough;
- }
- case 4: {
- }
- default: {
- a = 2;
- }
- }
- let x_31 : i32 = a;
- if ((x_31 == 2)) {
- x_GLF_color = vec4<f32>(1.0, 0.0, 0.0, 1.0);
- } else {
- x_GLF_color = vec4<f32>(0.0, 0.0, 0.0, 0.0);
- }
- return;
-}
-
-struct main_out {
- @location(0)
- x_GLF_color_1 : vec4<f32>,
-}
-
-@fragment
-fn main() -> main_out {
- main_1();
- return main_out(x_GLF_color);
-}
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/if-and-switch/0.spvasm b/test/tint/vk-gl-cts/graphicsfuzz/if-and-switch/0.spvasm
deleted file mode 100644
index 28e20bd..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/if-and-switch/0.spvasm
+++ /dev/null
@@ -1,74 +0,0 @@
- OpCapability Shader
- %1 = OpExtInstImport "GLSL.std.450"
- OpMemoryModel Logical GLSL450
- OpEntryPoint Fragment %main "main" %_GLF_color
- OpExecutionMode %main OriginUpperLeft
- OpSource ESSL 310
- OpName %main "main"
- OpName %data "data"
- OpName %buf0 "buf0"
- OpMemberName %buf0 0 "injectionSwitch"
- OpName %_ ""
- OpName %_GLF_color "_GLF_color"
- OpMemberDecorate %buf0 0 Offset 0
- OpDecorate %buf0 Block
- OpDecorate %_ DescriptorSet 0
- OpDecorate %_ Binding 0
- OpDecorate %_GLF_color Location 0
- %void = OpTypeVoid
- %8 = OpTypeFunction %void
- %float = OpTypeFloat 32
- %uint = OpTypeInt 32 0
- %uint_2 = OpConstant %uint 2
-%_arr_float_uint_2 = OpTypeArray %float %uint_2
-%_ptr_Function__arr_float_uint_2 = OpTypePointer Function %_arr_float_uint_2
- %int = OpTypeInt 32 1
- %int_0 = OpConstant %int 0
- %v2float = OpTypeVector %float 2
- %buf0 = OpTypeStruct %v2float
-%_ptr_Uniform_buf0 = OpTypePointer Uniform %buf0
- %_ = OpVariable %_ptr_Uniform_buf0 Uniform
- %uint_0 = OpConstant %uint 0
-%_ptr_Uniform_float = OpTypePointer Uniform %float
-%_ptr_Function_float = OpTypePointer Function %float
- %int_1 = OpConstant %int 1
- %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
- %_GLF_color = OpVariable %_ptr_Output_v4float Output
- %float_1 = OpConstant %float 1
- %float_0 = OpConstant %float 0
- %26 = OpConstantComposite %v4float %float_1 %float_0 %float_0 %float_1
- %bool = OpTypeBool
- %28 = OpConstantComposite %v4float %float_0 %float_0 %float_0 %float_0
- %29 = OpUndef %float
- %main = OpFunction %void None %8
- %30 = OpLabel
- %data = OpVariable %_ptr_Function__arr_float_uint_2 Function
- %31 = OpAccessChain %_ptr_Uniform_float %_ %int_0 %uint_0
- %32 = OpLoad %float %31
- %33 = OpAccessChain %_ptr_Function_float %data %int_0
- OpStore %33 %32
- %34 = OpAccessChain %_ptr_Function_float %data %int_1
- OpStore %34 %32
- OpStore %_GLF_color %26
- %35 = OpLoad %float %34
- %36 = OpFOrdGreaterThan %bool %35 %float_1
- OpSelectionMerge %37 None
- OpBranchConditional %36 %38 %37
- %38 = OpLabel
- %39 = OpConvertFToS %int %32
- OpSelectionMerge %40 None
- OpSwitch %39 %40 0 %41 1 %42
- %40 = OpLabel
- OpBranch %37
- %41 = OpLabel
- OpBranch %42
- %42 = OpLabel
- %43 = OpPhi %float %29 %38 %float_1 %41
- %44 = OpAccessChain %_ptr_Function_float %data %39
- OpStore %44 %43
- OpStore %_GLF_color %28
- OpBranch %40
- %37 = OpLabel
- OpReturn
- OpFunctionEnd
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/if-and-switch/0.wgsl b/test/tint/vk-gl-cts/graphicsfuzz/if-and-switch/0.wgsl
deleted file mode 100644
index d9cf1c7..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/if-and-switch/0.wgsl
+++ /dev/null
@@ -1,47 +0,0 @@
-struct buf0 {
- injectionSwitch : vec2<f32>,
-}
-
-@group(0) @binding(0) var<uniform> x_6 : buf0;
-
-var<private> x_GLF_color : vec4<f32>;
-
-fn main_1() {
- var data : array<f32, 2u>;
- var x_32 : f32;
- x_32 = x_6.injectionSwitch.x;
- data[0] = x_32;
- data[1] = x_32;
- x_GLF_color = vec4<f32>(1.0, 0.0, 0.0, 1.0);
- let x_35 : f32 = data[1];
- if ((x_35 > 1.0)) {
- var x_43_phi : f32;
- let x_39 : i32 = i32(x_32);
- x_43_phi = 0.0;
- switch(x_39) {
- case 0: {
- x_43_phi = 1.0;
- fallthrough;
- }
- case 1: {
- let x_43 : f32 = x_43_phi;
- data[x_39] = x_43;
- x_GLF_color = vec4<f32>(0.0, 0.0, 0.0, 0.0);
- }
- default: {
- }
- }
- }
- return;
-}
-
-struct main_out {
- @location(0)
- x_GLF_color_1 : vec4<f32>,
-}
-
-@fragment
-fn main() -> main_out {
- main_1();
- return main_out(x_GLF_color);
-}
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/nested-switch-break-discard/0-opt.spvasm b/test/tint/vk-gl-cts/graphicsfuzz/nested-switch-break-discard/0-opt.spvasm
deleted file mode 100644
index 7a18a05..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/nested-switch-break-discard/0-opt.spvasm
+++ /dev/null
@@ -1,72 +0,0 @@
- OpCapability Shader
- %1 = OpExtInstImport "GLSL.std.450"
- OpMemoryModel Logical GLSL450
- OpEntryPoint Fragment %main "main" %_GLF_color %gl_FragCoord
- OpExecutionMode %main OriginUpperLeft
- OpSource ESSL 310
- OpName %main "main"
- OpName %_GLF_color "_GLF_color"
- OpName %buf0 "buf0"
- OpMemberName %buf0 0 "injectionSwitch"
- OpName %_ ""
- OpName %gl_FragCoord "gl_FragCoord"
- OpDecorate %_GLF_color Location 0
- OpMemberDecorate %buf0 0 Offset 0
- OpDecorate %buf0 Block
- OpDecorate %_ DescriptorSet 0
- OpDecorate %_ Binding 0
- OpDecorate %gl_FragCoord BuiltIn FragCoord
- %void = OpTypeVoid
- %8 = OpTypeFunction %void
- %float = OpTypeFloat 32
- %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
- %_GLF_color = OpVariable %_ptr_Output_v4float Output
- %float_0 = OpConstant %float 0
- %float_1 = OpConstant %float 1
- %14 = OpConstantComposite %v4float %float_0 %float_0 %float_0 %float_1
- %v2float = OpTypeVector %float 2
- %buf0 = OpTypeStruct %v2float
-%_ptr_Uniform_buf0 = OpTypePointer Uniform %buf0
- %_ = OpVariable %_ptr_Uniform_buf0 Uniform
- %int = OpTypeInt 32 1
- %int_0 = OpConstant %int 0
- %uint = OpTypeInt 32 0
- %uint_0 = OpConstant %uint 0
-%_ptr_Uniform_float = OpTypePointer Uniform %float
- %int_1 = OpConstant %int 1
-%_ptr_Input_v4float = OpTypePointer Input %v4float
-%gl_FragCoord = OpVariable %_ptr_Input_v4float Input
- %uint_1 = OpConstant %uint 1
-%_ptr_Input_float = OpTypePointer Input %float
- %bool = OpTypeBool
- %27 = OpConstantComposite %v4float %float_1 %float_0 %float_0 %float_1
- %main = OpFunction %void None %8
- %28 = OpLabel
- OpStore %_GLF_color %14
- %29 = OpAccessChain %_ptr_Uniform_float %_ %int_0 %uint_0
- %30 = OpLoad %float %29
- %31 = OpConvertFToS %int %30
- OpSelectionMerge %32 None
- OpSwitch %31 %32 0 %33 42 %34
- %33 = OpLabel
- OpSelectionMerge %35 None
- OpSwitch %int_1 %35 1 %36
- %36 = OpLabel
- %37 = OpAccessChain %_ptr_Input_float %gl_FragCoord %uint_1
- %38 = OpLoad %float %37
- %39 = OpFOrdGreaterThanEqual %bool %38 %float_0
- OpSelectionMerge %40 None
- OpBranchConditional %39 %41 %40
- %41 = OpLabel
- OpStore %_GLF_color %27
- OpBranch %35
- %40 = OpLabel
- OpKill
- %35 = OpLabel
- OpBranch %34
- %34 = OpLabel
- OpBranch %32
- %32 = OpLabel
- OpReturn
- OpFunctionEnd
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/nested-switch-break-discard/0-opt.wgsl b/test/tint/vk-gl-cts/graphicsfuzz/nested-switch-break-discard/0-opt.wgsl
deleted file mode 100644
index f932ed9..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/nested-switch-break-discard/0-opt.wgsl
+++ /dev/null
@@ -1,48 +0,0 @@
-struct buf0 {
- injectionSwitch : vec2<f32>,
-}
-
-var<private> x_GLF_color : vec4<f32>;
-
-@group(0) @binding(0) var<uniform> x_6 : buf0;
-
-var<private> gl_FragCoord : vec4<f32>;
-
-fn main_1() {
- x_GLF_color = vec4<f32>(0.0, 0.0, 0.0, 1.0);
- let x_30 : f32 = x_6.injectionSwitch.x;
- switch(i32(x_30)) {
- case 0: {
- switch(1) {
- case 1: {
- let x_38 : f32 = gl_FragCoord.y;
- if ((x_38 >= 0.0)) {
- x_GLF_color = vec4<f32>(1.0, 0.0, 0.0, 1.0);
- break;
- }
- discard;
- }
- default: {
- }
- }
- fallthrough;
- }
- case 42: {
- }
- default: {
- }
- }
- return;
-}
-
-struct main_out {
- @location(0)
- x_GLF_color_1 : vec4<f32>,
-}
-
-@fragment
-fn main(@builtin(position) gl_FragCoord_param : vec4<f32>) -> main_out {
- gl_FragCoord = gl_FragCoord_param;
- main_1();
- return main_out(x_GLF_color);
-}
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/switch-loop-switch-if/0-opt.spvasm b/test/tint/vk-gl-cts/graphicsfuzz/switch-loop-switch-if/0-opt.spvasm
deleted file mode 100644
index 099b620..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/switch-loop-switch-if/0-opt.spvasm
+++ /dev/null
@@ -1,134 +0,0 @@
- OpCapability Shader
- %1 = OpExtInstImport "GLSL.std.450"
- OpMemoryModel Logical GLSL450
- OpEntryPoint Fragment %main "main" %_GLF_color
- OpExecutionMode %main OriginUpperLeft
- OpSource ESSL 310
- OpName %main "main"
- OpName %i "i"
- OpName %buf0 "buf0"
- OpMemberName %buf0 0 "injectionSwitch"
- OpName %_ ""
- OpName %_GLF_color "_GLF_color"
- OpDecorate %i RelaxedPrecision
- OpMemberDecorate %buf0 0 Offset 0
- OpDecorate %buf0 Block
- OpDecorate %_ DescriptorSet 0
- OpDecorate %_ Binding 0
- OpDecorate %7 RelaxedPrecision
- OpDecorate %8 RelaxedPrecision
- OpDecorate %9 RelaxedPrecision
- OpDecorate %10 RelaxedPrecision
- OpDecorate %11 RelaxedPrecision
- OpDecorate %12 RelaxedPrecision
- OpDecorate %13 RelaxedPrecision
- OpDecorate %14 RelaxedPrecision
- OpDecorate %15 RelaxedPrecision
- OpDecorate %16 RelaxedPrecision
- OpDecorate %17 RelaxedPrecision
- OpDecorate %18 RelaxedPrecision
- OpDecorate %19 RelaxedPrecision
- OpDecorate %20 RelaxedPrecision
- OpDecorate %21 RelaxedPrecision
- OpDecorate %22 RelaxedPrecision
- OpDecorate %_GLF_color Location 0
- %void = OpTypeVoid
- %24 = OpTypeFunction %void
- %int = OpTypeInt 32 1
-%_ptr_Function_int = OpTypePointer Function %int
- %float = OpTypeFloat 32
- %v2float = OpTypeVector %float 2
- %buf0 = OpTypeStruct %v2float
-%_ptr_Uniform_buf0 = OpTypePointer Uniform %buf0
- %_ = OpVariable %_ptr_Uniform_buf0 Uniform
- %int_0 = OpConstant %int 0
- %uint = OpTypeInt 32 0
- %uint_0 = OpConstant %uint 0
-%_ptr_Uniform_float = OpTypePointer Uniform %float
- %int_1 = OpConstant %int 1
- %int_5 = OpConstant %int 5
- %int_7 = OpConstant %int 7
- %int_200 = OpConstant %int 200
- %bool = OpTypeBool
- %int_100 = OpConstant %int 100
- %int_2 = OpConstant %int 2
- %int_3 = OpConstant %int 3
- %int_n2 = OpConstant %int -2
- %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
- %_GLF_color = OpVariable %_ptr_Output_v4float Output
- %float_1 = OpConstant %float 1
- %float_0 = OpConstant %float 0
- %47 = OpConstantComposite %v4float %float_1 %float_0 %float_0 %float_1
- %48 = OpConstantComposite %v4float %float_0 %float_0 %float_0 %float_1
- %main = OpFunction %void None %24
- %49 = OpLabel
- %i = OpVariable %_ptr_Function_int Function
- %50 = OpAccessChain %_ptr_Uniform_float %_ %int_0 %uint_0
- %51 = OpLoad %float %50
- %7 = OpConvertFToS %int %51
- OpStore %i %7
- %8 = OpLoad %int %i
- OpSelectionMerge %52 None
- OpSwitch %8 %53 0 %54
- %53 = OpLabel
- %20 = OpLoad %int %i
- %21 = OpISub %int %20 %int_3
- OpStore %i %21
- OpBranch %52
- %54 = OpLabel
- OpBranch %55
- %55 = OpLabel
- OpLoopMerge %56 %57 None
- OpBranch %58
- %58 = OpLabel
- %9 = OpLoad %int %i
- %10 = OpIAdd %int %9 %int_1
- OpStore %i %10
- %11 = OpLoad %int %i
- OpSelectionMerge %59 None
- OpSwitch %11 %60 1 %61 2 %62
- %60 = OpLabel
- %14 = OpLoad %int %i
- %15 = OpIAdd %int %14 %int_7
- OpStore %i %15
- OpBranch %59
- %61 = OpLabel
- OpBranch %57
- %62 = OpLabel
- %12 = OpLoad %int %i
- %13 = OpIAdd %int %12 %int_5
- OpStore %i %13
- OpBranch %59
- %59 = OpLabel
- OpBranch %57
- %57 = OpLabel
- %16 = OpLoad %int %i
- %63 = OpSGreaterThan %bool %16 %int_200
- OpBranchConditional %63 %55 %56
- %56 = OpLabel
- %17 = OpLoad %int %i
- %64 = OpSGreaterThan %bool %17 %int_100
- OpSelectionMerge %65 None
- OpBranchConditional %64 %66 %65
- %66 = OpLabel
- %18 = OpLoad %int %i
- %19 = OpISub %int %18 %int_2
- OpStore %i %19
- OpBranch %52
- %65 = OpLabel
- OpBranch %53
- %52 = OpLabel
- %22 = OpLoad %int %i
- %67 = OpIEqual %bool %22 %int_n2
- OpSelectionMerge %68 None
- OpBranchConditional %67 %69 %70
- %69 = OpLabel
- OpStore %_GLF_color %47
- OpBranch %68
- %70 = OpLabel
- OpStore %_GLF_color %48
- OpBranch %68
- %68 = OpLabel
- OpReturn
- OpFunctionEnd
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/switch-loop-switch-if/0-opt.wgsl b/test/tint/vk-gl-cts/graphicsfuzz/switch-loop-switch-if/0-opt.wgsl
deleted file mode 100644
index 41e5cfe..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/switch-loop-switch-if/0-opt.wgsl
+++ /dev/null
@@ -1,73 +0,0 @@
-struct buf0 {
- injectionSwitch : vec2<f32>,
-}
-
-@group(0) @binding(0) var<uniform> x_6 : buf0;
-
-var<private> x_GLF_color : vec4<f32>;
-
-fn main_1() {
- var i : i32;
- let x_51 : f32 = x_6.injectionSwitch.x;
- i = i32(x_51);
- let x_8 : i32 = i;
- switch(x_8) {
- case 0: {
- loop {
- let x_9 : i32 = i;
- i = (x_9 + 1);
- let x_11 : i32 = i;
- switch(x_11) {
- case 2: {
- let x_12 : i32 = i;
- i = (x_12 + 5);
- }
- case 1: {
- continue;
- }
- default: {
- let x_14 : i32 = i;
- i = (x_14 + 7);
- }
- }
-
- continuing {
- let x_16 : i32 = i;
- if ((x_16 > 200)) {
- } else {
- break;
- }
- }
- }
- let x_17 : i32 = i;
- if ((x_17 > 100)) {
- let x_18 : i32 = i;
- i = (x_18 - 2);
- break;
- }
- fallthrough;
- }
- default: {
- let x_20 : i32 = i;
- i = (x_20 - 3);
- }
- }
- let x_22 : i32 = i;
- if ((x_22 == -2)) {
- x_GLF_color = vec4<f32>(1.0, 0.0, 0.0, 1.0);
- } else {
- x_GLF_color = vec4<f32>(0.0, 0.0, 0.0, 1.0);
- }
- return;
-}
-
-struct main_out {
- @location(0)
- x_GLF_color_1 : vec4<f32>,
-}
-
-@fragment
-fn main() -> main_out {
- main_1();
- return main_out(x_GLF_color);
-}
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/switch-with-fall-through-cases/0-opt.spvasm b/test/tint/vk-gl-cts/graphicsfuzz/switch-with-fall-through-cases/0-opt.spvasm
deleted file mode 100644
index 20a942d..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/switch-with-fall-through-cases/0-opt.spvasm
+++ /dev/null
@@ -1,96 +0,0 @@
- OpCapability Shader
- %1 = OpExtInstImport "GLSL.std.450"
- OpMemoryModel Logical GLSL450
- OpEntryPoint Fragment %main "main" %_GLF_color
- OpExecutionMode %main OriginUpperLeft
- OpSource ESSL 310
- OpName %main "main"
- OpName %i "i"
- OpName %buf0 "buf0"
- OpMemberName %buf0 0 "injectionSwitch"
- OpName %_ ""
- OpName %value "value"
- OpName %y "y"
- OpName %_GLF_color "_GLF_color"
- OpMemberDecorate %buf0 0 Offset 0
- OpDecorate %buf0 Block
- OpDecorate %_ DescriptorSet 0
- OpDecorate %_ Binding 0
- OpDecorate %_GLF_color Location 0
- %void = OpTypeVoid
- %10 = OpTypeFunction %void
- %int = OpTypeInt 32 1
-%_ptr_Function_int = OpTypePointer Function %int
- %int_0 = OpConstant %int 0
- %int_2 = OpConstant %int 2
- %float = OpTypeFloat 32
- %v2float = OpTypeVector %float 2
- %buf0 = OpTypeStruct %v2float
-%_ptr_Uniform_buf0 = OpTypePointer Uniform %buf0
- %_ = OpVariable %_ptr_Uniform_buf0 Uniform
- %uint = OpTypeInt 32 0
- %uint_0 = OpConstant %uint 0
-%_ptr_Uniform_float = OpTypePointer Uniform %float
- %bool = OpTypeBool
-%_ptr_Function_float = OpTypePointer Function %float
- %float_0_5 = OpConstant %float 0.5
- %float_1 = OpConstant %float 1
- %v4float = OpTypeVector %float 4
-%_ptr_Output_v4float = OpTypePointer Output %v4float
- %_GLF_color = OpVariable %_ptr_Output_v4float Output
- %int_1 = OpConstant %int 1
- %float_0 = OpConstant %float 0
- %main = OpFunction %void None %10
- %29 = OpLabel
- %i = OpVariable %_ptr_Function_int Function
- %value = OpVariable %_ptr_Function_int Function
- %y = OpVariable %_ptr_Function_float Function
- OpStore %i %int_0
- OpBranch %30
- %30 = OpLabel
- %31 = OpPhi %int %int_0 %29 %32 %33
- OpLoopMerge %34 %33 None
- OpBranch %35
- %35 = OpLabel
- %36 = OpAccessChain %_ptr_Uniform_float %_ %int_0 %uint_0
- %37 = OpLoad %float %36
- %38 = OpConvertFToS %int %37
- %39 = OpIAdd %int %int_2 %38
- %40 = OpSLessThan %bool %31 %39
- OpBranchConditional %40 %41 %34
- %41 = OpLabel
- OpStore %value %31
- OpStore %y %float_0_5
- OpSelectionMerge %42 None
- OpSwitch %31 %43 0 %44 1 %45 2 %43
- %43 = OpLabel
- %46 = OpPhi %float %float_0_5 %41 %47 %45
- %48 = OpFOrdEqual %bool %46 %float_1
- OpSelectionMerge %49 None
- OpBranchConditional %48 %50 %49
- %50 = OpLabel
- %51 = OpIAdd %int %31 %int_1
- %52 = OpConvertSToF %float %51
- %53 = OpCompositeConstruct %v4float %52 %float_0 %float_0 %float_1
- OpStore %_GLF_color %53
- OpReturn
- %49 = OpLabel
- OpBranch %42
- %44 = OpLabel
- %54 = OpFAdd %float %float_0_5 %float_0_5
- OpStore %y %54
- OpBranch %45
- %45 = OpLabel
- %55 = OpPhi %float %float_0_5 %41 %54 %44
- %47 = OpExtInst %float %1 FClamp %float_1 %float_0_5 %55
- OpStore %y %47
- OpBranch %43
- %42 = OpLabel
- OpBranch %33
- %33 = OpLabel
- %32 = OpIAdd %int %31 %int_1
- OpStore %i %32
- OpBranch %30
- %34 = OpLabel
- OpReturn
- OpFunctionEnd
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/switch-with-fall-through-cases/0-opt.wgsl b/test/tint/vk-gl-cts/graphicsfuzz/switch-with-fall-through-cases/0-opt.wgsl
deleted file mode 100644
index 44c1c31..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/switch-with-fall-through-cases/0-opt.wgsl
+++ /dev/null
@@ -1,73 +0,0 @@
-struct buf0 {
- injectionSwitch : vec2<f32>,
-}
-
-@group(0) @binding(0) var<uniform> x_6 : buf0;
-
-var<private> x_GLF_color : vec4<f32>;
-
-fn main_1() {
- var i : i32;
- var value : i32;
- var y : f32;
- var x_31_phi : i32;
- i = 0;
- x_31_phi = 0;
- loop {
- let x_31 : i32 = x_31_phi;
- let x_37 : f32 = x_6.injectionSwitch.x;
- if ((x_31 < (2 + i32(x_37)))) {
- } else {
- break;
- }
- var x_55_phi : f32;
- var x_46_phi : f32;
- value = x_31;
- y = 0.5;
- x_55_phi = 0.5;
- x_46_phi = 0.5;
- switch(x_31) {
- case 0: {
- let x_54 : f32 = (0.5 + 0.5);
- y = x_54;
- x_55_phi = x_54;
- fallthrough;
- }
- case 1: {
- let x_55 : f32 = x_55_phi;
- let x_47 : f32 = clamp(1.0, 0.5, x_55);
- y = x_47;
- x_46_phi = x_47;
- fallthrough;
- }
- default: {
- fallthrough;
- }
- case 2: {
- let x_46 : f32 = x_46_phi;
- if ((x_46 == 1.0)) {
- x_GLF_color = vec4<f32>(f32((x_31 + 1)), 0.0, 0.0, 1.0);
- return;
- }
- }
- }
-
- continuing {
- let x_32 : i32 = (x_31 + 1);
- i = x_32;
- x_31_phi = x_32;
- }
- }
- return;
-}
-
-struct main_out {
- @location(0)
- x_GLF_color_1 : vec4<f32>,
-}
-
-@fragment
-fn main() -> main_out {
- main_1();
- return main_out(x_GLF_color);
-}
diff --git a/test/tint/vk-gl-cts/graphicsfuzz/unreachable-loops-in-switch/0.wgsl b/test/tint/vk-gl-cts/graphicsfuzz/unreachable-loops-in-switch/0.wgsl
deleted file mode 100644
index 5b160e1..0000000
--- a/test/tint/vk-gl-cts/graphicsfuzz/unreachable-loops-in-switch/0.wgsl
+++ /dev/null
@@ -1,63 +0,0 @@
-var<private> x_GLF_color : vec4<f32>;
-
-fn main_1() {
- var i : i32;
- var data : array<f32, 1>;
- x_GLF_color = vec4<f32>(1.0, 0.0, 0.0, 1.0);
- i = 0;
- loop {
- let x_6 : i32 = i;
- if ((x_6 < 1)) {
- } else {
- break;
- }
- let x_7 : i32 = i;
- let x_40 : f32 = data[x_7];
- let x_42 : f32 = data[0];
- if ((x_40 < x_42)) {
- if (false) {
- let x_8 : i32 = i;
- if ((f32(x_8) >= 1.0)) {
- }
- }
- switch(0) {
- case 1: {
- loop {
- if (true) {
- } else {
- break;
- }
- }
- loop {
- break;
- }
- fallthrough;
- }
- case 0: {
- data[0] = 2.0;
- }
- default: {
- }
- }
-
- break;
- }
-
- continuing {
- let x_9 : i32 = i;
- i = (x_9 + 1);
- }
- }
- return;
-}
-
-struct main_out {
- @location(0)
- x_GLF_color_1 : vec4<f32>,
-};
-
-@fragment
-fn main() -> main_out {
- main_1();
- return main_out(x_GLF_color);
-}
