encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

Remove fallthrough support from SPIRV-Reader. 

This CL removes support for fallthrough from the SPIRV-Reader.

Bug: tint:1644
Change-Id: I80b63d627960a82ba90de83af407c539b0442080
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/109004
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
This commit is contained in:
dan sinclair 2022-11-15 00:19:05 +00:00 committed by Dawn LUCI CQ
parent f0e6f57959
commit d8a986beae
48 changed files with 38 additions and 4124 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

72
src/tint/reader/spirv/function.cc
View File

@ -25,7 +25,6 @@
#include "src/tint/ast/call_statement.h" #include "src/tint/ast/call_statement.h"
#include "src/tint/ast/continue_statement.h" #include "src/tint/ast/continue_statement.h"
#include "src/tint/ast/discard_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/if_statement.h"
#include "src/tint/ast/loop_statement.h" #include "src/tint/ast/loop_statement.h"
#include "src/tint/ast/return_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: // For each branch encountered, classify each edge (S,T) as:
// - a back-edge // - a back-edge
// - a structured exit (specific ways of branching to enclosing construct) // - a structured exit (specific ways of branching to enclosing construct)
// - a normal (forward) edge, either natural control flow or a case // - a normal (forward) edge, either natural control flow or a case fallthrough
// fallthrough
// //
// If more than one block is targeted by a normal edge, then S must be a // If more than one block is targeted by a normal edge, then S must be a
// structured header. // structured header.
@ -2194,11 +2192,10 @@ bool FunctionEmitter::ClassifyCFGEdges() {
// There should only be one backedge per backedge block. // There should only be one backedge per backedge block.
uint32_t num_backedges = 0; uint32_t num_backedges = 0;
// Track destinations for normal forward edges, either kForward // Track destinations for normal forward edges, either kForward or kCaseFallThrough.
// or kCaseFallThrough. These count toward the need // These count toward the need to have a merge instruction. We also track kIfBreak edges
// 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
// because when used with normal forward edges, we'll need // variable.
// to generate a flow guard variable.
utils::Vector<uint32_t, 4> normal_forward_edges; utils::Vector<uint32_t, 4> normal_forward_edges;
utils::Vector<uint32_t, 4> if_break_edges; utils::Vector<uint32_t, 4> if_break_edges;
@ -2376,6 +2373,12 @@ bool FunctionEmitter::ClassifyCFGEdges() {
<< dest_construct.begin_id << " (dominance rule violated)"; << 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 forward edge
} // end successor } // end successor
@ -3261,11 +3264,9 @@ bool FunctionEmitter::EmitNormalTerminator(const BlockInfo& block_info) {
// The fallthrough case is special because WGSL requires the fallthrough // The fallthrough case is special because WGSL requires the fallthrough
// statement to be last in the case clause. // statement to be last in the case clause.
if (true_kind == EdgeKind::kCaseFallThrough) { if (true_kind == EdgeKind::kCaseFallThrough) {
return EmitConditionalCaseFallThrough(block_info, cond, false_kind, *false_info, return Fail() << "Fallthrough not supported in WGSL";
true);
} else if (false_kind == EdgeKind::kCaseFallThrough) { } else if (false_kind == EdgeKind::kCaseFallThrough) {
return EmitConditionalCaseFallThrough(block_info, cond, true_kind, *true_info, return Fail() << "Fallthrough not supported in WGSL";
false);
} }
// At this point, at most one edge is kForward or kIfBreak. // 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 ast::Statement* FunctionEmitter::MakeBranchDetailed(const BlockInfo& src_info,
const BlockInfo& dest_info, const BlockInfo& dest_info,
bool forced, 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; auto kind = src_info.succ_edge.find(dest_info.id)->second;
switch (kind) { switch (kind) {
case EdgeKind::kBack: case EdgeKind::kBack:
@ -3367,8 +3368,10 @@ const ast::Statement* FunctionEmitter::MakeBranchDetailed(const BlockInfo& src_i
// merge block is implicit. // merge block is implicit.
break; break;
} }
case EdgeKind::kCaseFallThrough: case EdgeKind::kCaseFallThrough: {
return create<ast::FallthroughStatement>(Source{}); Fail() << "Fallthrough not supported in WGSL";
return nullptr;
}
case EdgeKind::kForward: case EdgeKind::kForward:
// Unconditional forward branch is implicit. // Unconditional forward branch is implicit.
break; break;
@ -3398,45 +3401,6 @@ const ast::Statement* FunctionEmitter::MakeSimpleIf(const ast::Expression* condi
return if_stmt; return if_stmt;
} }
bool FunctionEmitter::EmitConditionalCaseFallThrough(const BlockInfo& src_info,
const ast::Expression* cond,
EdgeKind other_edge_kind,
const BlockInfo& other_dest,
bool fall_through_is_true_branch) {
// In WGSL, the fallthrough statement must come last in the case clause.
// So we'll emit an if statement for the other branch, and then emit
// the fallthrough.
// We have two distinct destinations. But we only get here if this
// is a normal terminator; in particular the source block is *not* the
// start of an if-selection. So at most one branch is a kForward or
// kCaseFallThrough.
if (other_edge_kind == EdgeKind::kForward) {
return Fail() << "internal error: normal terminator OpBranchConditional has "
"both forward and fallthrough edges";
}
if (other_edge_kind == EdgeKind::kIfBreak) {
return Fail() << "internal error: normal terminator OpBranchConditional has "
"both IfBreak and fallthrough edges. Violates nesting rule";
}
if (other_edge_kind == EdgeKind::kBack) {
return Fail() << "internal error: normal terminator OpBranchConditional has "
"both backedge and fallthrough edges. Violates nesting rule";
}
auto* other_branch = MakeForcedBranch(src_info, other_dest);
if (other_branch == nullptr) {
return Fail() << "internal error: expected a branch for edge-kind " << int(other_edge_kind);
}
if (fall_through_is_true_branch) {
AddStatement(MakeSimpleIf(cond, nullptr, other_branch));
} else {
AddStatement(MakeSimpleIf(cond, other_branch, nullptr));
}
AddStatement(create<ast::FallthroughStatement>(Source{}));
return success();
}
bool FunctionEmitter::EmitStatementsInBasicBlock(const BlockInfo& block_info, bool FunctionEmitter::EmitStatementsInBasicBlock(const BlockInfo& block_info,
bool* already_emitted) { bool* already_emitted) {
if (*already_emitted) { if (*already_emitted) {

29
src/tint/reader/spirv/function.h
View File

@ -34,15 +34,13 @@ namespace tint::reader::spirv {
// //
// The edge kinds are used in many ways. // The edge kinds are used in many ways.
// //
// For example, consider the edges leaving a basic block and going to distinct // For example, consider the edges leaving a basic block and going to distinct targets. If the
// targets. If the total number of kForward + kIfBreak + kCaseFallThrough edges // total number of kForward + kIfBreak + kCaseFallThrough edges is more than 1, then the block must
// is more than 1, then the block must be a structured header, i.e. it needs // be a structured header, i.e. it needs a merge instruction to declare the control flow divergence
// a merge instruction to declare the control flow divergence and associated // and associated reconvergence point. Those those edge kinds count toward divergence because
// reconvergence point. Those those edge kinds count toward divergence // SPIR-V is designed to easily map back to structured control flow in GLSL (and C). In GLSL and C,
// because SPIR-v is designed to easily map back to structured control flow // those forward-flow edges don't have a special statement to express them. The other forward
// in GLSL (and C). In GLSL and C, those forward-flow edges don't have a // edges: kSwitchBreak, kLoopBreak, and kLoopContinue directly map to 'break', 'break', and
// special statement to express them. The other forward edges: kSwitchBreak,
// kLoopBreak, and kLoopContinue directly map to 'break', 'break', and
// 'continue', respectively. // 'continue', respectively.
enum class EdgeKind { enum class EdgeKind {
// A back-edge: An edge from a node to one of its ancestors in a depth-first // 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 // This can only occur for an "if" selection, i.e. where the selection
// header ends in OpBranchConditional. // header ends in OpBranchConditional.
kIfBreak, 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, kCaseFallThrough,
// None of the above. // None of the above.
kForward kForward
@ -708,8 +707,7 @@ class FunctionEmitter {
/// Emits code for terminators, but that aren't part of entering or /// Emits code for terminators, but that aren't part of entering or
/// resolving structured control flow. That is, if the basic block /// resolving structured control flow. That is, if the basic block
/// terminator calls for it, emit the fallthrough, break, continue, return, /// terminator calls for it, emit the fallthrough break, continue, return, or kill commands.
/// or kill commands.
/// @param block_info the block with the terminator to emit (if any) /// @param block_info the block with the terminator to emit (if any)
/// @returns false if emission failed /// @returns false if emission failed
bool EmitNormalTerminator(const BlockInfo& block_info); bool EmitNormalTerminator(const BlockInfo& block_info);
@ -722,7 +720,7 @@ class FunctionEmitter {
/// @param src_info the source block /// @param src_info the source block
/// @param dest_info the destination block /// @param dest_info the destination block
/// @returns the new statement, or a null statement /// @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); return MakeBranchDetailed(src_info, dest_info, false, nullptr);
} }
@ -732,8 +730,7 @@ class FunctionEmitter {
/// @param src_info the source block /// @param src_info the source block
/// @param dest_info the destination block /// @param dest_info the destination block
/// @returns the new statement, or a null statement /// @returns the new statement, or a null statement
const ast::Statement* MakeForcedBranch(const BlockInfo& src_info, const ast::Statement* MakeForcedBranch(const BlockInfo& src_info, const BlockInfo& dest_info) {
const BlockInfo& dest_info) const {
return MakeBranchDetailed(src_info, dest_info, true, nullptr); return MakeBranchDetailed(src_info, dest_info, true, nullptr);
} }
@ -754,7 +751,7 @@ class FunctionEmitter {
const ast::Statement* MakeBranchDetailed(const BlockInfo& src_info, const ast::Statement* MakeBranchDetailed(const BlockInfo& src_info,
const BlockInfo& dest_info, const BlockInfo& dest_info,
bool forced, 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 /// 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 /// and the else-clause. Either or both clauses might be nullptr. If both

1095
src/tint/reader/spirv/function_cfg_test.cc
View File

File diff suppressed because it is too large Load Diff

67
src/tint/reader/spirv/function_var_test.cc
View File

@ -1344,73 +1344,6 @@ return;
EXPECT_EQ(expect, got) << got; 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) { TEST_F(SpvParserFunctionVarTest, EmitStatement_Phi_UseInPhiCountsAsUse) {
// From crbug.com/215 // From crbug.com/215
// If the only use of a combinatorially computed ID is as the value // If the only use of a combinatorially computed ID is as the value

12
test/tint/statements/switch/fallthrough.wgsl
View File

@ -1,12 +0,0 @@
@compute @workgroup_size(1)
fn f() {
var i : i32;
switch(i) {
case 0: {
fallthrough;
}
default: {
break;
}
}
}

20
test/tint/statements/switch/fallthrough.wgsl.expected.dxc.hlsl
View File

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

20
test/tint/statements/switch/fallthrough.wgsl.expected.fxc.hlsl
View File

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

23
test/tint/statements/switch/fallthrough.wgsl.expected.glsl
View File

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

20
test/tint/statements/switch/fallthrough.wgsl.expected.msl
View File

@ -1,20 +0,0 @@
statements/switch/fallthrough.wgsl:6:13 warning: use of deprecated language feature: fallthrough is set to be removed from WGSL. Case can accept multiple selectors if the existing case bodies are empty. (e.g. `case 1, 2, 3:`) `default` is a valid case selector value. (e.g. `case 1, default:`)
fallthrough;
^^^^^^^^^^^
#include <metal_stdlib>
using namespace metal;
kernel void f() {
int i = 0;
switch(i) {
case 0: {
/* fallthrough */
}
default: {
break;
}
}
return;
}

33
test/tint/statements/switch/fallthrough.wgsl.expected.spvasm
View File

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

16
test/tint/statements/switch/fallthrough.wgsl.expected.wgsl
View File

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

55
test/tint/unittest/reader/spirv/SpvParserCFGTest_ClassifyCFGEdges_Fallthrough_CaseTailToCase.spvasm
View File

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

55
test/tint/unittest/reader/spirv/SpvParserCFGTest_ClassifyCFGEdges_Fallthrough_CaseTailToDefaultNotMerge.spvasm
View File

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

55
test/tint/unittest/reader/spirv/SpvParserCFGTest_ClassifyCFGEdges_Fallthrough_DefaultToCase.spvasm
View File

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

67
test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_Nest_IfFallthrough_In_SwitchCase.spvasm
View File

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

57
test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough.spvasm
View File

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

55
test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough_FromCaseToDefaultToCase.spvasm
View File

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

57
test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough_FromDefault.spvasm
View File

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

61
test/tint/unittest/reader/spirv/SpvParserCFGTest_ComputeBlockOrder_RespectSwitchCaseFallthrough_Interleaved.spvasm
View File

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

57
test/tint/unittest/reader/spirv/SpvParserCFGTest_EmitBody_BranchConditional_Fallthrough_Fallthrough_Same.spvasm
View File

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

57
test/tint/unittest/reader/spirv/SpvParserCFGTest_EmitBody_BranchConditional_SwitchBreak_Fallthrough_OnFalse.spvasm
View File

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

57
test/tint/unittest/reader/spirv/SpvParserCFGTest_EmitBody_BranchConditional_SwitchBreak_Fallthrough_OnTrue.spvasm
View File

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

57
test/tint/unittest/reader/spirv/SpvParserCFGTest_EmitBody_Branch_Fallthrough.spvasm
View File

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

55
test/tint/unittest/reader/spirv/SpvParserCFGTest_TerminatorsAreValid_Switch.spvasm
View File

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

66
test/tint/unittest/reader/spirv/SpvParserFunctionVarTest_EmitStatement_Phi_InMerge_PredecessorsDominatdByNestedSwitchCase.spvasm
View File

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

93
test/tint/vk-gl-cts/graphicsfuzz/call-if-while-switch/0-opt.spvasm
View File

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

72
test/tint/vk-gl-cts/graphicsfuzz/call-if-while-switch/0-opt.wgsl
View File

@ -1,72 +0,0 @@
struct buf0 {
injectionSwitch : vec2<f32>,
}
@group(0) @binding(0) var<uniform> x_6 : buf0;
var<private> x_GLF_color : vec4<f32>;
fn main_1() {
var data : array<i32, 10u>;
var x_40 : i32;
var x_40_phi : i32;
var x_11_phi : i32;
let x_7 : i32 = data[1];
let x_10 : i32 = select(1, 2, (1 < x_7));
x_40_phi = 1;
x_11_phi = x_10;
loop {
var x_54 : i32;
var x_41 : i32;
var x_41_phi : i32;
x_40 = x_40_phi;
let x_11 : i32 = x_11_phi;
if ((x_11 < 3)) {
} else {
break;
}
var x_54_phi : i32;
let x_8 : i32 = (x_11 + 1);
let x_47 : f32 = x_6.injectionSwitch.x;
x_54_phi = x_40;
switch(i32(x_47)) {
case 78: {
x_GLF_color = vec4<f32>(1.0, 1.0, 1.0, 1.0);
fallthrough;
}
case 19: {
x_54_phi = bitcast<i32>((x_40 + bitcast<i32>(1)));
fallthrough;
}
case 23, 38: {
x_54 = x_54_phi;
x_41_phi = x_54;
continue;
}
default: {
x_41_phi = x_40;
continue;
}
}
continuing {
x_41 = x_41_phi;
x_40_phi = x_41;
x_11_phi = x_8;
}
}
data[x_40] = 1;
x_GLF_color = vec4<f32>(1.0, 0.0, 0.0, 1.0);
return;
}
struct main_out {
@location(0)
x_GLF_color_1 : vec4<f32>,
}
@fragment
fn main() -> main_out {
main_1();
return main_out(x_GLF_color);
}

154
test/tint/vk-gl-cts/graphicsfuzz/cov-dead-code-unreachable-merge/0-opt.spvasm
View File

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

93
test/tint/vk-gl-cts/graphicsfuzz/cov-dead-code-unreachable-merge/0-opt.spvasm.expected.dxc.hlsl
View File

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

84
test/tint/vk-gl-cts/graphicsfuzz/cov-dead-code-unreachable-merge/0-opt.wgsl
View File

@ -1,84 +0,0 @@
struct buf0 {
injectionSwitch : vec2<f32>,
}
var<private> gl_FragCoord : vec4<f32>;
var<private> array0 : array<f32, 3u>;
var<private> array1 : array<f32, 3u>;
@group(0) @binding(0) var<uniform> x_11 : buf0;
var<private> x_GLF_color : vec4<f32>;
fn main_1() {
var q : i32;
var i : i32;
var c : i32;
q = 0;
let x_55 : f32 = gl_FragCoord.x;
i = (i32(x_55) % 3);
c = 0;
loop {
let x_14 : i32 = c;
if ((x_14 < 3)) {
} else {
break;
}
let x_15 : i32 = c;
array0[x_15] = 0.0;
let x_16 : i32 = c;
array1[x_16] = 0.0;
let x_65 : f32 = x_11.injectionSwitch.x;
let x_18 : i32 = q;
switch((i32(x_65) + x_18)) {
case 51: {
loop {
if (true) {
} else {
break;
}
}
let x_20 : i32 = c;
array0[x_20] = 1.0;
fallthrough;
}
case 61: {
array1[0] = 1.0;
let x_21 : i32 = c;
array1[x_21] = 1.0;
}
case 0: {
q = 61;
}
default: {
}
}
continuing {
let x_22 : i32 = c;
c = (x_22 + 1);
}
}
let x_24 : i32 = i;
let x_79 : f32 = array1[x_24];
let x_25 : i32 = i;
let x_81 : f32 = array0[x_25];
let x_26 : i32 = i;
let x_83 : f32 = array0[x_26];
x_GLF_color = vec4<f32>(x_79, x_81, x_83, 1.0);
return;
}
struct main_out {
@location(0)
x_GLF_color_1 : vec4<f32>,
}
@fragment
fn main(@builtin(position) gl_FragCoord_param : vec4<f32>) -> main_out {
gl_FragCoord = gl_FragCoord_param;
main_1();
return main_out(x_GLF_color);
}

93
test/tint/vk-gl-cts/graphicsfuzz/cov-dead-code-unreachable-merge/0-opt.wgsl.expected.dxc.hlsl
View File

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

133
test/tint/vk-gl-cts/graphicsfuzz/cov-inst-combine-and-or-xor-switch/0-opt.spvasm
View File

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

85
test/tint/vk-gl-cts/graphicsfuzz/cov-inst-combine-and-or-xor-switch/0-opt.wgsl
View File

@ -1,85 +0,0 @@
struct strided_arr {
@size(16)
el : i32,
}
type Arr = array<strided_arr, 5u>;
struct buf0 {
x_GLF_uniform_int_values : Arr,
}
@group(0) @binding(0) var<uniform> x_6 : buf0;
var<private> x_GLF_color : vec4<f32>;
fn main_1() {
var count0 : i32;
var count1 : i32;
var i : i32;
let x_29 : i32 = x_6.x_GLF_uniform_int_values[2].el;
count0 = x_29;
let x_31 : i32 = x_6.x_GLF_uniform_int_values[2].el;
count1 = x_31;
let x_33 : i32 = x_6.x_GLF_uniform_int_values[2].el;
i = x_33;
loop {
let x_38 : i32 = i;
let x_40 : i32 = x_6.x_GLF_uniform_int_values[4].el;
if ((x_38 < x_40)) {
} else {
break;
}
let x_43 : i32 = i;
switch(x_43) {
case 0, 1: {
let x_47 : i32 = count0;
count0 = (x_47 + 1);
fallthrough;
}
case 2, 3: {
let x_49 : i32 = count1;
count1 = (x_49 + 1);
}
default: {
}
}
continuing {
let x_51 : i32 = i;
i = (x_51 + 1);
}
}
let x_53 : i32 = count1;
let x_55 : i32 = x_6.x_GLF_uniform_int_values[0].el;
if ((x_53 == x_55)) {
let x_61 : i32 = x_6.x_GLF_uniform_int_values[3].el;
let x_64 : i32 = x_6.x_GLF_uniform_int_values[2].el;
let x_67 : i32 = x_6.x_GLF_uniform_int_values[2].el;
let x_70 : i32 = x_6.x_GLF_uniform_int_values[3].el;
x_GLF_color = vec4<f32>(f32(x_61), f32(x_64), f32(x_67), f32(x_70));
} else {
let x_74 : i32 = x_6.x_GLF_uniform_int_values[2].el;
let x_75 : f32 = f32(x_74);
x_GLF_color = vec4<f32>(x_75, x_75, x_75, x_75);
}
let x_77 : i32 = count0;
let x_79 : i32 = x_6.x_GLF_uniform_int_values[1].el;
if ((x_77 != x_79)) {
let x_84 : i32 = x_6.x_GLF_uniform_int_values[2].el;
let x_85 : f32 = f32(x_84);
x_GLF_color = vec4<f32>(x_85, x_85, x_85, x_85);
}
return;
}
struct main_out {
@location(0)
x_GLF_color_1 : vec4<f32>,
}
@fragment
fn main() -> main_out {
main_1();
return main_out(x_GLF_color);
}

117
test/tint/vk-gl-cts/graphicsfuzz/cov-loop-returns-behind-true-and-false/0-opt.spvasm
View File

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

78
test/tint/vk-gl-cts/graphicsfuzz/cov-loop-returns-behind-true-and-false/0-opt.wgsl
View File

@ -1,78 +0,0 @@
struct strided_arr {
@size(16)
el : i32,
}
type Arr = array<strided_arr, 2u>;
struct buf0 {
x_GLF_uniform_int_values : Arr,
}
var<private> x_GLF_global_loop_count : i32;
@group(0) @binding(0) var<uniform> x_6 : buf0;
var<private> x_GLF_color : vec4<f32>;
fn main_1() {
x_GLF_global_loop_count = 0;
let x_27 : i32 = x_6.x_GLF_uniform_int_values[0].el;
switch(x_27) {
case 0: {
if (true) {
let x_34 : i32 = x_6.x_GLF_uniform_int_values[1].el;
let x_35 : f32 = f32(x_34);
x_GLF_color = vec4<f32>(x_35, x_35, x_35, x_35);
return;
}
fallthrough;
}
case 1: {
if (true) {
let x_40 : i32 = x_6.x_GLF_uniform_int_values[0].el;
let x_43 : i32 = x_6.x_GLF_uniform_int_values[1].el;
let x_46 : i32 = x_6.x_GLF_uniform_int_values[1].el;
let x_49 : i32 = x_6.x_GLF_uniform_int_values[0].el;
x_GLF_color = vec4<f32>(f32(x_40), f32(x_43), f32(x_46), f32(x_49));
if (false) {
let x_55 : i32 = x_6.x_GLF_uniform_int_values[0].el;
let x_56 : f32 = f32(x_55);
x_GLF_color = vec4<f32>(x_56, x_56, x_56, x_56);
loop {
let x_62 : i32 = x_GLF_global_loop_count;
x_GLF_global_loop_count = (x_62 + 1);
if (false) {
return;
}
if (true) {
return;
}
continuing {
let x_68 : i32 = x_GLF_global_loop_count;
if ((x_68 < 100)) {
} else {
break;
}
}
}
}
}
}
default: {
}
}
return;
}
struct main_out {
@location(0)
x_GLF_color_1 : vec4<f32>,
}
@fragment
fn main() -> main_out {
main_1();
return main_out(x_GLF_color);
}

92
test/tint/vk-gl-cts/graphicsfuzz/cov-ssa-rewrite-case-with-default/0-opt.spvasm
View File

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

61
test/tint/vk-gl-cts/graphicsfuzz/cov-ssa-rewrite-case-with-default/0-opt.wgsl
View File

@ -1,61 +0,0 @@
struct buf0 {
one : i32,
}
@group(0) @binding(0) var<uniform> x_6 : buf0;
var<private> x_GLF_color : vec4<f32>;
fn func_() -> vec4<f32> {
let x_48 : i32 = x_6.one;
if ((x_48 == 1)) {
return vec4<f32>(1.0, 0.0, 0.0, 1.0);
} else {
return vec4<f32>(0.0, 0.0, 0.0, 0.0);
}
}
fn main_1() {
var i : i32;
x_GLF_color = vec4<f32>(0.0, 0.0, 0.0, 0.0);
i = 0;
loop {
let x_33 : i32 = i;
let x_35 : i32 = x_6.one;
if ((x_33 <= x_35)) {
} else {
break;
}
let x_38 : i32 = i;
switch(x_38) {
case 1: {
let x_43 : vec4<f32> = func_();
x_GLF_color = x_43;
fallthrough;
}
default: {
fallthrough;
}
case 0: {
x_GLF_color.y = 0.0;
}
}
continuing {
let x_44 : i32 = i;
i = (x_44 + 1);
}
}
return;
}
struct main_out {
@location(0)
x_GLF_color_1 : vec4<f32>,
}
@fragment
fn main() -> main_out {
main_1();
return main_out(x_GLF_color);
}

67
test/tint/vk-gl-cts/graphicsfuzz/cov-val-cfg-case-fallthrough/0-opt.spvasm
View File

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

42
test/tint/vk-gl-cts/graphicsfuzz/cov-val-cfg-case-fallthrough/0-opt.wgsl
View File

@ -1,42 +0,0 @@
struct buf0 {
one : i32,
}
@group(0) @binding(0) var<uniform> x_6 : buf0;
var<private> x_GLF_color : vec4<f32>;
fn main_1() {
var a : i32;
a = 0;
let x_26 : i32 = x_6.one;
switch(x_26) {
case 2, 3: {
a = 1;
fallthrough;
}
case 4: {
}
default: {
a = 2;
}
}
let x_31 : i32 = a;
if ((x_31 == 2)) {
x_GLF_color = vec4<f32>(1.0, 0.0, 0.0, 1.0);
} else {
x_GLF_color = vec4<f32>(0.0, 0.0, 0.0, 0.0);
}
return;
}
struct main_out {
@location(0)
x_GLF_color_1 : vec4<f32>,
}
@fragment
fn main() -> main_out {
main_1();
return main_out(x_GLF_color);
}

74
test/tint/vk-gl-cts/graphicsfuzz/if-and-switch/0.spvasm
View File

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

47
test/tint/vk-gl-cts/graphicsfuzz/if-and-switch/0.wgsl
View File

@ -1,47 +0,0 @@
struct buf0 {
injectionSwitch : vec2<f32>,
}
@group(0) @binding(0) var<uniform> x_6 : buf0;
var<private> x_GLF_color : vec4<f32>;
fn main_1() {
var data : array<f32, 2u>;
var x_32 : f32;
x_32 = x_6.injectionSwitch.x;
data[0] = x_32;
data[1] = x_32;
x_GLF_color = vec4<f32>(1.0, 0.0, 0.0, 1.0);
let x_35 : f32 = data[1];
if ((x_35 > 1.0)) {
var x_43_phi : f32;
let x_39 : i32 = i32(x_32);
x_43_phi = 0.0;
switch(x_39) {
case 0: {
x_43_phi = 1.0;
fallthrough;
}
case 1: {
let x_43 : f32 = x_43_phi;
data[x_39] = x_43;
x_GLF_color = vec4<f32>(0.0, 0.0, 0.0, 0.0);
}
default: {
}
}
}
return;
}
struct main_out {
@location(0)
x_GLF_color_1 : vec4<f32>,
}
@fragment
fn main() -> main_out {
main_1();
return main_out(x_GLF_color);
}

72
test/tint/vk-gl-cts/graphicsfuzz/nested-switch-break-discard/0-opt.spvasm
View File

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

48
test/tint/vk-gl-cts/graphicsfuzz/nested-switch-break-discard/0-opt.wgsl
View File

@ -1,48 +0,0 @@
struct buf0 {
injectionSwitch : vec2<f32>,
}
var<private> x_GLF_color : vec4<f32>;
@group(0) @binding(0) var<uniform> x_6 : buf0;
var<private> gl_FragCoord : vec4<f32>;
fn main_1() {
x_GLF_color = vec4<f32>(0.0, 0.0, 0.0, 1.0);
let x_30 : f32 = x_6.injectionSwitch.x;
switch(i32(x_30)) {
case 0: {
switch(1) {
case 1: {
let x_38 : f32 = gl_FragCoord.y;
if ((x_38 >= 0.0)) {
x_GLF_color = vec4<f32>(1.0, 0.0, 0.0, 1.0);
break;
}
discard;
}
default: {
}
}
fallthrough;
}
case 42: {
}
default: {
}
}
return;
}
struct main_out {
@location(0)
x_GLF_color_1 : vec4<f32>,
}
@fragment
fn main(@builtin(position) gl_FragCoord_param : vec4<f32>) -> main_out {
gl_FragCoord = gl_FragCoord_param;
main_1();
return main_out(x_GLF_color);
}

134
test/tint/vk-gl-cts/graphicsfuzz/switch-loop-switch-if/0-opt.spvasm
View File

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

73
test/tint/vk-gl-cts/graphicsfuzz/switch-loop-switch-if/0-opt.wgsl
View File

@ -1,73 +0,0 @@
struct buf0 {
injectionSwitch : vec2<f32>,
}
@group(0) @binding(0) var<uniform> x_6 : buf0;
var<private> x_GLF_color : vec4<f32>;
fn main_1() {
var i : i32;
let x_51 : f32 = x_6.injectionSwitch.x;
i = i32(x_51);
let x_8 : i32 = i;
switch(x_8) {
case 0: {
loop {
let x_9 : i32 = i;
i = (x_9 + 1);
let x_11 : i32 = i;
switch(x_11) {
case 2: {
let x_12 : i32 = i;
i = (x_12 + 5);
}
case 1: {
continue;
}
default: {
let x_14 : i32 = i;
i = (x_14 + 7);
}
}
continuing {
let x_16 : i32 = i;
if ((x_16 > 200)) {
} else {
break;
}
}
}
let x_17 : i32 = i;
if ((x_17 > 100)) {
let x_18 : i32 = i;
i = (x_18 - 2);
break;
}
fallthrough;
}
default: {
let x_20 : i32 = i;
i = (x_20 - 3);
}
}
let x_22 : i32 = i;
if ((x_22 == -2)) {
x_GLF_color = vec4<f32>(1.0, 0.0, 0.0, 1.0);
} else {
x_GLF_color = vec4<f32>(0.0, 0.0, 0.0, 1.0);
}
return;
}
struct main_out {
@location(0)
x_GLF_color_1 : vec4<f32>,
}
@fragment
fn main() -> main_out {
main_1();
return main_out(x_GLF_color);
}

96
test/tint/vk-gl-cts/graphicsfuzz/switch-with-fall-through-cases/0-opt.spvasm
View File

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

73
test/tint/vk-gl-cts/graphicsfuzz/switch-with-fall-through-cases/0-opt.wgsl
View File

@ -1,73 +0,0 @@
struct buf0 {
injectionSwitch : vec2<f32>,
}
@group(0) @binding(0) var<uniform> x_6 : buf0;
var<private> x_GLF_color : vec4<f32>;
fn main_1() {
var i : i32;
var value : i32;
var y : f32;
var x_31_phi : i32;
i = 0;
x_31_phi = 0;
loop {
let x_31 : i32 = x_31_phi;
let x_37 : f32 = x_6.injectionSwitch.x;
if ((x_31 < (2 + i32(x_37)))) {
} else {
break;
}
var x_55_phi : f32;
var x_46_phi : f32;
value = x_31;
y = 0.5;
x_55_phi = 0.5;
x_46_phi = 0.5;
switch(x_31) {
case 0: {
let x_54 : f32 = (0.5 + 0.5);
y = x_54;
x_55_phi = x_54;
fallthrough;
}
case 1: {
let x_55 : f32 = x_55_phi;
let x_47 : f32 = clamp(1.0, 0.5, x_55);
y = x_47;
x_46_phi = x_47;
fallthrough;
}
default: {
fallthrough;
}
case 2: {
let x_46 : f32 = x_46_phi;
if ((x_46 == 1.0)) {
x_GLF_color = vec4<f32>(f32((x_31 + 1)), 0.0, 0.0, 1.0);
return;
}
}
}
continuing {
let x_32 : i32 = (x_31 + 1);
i = x_32;
x_31_phi = x_32;
}
}
return;
}
struct main_out {
@location(0)
x_GLF_color_1 : vec4<f32>,
}
@fragment
fn main() -> main_out {
main_1();
return main_out(x_GLF_color);
}

63
test/tint/vk-gl-cts/graphicsfuzz/unreachable-loops-in-switch/0.wgsl
View File

@ -1,63 +0,0 @@
var<private> x_GLF_color : vec4<f32>;
fn main_1() {
var i : i32;
var data : array<f32, 1>;
x_GLF_color = vec4<f32>(1.0, 0.0, 0.0, 1.0);
i = 0;
loop {
let x_6 : i32 = i;
if ((x_6 < 1)) {
} else {
break;
}
let x_7 : i32 = i;
let x_40 : f32 = data[x_7];
let x_42 : f32 = data[0];
if ((x_40 < x_42)) {
if (false) {
let x_8 : i32 = i;
if ((f32(x_8) >= 1.0)) {
}
}
switch(0) {
case 1: {
loop {
if (true) {
} else {
break;
}
}
loop {
break;
}
fallthrough;
}
case 0: {
data[0] = 2.0;
}
default: {
}
}
break;
}
continuing {
let x_9 : i32 = i;
i = (x_9 + 1);
}
}
return;
}
struct main_out {
@location(0)
x_GLF_color_1 : vec4<f32>,
};
@fragment
fn main() -> main_out {
main_1();
return main_out(x_GLF_color);
}