diff --git a/src/tint/reader/spirv/function.cc b/src/tint/reader/spirv/function.cc index 07b5980c3d..a71aa42a66 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 @@ bool FunctionEmitter::ClassifyCFGEdges() { // 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 @@ bool FunctionEmitter::ClassifyCFGEdges() { // 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 normal_forward_edges; utils::Vector if_break_edges; @@ -2376,6 +2373,12 @@ bool FunctionEmitter::ClassifyCFGEdges() { << 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 @@ bool FunctionEmitter::EmitNormalTerminator(const BlockInfo& block_info) { // 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 @@ bool FunctionEmitter::EmitNormalTerminator(const BlockInfo& block_info) { 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 @@ const ast::Statement* FunctionEmitter::MakeBranchDetailed(const BlockInfo& src_i // merge block is implicit. break; } - case EdgeKind::kCaseFallThrough: - return create(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 @@ const ast::Statement* FunctionEmitter::MakeSimpleIf(const ast::Expression* condi 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(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 8d06735644..917b37485b 100644 --- a/src/tint/reader/spirv/function.h +++ b/src/tint/reader/spirv/function.h @@ -34,15 +34,13 @@ namespace tint::reader::spirv { // // 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 @@ enum class EdgeKind { // 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 @@ class FunctionEmitter { /// 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 @@ class FunctionEmitter { /// @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 @@ class FunctionEmitter { /// @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 @@ class FunctionEmitter { 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 538657b495..88dc78eba1 100644 --- a/src/tint/reader/spirv/function_cfg_test.cc +++ b/src/tint/reader/spirv/function_cfg_test.cc @@ -1288,43 +1288,7 @@ TEST_F(SpvParserCFGTest, ComputeBlockOrder_Switch_DefaultSameAsACase) { 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 @@ TEST_F(SpvParserCFGTest, ComputeBlockOrder_RespectSwitchCaseFallthrough_FromDefa 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 @@ TEST_F(SpvParserCFGTest, ComputeBlockOrder_Nest_If_In_SwitchCase) { 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 @@ TEST_F(SpvParserCFGTest, FindSwitchCaseHeaders_DefaultForTwoSwitches_AsCaseClaus 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_BypassesMerge_IsError) { } 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 @@ TEST_F(SpvParserCFGTest, ClassifyCFGEdges_LoopContinue_FromNestedLoopHeader_IsEr "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 @@ return; 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 @@ return; 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 @@ return; 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 @@ return; // 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 @@ return; // 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 @@ return; // 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 @@ return; // kLoopBreak: symmetry // kLoopContinue: dup general case // kIfBreak: TESTED -// kCaseFallThrough: TESTED // kForward: TESTED // // kIfBreak with: @@ -10543,25 +9858,14 @@ return; // 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 @@ return; 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 @@ return; 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 @@ return; 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 @@ return; 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 @@ TEST_F(SpvParserCFGTest, EmitBody_BranchConditional_IfBreak_IfBreak_DifferentIsE "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 e3d2acfb75..5b8310591a 100644 --- a/src/tint/reader/spirv/function_var_test.cc +++ b/src/tint/reader/spirv/function_var_test.cc @@ -1344,73 +1344,6 @@ return; 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 667a8f2f6e..0000000000 --- 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 de066b77ab..0000000000 --- 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 de066b77ab..0000000000 --- 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 6ca3be411a..0000000000 --- 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 d2cd679c2b..0000000000 --- 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 - -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 dc8c06387b..0000000000 --- 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 50254ab400..0000000000 --- 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 bf5da38858..0000000000 --- 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 e3f3d8ddaa..0000000000 --- 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 5ab7265d27..0000000000 --- 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 87bdbc830f..0000000000 --- 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 92feaa3487..0000000000 --- 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 a74bcc111e..0000000000 --- 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 d6e1cb3d3c..0000000000 --- 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 3f26450af2..0000000000 --- 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 0560ea2d94..0000000000 --- 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 7f1d386af8..0000000000 --- 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 5991963df8..0000000000 --- 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 6e9d1ebe3c..0000000000 --- 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 4269988a0e..0000000000 --- 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 83d008f39f..0000000000 --- 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 07f5e01b17..0000000000 --- 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 00cdb6388d..0000000000 --- a/test/tint/vk-gl-cts/graphicsfuzz/call-if-while-switch/0-opt.wgsl +++ /dev/null @@ -1,72 +0,0 @@ -struct buf0 { - injectionSwitch : vec2, -} - -@group(0) @binding(0) var x_6 : buf0; - -var x_GLF_color : vec4; - -fn main_1() { - var data : array; - 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(1.0, 1.0, 1.0, 1.0); - fallthrough; - } - case 19: { - x_54_phi = bitcast((x_40 + bitcast(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(1.0, 0.0, 0.0, 1.0); - return; -} - -struct main_out { - @location(0) - x_GLF_color_1 : vec4, -} - -@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 3ccb8a94cc..0000000000 --- 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 2e7ba7654a..0000000000 --- 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 6042980277..0000000000 --- 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, -} - -var gl_FragCoord : vec4; - -var array0 : array; - -var array1 : array; - -@group(0) @binding(0) var x_11 : buf0; - -var x_GLF_color : vec4; - -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(x_79, x_81, x_83, 1.0); - return; -} - -struct main_out { - @location(0) - x_GLF_color_1 : vec4, -} - -@fragment -fn main(@builtin(position) gl_FragCoord_param : vec4) -> 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 2e7ba7654a..0000000000 --- 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 07cab15f61..0000000000 --- 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 525654307b..0000000000 --- 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; - -struct buf0 { - x_GLF_uniform_int_values : Arr, -} - -@group(0) @binding(0) var x_6 : buf0; - -var x_GLF_color : vec4; - -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(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(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(x_85, x_85, x_85, x_85); - } - return; -} - -struct main_out { - @location(0) - x_GLF_color_1 : vec4, -} - -@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 97d37447ef..0000000000 --- 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 725431946c..0000000000 --- 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; - -struct buf0 { - x_GLF_uniform_int_values : Arr, -} - -var x_GLF_global_loop_count : i32; - -@group(0) @binding(0) var x_6 : buf0; - -var x_GLF_color : vec4; - -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(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(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(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, -} - -@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 5ef5478795..0000000000 --- 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 0cbcc97502..0000000000 --- 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 x_6 : buf0; - -var x_GLF_color : vec4; - -fn func_() -> vec4 { - let x_48 : i32 = x_6.one; - if ((x_48 == 1)) { - return vec4(1.0, 0.0, 0.0, 1.0); - } else { - return vec4(0.0, 0.0, 0.0, 0.0); - } -} - -fn main_1() { - var i : i32; - x_GLF_color = vec4(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 = 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, -} - -@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 55916643d8..0000000000 --- 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 e58d17ee86..0000000000 --- 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 x_6 : buf0; - -var x_GLF_color : vec4; - -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(1.0, 0.0, 0.0, 1.0); - } else { - x_GLF_color = vec4(0.0, 0.0, 0.0, 0.0); - } - return; -} - -struct main_out { - @location(0) - x_GLF_color_1 : vec4, -} - -@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 28e20bd1e5..0000000000 --- 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 d9cf1c740e..0000000000 --- a/test/tint/vk-gl-cts/graphicsfuzz/if-and-switch/0.wgsl +++ /dev/null @@ -1,47 +0,0 @@ -struct buf0 { - injectionSwitch : vec2, -} - -@group(0) @binding(0) var x_6 : buf0; - -var x_GLF_color : vec4; - -fn main_1() { - var data : array; - var x_32 : f32; - x_32 = x_6.injectionSwitch.x; - data[0] = x_32; - data[1] = x_32; - x_GLF_color = vec4(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(0.0, 0.0, 0.0, 0.0); - } - default: { - } - } - } - return; -} - -struct main_out { - @location(0) - x_GLF_color_1 : vec4, -} - -@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 7a18a0579e..0000000000 --- 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 f932ed929a..0000000000 --- a/test/tint/vk-gl-cts/graphicsfuzz/nested-switch-break-discard/0-opt.wgsl +++ /dev/null @@ -1,48 +0,0 @@ -struct buf0 { - injectionSwitch : vec2, -} - -var x_GLF_color : vec4; - -@group(0) @binding(0) var x_6 : buf0; - -var gl_FragCoord : vec4; - -fn main_1() { - x_GLF_color = vec4(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(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, -} - -@fragment -fn main(@builtin(position) gl_FragCoord_param : vec4) -> 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 099b620022..0000000000 --- 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 41e5cfeaad..0000000000 --- a/test/tint/vk-gl-cts/graphicsfuzz/switch-loop-switch-if/0-opt.wgsl +++ /dev/null @@ -1,73 +0,0 @@ -struct buf0 { - injectionSwitch : vec2, -} - -@group(0) @binding(0) var x_6 : buf0; - -var x_GLF_color : vec4; - -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(1.0, 0.0, 0.0, 1.0); - } else { - x_GLF_color = vec4(0.0, 0.0, 0.0, 1.0); - } - return; -} - -struct main_out { - @location(0) - x_GLF_color_1 : vec4, -} - -@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 20a942dccd..0000000000 --- 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 44c1c317a9..0000000000 --- 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, -} - -@group(0) @binding(0) var x_6 : buf0; - -var x_GLF_color : vec4; - -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((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, -} - -@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 5b160e165c..0000000000 --- a/test/tint/vk-gl-cts/graphicsfuzz/unreachable-loops-in-switch/0.wgsl +++ /dev/null @@ -1,63 +0,0 @@ -var x_GLF_color : vec4; - -fn main_1() { - var i : i32; - var data : array; - x_GLF_color = vec4(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, -}; - -@fragment -fn main() -> main_out { - main_1(); - return main_out(x_GLF_color); -}