tint/resolver: Limit scope depth / if-chains.

DXC will bail if these get too deeply nested (~256).
This is also a risk for stack-overflows, so apply a limit agreed by the
WGSL working group.

Fixed: tint:1518
Change-Id: Idacdba85b36b27a0a89a3a7958fd4c6cce7dc84d
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/105964
Auto-Submit: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: David Neto <dneto@google.com>
This commit is contained in:
Ben Clayton 2022-10-25 16:53:06 +00:00 committed by Dawn LUCI CQ
parent 3a6b5362d5
commit bd5bd247f0
10 changed files with 2723 additions and 1 deletions

View File

@ -93,6 +93,7 @@ namespace tint::resolver {
namespace { namespace {
constexpr int64_t kMaxArrayElementCount = 65536; constexpr int64_t kMaxArrayElementCount = 65536;
constexpr uint32_t kMaxStatementDepth = 127;
} // namespace } // namespace
@ -3400,6 +3401,14 @@ SEM* Resolver::StatementScope(const ast::Statement* ast, SEM* sem, F&& callback)
TINT_SCOPED_ASSIGNMENT(current_statement_, sem); TINT_SCOPED_ASSIGNMENT(current_statement_, sem);
TINT_SCOPED_ASSIGNMENT(current_compound_statement_, TINT_SCOPED_ASSIGNMENT(current_compound_statement_,
as_compound ? as_compound : current_compound_statement_); as_compound ? as_compound : current_compound_statement_);
TINT_SCOPED_ASSIGNMENT(current_scoping_depth_, current_scoping_depth_ + 1);
if (current_scoping_depth_ > kMaxStatementDepth) {
AddError("statement nesting depth / chaining length exceeds limit of " +
std::to_string(kMaxStatementDepth),
ast->source);
return nullptr;
}
if (!callback()) { if (!callback()) {
return nullptr; return nullptr;

View File

@ -377,7 +377,8 @@ class Resolver {
/// * Assigns `sem` to #current_compound_statement_ if `sem` derives from /// * Assigns `sem` to #current_compound_statement_ if `sem` derives from
/// sem::CompoundStatement. /// sem::CompoundStatement.
/// * Then calls `callback`. /// * Then calls `callback`.
/// * Before returning #current_statement_ and #current_compound_statement_ are restored to their original values. /// * Before returning #current_statement_ and #current_compound_statement_ are restored to
/// their original values.
/// @returns `sem` if `callback` returns true, otherwise `nullptr`. /// @returns `sem` if `callback` returns true, otherwise `nullptr`.
template <typename SEM, typename F> template <typename SEM, typename F>
SEM* StatementScope(const ast::Statement* ast, SEM* sem, F&& callback); SEM* StatementScope(const ast::Statement* ast, SEM* sem, F&& callback);
@ -439,6 +440,7 @@ class Resolver {
sem::Function* current_function_ = nullptr; sem::Function* current_function_ = nullptr;
sem::Statement* current_statement_ = nullptr; sem::Statement* current_statement_ = nullptr;
sem::CompoundStatement* current_compound_statement_ = nullptr; sem::CompoundStatement* current_compound_statement_ = nullptr;
uint32_t current_scoping_depth_ = 0;
}; };
} // namespace tint::resolver } // namespace tint::resolver

View File

@ -2331,5 +2331,47 @@ TEST_F(ResolverTest, Literal_F16WithExtension) {
EXPECT_TRUE(r()->Resolve()); EXPECT_TRUE(r()->Resolve());
} }
// Windows debug builds have significantly smaller stack than other builds, and these tests will stack
// overflow.
#if !defined(NDEBUG)
TEST_F(ResolverTest, ScopeDepth_NestedBlocks) {
const ast::Statement* stmt = Return();
for (size_t i = 0; i < 150; i++) {
stmt = Block(Source{{i, 1}}, stmt);
}
WrapInFunction(stmt);
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
"23:1 error: statement nesting depth / chaining length exceeds limit of 127");
}
TEST_F(ResolverTest, ScopeDepth_NestedIf) {
const ast::Statement* stmt = Return();
for (size_t i = 0; i < 150; i++) {
stmt = If(Source{{i, 1}}, false, Block(Source{{i, 2}}, stmt));
}
WrapInFunction(stmt);
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
"86:1 error: statement nesting depth / chaining length exceeds limit of 127");
}
TEST_F(ResolverTest, ScopeDepth_IfElseChain) {
const ast::Statement* stmt = nullptr;
for (size_t i = 0; i < 150; i++) {
stmt = If(Source{{i, 1}}, false, Block(Source{{i, 2}}), Else(stmt));
}
WrapInFunction(stmt);
EXPECT_FALSE(r()->Resolve());
EXPECT_EQ(r()->error(),
"24:2 error: statement nesting depth / chaining length exceeds limit of 127");
}
#endif // !defined(NDEBUG)
} // namespace } // namespace
} // namespace tint::resolver } // namespace tint::resolver

View File

@ -0,0 +1,128 @@
@compute @workgroup_size(1,1,1)
fn main() {
if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
else if (false) {}
}

View File

@ -0,0 +1,378 @@
[numthreads(1, 1, 1)]
void main() {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
return;
}

View File

@ -0,0 +1,378 @@
[numthreads(1, 1, 1)]
void main() {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
return;
}

View File

@ -0,0 +1,384 @@
#version 310 es
void tint_symbol() {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
void main() {
tint_symbol();
return;
}

View File

@ -0,0 +1,381 @@
#include <metal_stdlib>
using namespace metal;
kernel void tint_symbol() {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
} else {
if (false) {
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
}
return;
}

View File

@ -0,0 +1,891 @@
; SPIR-V
; Version: 1.3
; Generator: Google Tint Compiler; 0
; Bound: 381
; Schema: 0
OpCapability Shader
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %main "main"
OpExecutionMode %main LocalSize 1 1 1
OpName %main "main"
%void = OpTypeVoid
%1 = OpTypeFunction %void
%bool = OpTypeBool
%6 = OpConstantNull %bool
%main = OpFunction %void None %1
%4 = OpLabel
OpSelectionMerge %7 None
OpBranchConditional %6 %8 %9
%8 = OpLabel
OpBranch %7
%9 = OpLabel
OpSelectionMerge %10 None
OpBranchConditional %6 %11 %12
%11 = OpLabel
OpBranch %10
%12 = OpLabel
OpSelectionMerge %13 None
OpBranchConditional %6 %14 %15
%14 = OpLabel
OpBranch %13
%15 = OpLabel
OpSelectionMerge %16 None
OpBranchConditional %6 %17 %18
%17 = OpLabel
OpBranch %16
%18 = OpLabel
OpSelectionMerge %19 None
OpBranchConditional %6 %20 %21
%20 = OpLabel
OpBranch %19
%21 = OpLabel
OpSelectionMerge %22 None
OpBranchConditional %6 %23 %24
%23 = OpLabel
OpBranch %22
%24 = OpLabel
OpSelectionMerge %25 None
OpBranchConditional %6 %26 %27
%26 = OpLabel
OpBranch %25
%27 = OpLabel
OpSelectionMerge %28 None
OpBranchConditional %6 %29 %30
%29 = OpLabel
OpBranch %28
%30 = OpLabel
OpSelectionMerge %31 None
OpBranchConditional %6 %32 %33
%32 = OpLabel
OpBranch %31
%33 = OpLabel
OpSelectionMerge %34 None
OpBranchConditional %6 %35 %36
%35 = OpLabel
OpBranch %34
%36 = OpLabel
OpSelectionMerge %37 None
OpBranchConditional %6 %38 %39
%38 = OpLabel
OpBranch %37
%39 = OpLabel
OpSelectionMerge %40 None
OpBranchConditional %6 %41 %42
%41 = OpLabel
OpBranch %40
%42 = OpLabel
OpSelectionMerge %43 None
OpBranchConditional %6 %44 %45
%44 = OpLabel
OpBranch %43
%45 = OpLabel
OpSelectionMerge %46 None
OpBranchConditional %6 %47 %48
%47 = OpLabel
OpBranch %46
%48 = OpLabel
OpSelectionMerge %49 None
OpBranchConditional %6 %50 %51
%50 = OpLabel
OpBranch %49
%51 = OpLabel
OpSelectionMerge %52 None
OpBranchConditional %6 %53 %54
%53 = OpLabel
OpBranch %52
%54 = OpLabel
OpSelectionMerge %55 None
OpBranchConditional %6 %56 %57
%56 = OpLabel
OpBranch %55
%57 = OpLabel
OpSelectionMerge %58 None
OpBranchConditional %6 %59 %60
%59 = OpLabel
OpBranch %58
%60 = OpLabel
OpSelectionMerge %61 None
OpBranchConditional %6 %62 %63
%62 = OpLabel
OpBranch %61
%63 = OpLabel
OpSelectionMerge %64 None
OpBranchConditional %6 %65 %66
%65 = OpLabel
OpBranch %64
%66 = OpLabel
OpSelectionMerge %67 None
OpBranchConditional %6 %68 %69
%68 = OpLabel
OpBranch %67
%69 = OpLabel
OpSelectionMerge %70 None
OpBranchConditional %6 %71 %72
%71 = OpLabel
OpBranch %70
%72 = OpLabel
OpSelectionMerge %73 None
OpBranchConditional %6 %74 %75
%74 = OpLabel
OpBranch %73
%75 = OpLabel
OpSelectionMerge %76 None
OpBranchConditional %6 %77 %78
%77 = OpLabel
OpBranch %76
%78 = OpLabel
OpSelectionMerge %79 None
OpBranchConditional %6 %80 %81
%80 = OpLabel
OpBranch %79
%81 = OpLabel
OpSelectionMerge %82 None
OpBranchConditional %6 %83 %84
%83 = OpLabel
OpBranch %82
%84 = OpLabel
OpSelectionMerge %85 None
OpBranchConditional %6 %86 %87
%86 = OpLabel
OpBranch %85
%87 = OpLabel
OpSelectionMerge %88 None
OpBranchConditional %6 %89 %90
%89 = OpLabel
OpBranch %88
%90 = OpLabel
OpSelectionMerge %91 None
OpBranchConditional %6 %92 %93
%92 = OpLabel
OpBranch %91
%93 = OpLabel
OpSelectionMerge %94 None
OpBranchConditional %6 %95 %96
%95 = OpLabel
OpBranch %94
%96 = OpLabel
OpSelectionMerge %97 None
OpBranchConditional %6 %98 %99
%98 = OpLabel
OpBranch %97
%99 = OpLabel
OpSelectionMerge %100 None
OpBranchConditional %6 %101 %102
%101 = OpLabel
OpBranch %100
%102 = OpLabel
OpSelectionMerge %103 None
OpBranchConditional %6 %104 %105
%104 = OpLabel
OpBranch %103
%105 = OpLabel
OpSelectionMerge %106 None
OpBranchConditional %6 %107 %108
%107 = OpLabel
OpBranch %106
%108 = OpLabel
OpSelectionMerge %109 None
OpBranchConditional %6 %110 %111
%110 = OpLabel
OpBranch %109
%111 = OpLabel
OpSelectionMerge %112 None
OpBranchConditional %6 %113 %114
%113 = OpLabel
OpBranch %112
%114 = OpLabel
OpSelectionMerge %115 None
OpBranchConditional %6 %116 %117
%116 = OpLabel
OpBranch %115
%117 = OpLabel
OpSelectionMerge %118 None
OpBranchConditional %6 %119 %120
%119 = OpLabel
OpBranch %118
%120 = OpLabel
OpSelectionMerge %121 None
OpBranchConditional %6 %122 %123
%122 = OpLabel
OpBranch %121
%123 = OpLabel
OpSelectionMerge %124 None
OpBranchConditional %6 %125 %126
%125 = OpLabel
OpBranch %124
%126 = OpLabel
OpSelectionMerge %127 None
OpBranchConditional %6 %128 %129
%128 = OpLabel
OpBranch %127
%129 = OpLabel
OpSelectionMerge %130 None
OpBranchConditional %6 %131 %132
%131 = OpLabel
OpBranch %130
%132 = OpLabel
OpSelectionMerge %133 None
OpBranchConditional %6 %134 %135
%134 = OpLabel
OpBranch %133
%135 = OpLabel
OpSelectionMerge %136 None
OpBranchConditional %6 %137 %138
%137 = OpLabel
OpBranch %136
%138 = OpLabel
OpSelectionMerge %139 None
OpBranchConditional %6 %140 %141
%140 = OpLabel
OpBranch %139
%141 = OpLabel
OpSelectionMerge %142 None
OpBranchConditional %6 %143 %144
%143 = OpLabel
OpBranch %142
%144 = OpLabel
OpSelectionMerge %145 None
OpBranchConditional %6 %146 %147
%146 = OpLabel
OpBranch %145
%147 = OpLabel
OpSelectionMerge %148 None
OpBranchConditional %6 %149 %150
%149 = OpLabel
OpBranch %148
%150 = OpLabel
OpSelectionMerge %151 None
OpBranchConditional %6 %152 %153
%152 = OpLabel
OpBranch %151
%153 = OpLabel
OpSelectionMerge %154 None
OpBranchConditional %6 %155 %156
%155 = OpLabel
OpBranch %154
%156 = OpLabel
OpSelectionMerge %157 None
OpBranchConditional %6 %158 %159
%158 = OpLabel
OpBranch %157
%159 = OpLabel
OpSelectionMerge %160 None
OpBranchConditional %6 %161 %162
%161 = OpLabel
OpBranch %160
%162 = OpLabel
OpSelectionMerge %163 None
OpBranchConditional %6 %164 %165
%164 = OpLabel
OpBranch %163
%165 = OpLabel
OpSelectionMerge %166 None
OpBranchConditional %6 %167 %168
%167 = OpLabel
OpBranch %166
%168 = OpLabel
OpSelectionMerge %169 None
OpBranchConditional %6 %170 %171
%170 = OpLabel
OpBranch %169
%171 = OpLabel
OpSelectionMerge %172 None
OpBranchConditional %6 %173 %174
%173 = OpLabel
OpBranch %172
%174 = OpLabel
OpSelectionMerge %175 None
OpBranchConditional %6 %176 %177
%176 = OpLabel
OpBranch %175
%177 = OpLabel
OpSelectionMerge %178 None
OpBranchConditional %6 %179 %180
%179 = OpLabel
OpBranch %178
%180 = OpLabel
OpSelectionMerge %181 None
OpBranchConditional %6 %182 %183
%182 = OpLabel
OpBranch %181
%183 = OpLabel
OpSelectionMerge %184 None
OpBranchConditional %6 %185 %186
%185 = OpLabel
OpBranch %184
%186 = OpLabel
OpSelectionMerge %187 None
OpBranchConditional %6 %188 %189
%188 = OpLabel
OpBranch %187
%189 = OpLabel
OpSelectionMerge %190 None
OpBranchConditional %6 %191 %192
%191 = OpLabel
OpBranch %190
%192 = OpLabel
OpSelectionMerge %193 None
OpBranchConditional %6 %194 %195
%194 = OpLabel
OpBranch %193
%195 = OpLabel
OpSelectionMerge %196 None
OpBranchConditional %6 %197 %198
%197 = OpLabel
OpBranch %196
%198 = OpLabel
OpSelectionMerge %199 None
OpBranchConditional %6 %200 %201
%200 = OpLabel
OpBranch %199
%201 = OpLabel
OpSelectionMerge %202 None
OpBranchConditional %6 %203 %204
%203 = OpLabel
OpBranch %202
%204 = OpLabel
OpSelectionMerge %205 None
OpBranchConditional %6 %206 %207
%206 = OpLabel
OpBranch %205
%207 = OpLabel
OpSelectionMerge %208 None
OpBranchConditional %6 %209 %210
%209 = OpLabel
OpBranch %208
%210 = OpLabel
OpSelectionMerge %211 None
OpBranchConditional %6 %212 %213
%212 = OpLabel
OpBranch %211
%213 = OpLabel
OpSelectionMerge %214 None
OpBranchConditional %6 %215 %216
%215 = OpLabel
OpBranch %214
%216 = OpLabel
OpSelectionMerge %217 None
OpBranchConditional %6 %218 %219
%218 = OpLabel
OpBranch %217
%219 = OpLabel
OpSelectionMerge %220 None
OpBranchConditional %6 %221 %222
%221 = OpLabel
OpBranch %220
%222 = OpLabel
OpSelectionMerge %223 None
OpBranchConditional %6 %224 %225
%224 = OpLabel
OpBranch %223
%225 = OpLabel
OpSelectionMerge %226 None
OpBranchConditional %6 %227 %228
%227 = OpLabel
OpBranch %226
%228 = OpLabel
OpSelectionMerge %229 None
OpBranchConditional %6 %230 %231
%230 = OpLabel
OpBranch %229
%231 = OpLabel
OpSelectionMerge %232 None
OpBranchConditional %6 %233 %234
%233 = OpLabel
OpBranch %232
%234 = OpLabel
OpSelectionMerge %235 None
OpBranchConditional %6 %236 %237
%236 = OpLabel
OpBranch %235
%237 = OpLabel
OpSelectionMerge %238 None
OpBranchConditional %6 %239 %240
%239 = OpLabel
OpBranch %238
%240 = OpLabel
OpSelectionMerge %241 None
OpBranchConditional %6 %242 %243
%242 = OpLabel
OpBranch %241
%243 = OpLabel
OpSelectionMerge %244 None
OpBranchConditional %6 %245 %246
%245 = OpLabel
OpBranch %244
%246 = OpLabel
OpSelectionMerge %247 None
OpBranchConditional %6 %248 %249
%248 = OpLabel
OpBranch %247
%249 = OpLabel
OpSelectionMerge %250 None
OpBranchConditional %6 %251 %252
%251 = OpLabel
OpBranch %250
%252 = OpLabel
OpSelectionMerge %253 None
OpBranchConditional %6 %254 %255
%254 = OpLabel
OpBranch %253
%255 = OpLabel
OpSelectionMerge %256 None
OpBranchConditional %6 %257 %258
%257 = OpLabel
OpBranch %256
%258 = OpLabel
OpSelectionMerge %259 None
OpBranchConditional %6 %260 %261
%260 = OpLabel
OpBranch %259
%261 = OpLabel
OpSelectionMerge %262 None
OpBranchConditional %6 %263 %264
%263 = OpLabel
OpBranch %262
%264 = OpLabel
OpSelectionMerge %265 None
OpBranchConditional %6 %266 %267
%266 = OpLabel
OpBranch %265
%267 = OpLabel
OpSelectionMerge %268 None
OpBranchConditional %6 %269 %270
%269 = OpLabel
OpBranch %268
%270 = OpLabel
OpSelectionMerge %271 None
OpBranchConditional %6 %272 %273
%272 = OpLabel
OpBranch %271
%273 = OpLabel
OpSelectionMerge %274 None
OpBranchConditional %6 %275 %276
%275 = OpLabel
OpBranch %274
%276 = OpLabel
OpSelectionMerge %277 None
OpBranchConditional %6 %278 %279
%278 = OpLabel
OpBranch %277
%279 = OpLabel
OpSelectionMerge %280 None
OpBranchConditional %6 %281 %282
%281 = OpLabel
OpBranch %280
%282 = OpLabel
OpSelectionMerge %283 None
OpBranchConditional %6 %284 %285
%284 = OpLabel
OpBranch %283
%285 = OpLabel
OpSelectionMerge %286 None
OpBranchConditional %6 %287 %288
%287 = OpLabel
OpBranch %286
%288 = OpLabel
OpSelectionMerge %289 None
OpBranchConditional %6 %290 %291
%290 = OpLabel
OpBranch %289
%291 = OpLabel
OpSelectionMerge %292 None
OpBranchConditional %6 %293 %294
%293 = OpLabel
OpBranch %292
%294 = OpLabel
OpSelectionMerge %295 None
OpBranchConditional %6 %296 %297
%296 = OpLabel
OpBranch %295
%297 = OpLabel
OpSelectionMerge %298 None
OpBranchConditional %6 %299 %300
%299 = OpLabel
OpBranch %298
%300 = OpLabel
OpSelectionMerge %301 None
OpBranchConditional %6 %302 %303
%302 = OpLabel
OpBranch %301
%303 = OpLabel
OpSelectionMerge %304 None
OpBranchConditional %6 %305 %306
%305 = OpLabel
OpBranch %304
%306 = OpLabel
OpSelectionMerge %307 None
OpBranchConditional %6 %308 %309
%308 = OpLabel
OpBranch %307
%309 = OpLabel
OpSelectionMerge %310 None
OpBranchConditional %6 %311 %312
%311 = OpLabel
OpBranch %310
%312 = OpLabel
OpSelectionMerge %313 None
OpBranchConditional %6 %314 %315
%314 = OpLabel
OpBranch %313
%315 = OpLabel
OpSelectionMerge %316 None
OpBranchConditional %6 %317 %318
%317 = OpLabel
OpBranch %316
%318 = OpLabel
OpSelectionMerge %319 None
OpBranchConditional %6 %320 %321
%320 = OpLabel
OpBranch %319
%321 = OpLabel
OpSelectionMerge %322 None
OpBranchConditional %6 %323 %324
%323 = OpLabel
OpBranch %322
%324 = OpLabel
OpSelectionMerge %325 None
OpBranchConditional %6 %326 %327
%326 = OpLabel
OpBranch %325
%327 = OpLabel
OpSelectionMerge %328 None
OpBranchConditional %6 %329 %330
%329 = OpLabel
OpBranch %328
%330 = OpLabel
OpSelectionMerge %331 None
OpBranchConditional %6 %332 %333
%332 = OpLabel
OpBranch %331
%333 = OpLabel
OpSelectionMerge %334 None
OpBranchConditional %6 %335 %336
%335 = OpLabel
OpBranch %334
%336 = OpLabel
OpSelectionMerge %337 None
OpBranchConditional %6 %338 %339
%338 = OpLabel
OpBranch %337
%339 = OpLabel
OpSelectionMerge %340 None
OpBranchConditional %6 %341 %342
%341 = OpLabel
OpBranch %340
%342 = OpLabel
OpSelectionMerge %343 None
OpBranchConditional %6 %344 %345
%344 = OpLabel
OpBranch %343
%345 = OpLabel
OpSelectionMerge %346 None
OpBranchConditional %6 %347 %348
%347 = OpLabel
OpBranch %346
%348 = OpLabel
OpSelectionMerge %349 None
OpBranchConditional %6 %350 %351
%350 = OpLabel
OpBranch %349
%351 = OpLabel
OpSelectionMerge %352 None
OpBranchConditional %6 %353 %354
%353 = OpLabel
OpBranch %352
%354 = OpLabel
OpSelectionMerge %355 None
OpBranchConditional %6 %356 %357
%356 = OpLabel
OpBranch %355
%357 = OpLabel
OpSelectionMerge %358 None
OpBranchConditional %6 %359 %360
%359 = OpLabel
OpBranch %358
%360 = OpLabel
OpSelectionMerge %361 None
OpBranchConditional %6 %362 %363
%362 = OpLabel
OpBranch %361
%363 = OpLabel
OpSelectionMerge %364 None
OpBranchConditional %6 %365 %366
%365 = OpLabel
OpBranch %364
%366 = OpLabel
OpSelectionMerge %367 None
OpBranchConditional %6 %368 %369
%368 = OpLabel
OpBranch %367
%369 = OpLabel
OpSelectionMerge %370 None
OpBranchConditional %6 %371 %372
%371 = OpLabel
OpBranch %370
%372 = OpLabel
OpSelectionMerge %373 None
OpBranchConditional %6 %374 %375
%374 = OpLabel
OpBranch %373
%375 = OpLabel
OpSelectionMerge %376 None
OpBranchConditional %6 %377 %378
%377 = OpLabel
OpBranch %376
%378 = OpLabel
OpSelectionMerge %379 None
OpBranchConditional %6 %380 %379
%380 = OpLabel
OpBranch %379
%379 = OpLabel
OpBranch %376
%376 = OpLabel
OpBranch %373
%373 = OpLabel
OpBranch %370
%370 = OpLabel
OpBranch %367
%367 = OpLabel
OpBranch %364
%364 = OpLabel
OpBranch %361
%361 = OpLabel
OpBranch %358
%358 = OpLabel
OpBranch %355
%355 = OpLabel
OpBranch %352
%352 = OpLabel
OpBranch %349
%349 = OpLabel
OpBranch %346
%346 = OpLabel
OpBranch %343
%343 = OpLabel
OpBranch %340
%340 = OpLabel
OpBranch %337
%337 = OpLabel
OpBranch %334
%334 = OpLabel
OpBranch %331
%331 = OpLabel
OpBranch %328
%328 = OpLabel
OpBranch %325
%325 = OpLabel
OpBranch %322
%322 = OpLabel
OpBranch %319
%319 = OpLabel
OpBranch %316
%316 = OpLabel
OpBranch %313
%313 = OpLabel
OpBranch %310
%310 = OpLabel
OpBranch %307
%307 = OpLabel
OpBranch %304
%304 = OpLabel
OpBranch %301
%301 = OpLabel
OpBranch %298
%298 = OpLabel
OpBranch %295
%295 = OpLabel
OpBranch %292
%292 = OpLabel
OpBranch %289
%289 = OpLabel
OpBranch %286
%286 = OpLabel
OpBranch %283
%283 = OpLabel
OpBranch %280
%280 = OpLabel
OpBranch %277
%277 = OpLabel
OpBranch %274
%274 = OpLabel
OpBranch %271
%271 = OpLabel
OpBranch %268
%268 = OpLabel
OpBranch %265
%265 = OpLabel
OpBranch %262
%262 = OpLabel
OpBranch %259
%259 = OpLabel
OpBranch %256
%256 = OpLabel
OpBranch %253
%253 = OpLabel
OpBranch %250
%250 = OpLabel
OpBranch %247
%247 = OpLabel
OpBranch %244
%244 = OpLabel
OpBranch %241
%241 = OpLabel
OpBranch %238
%238 = OpLabel
OpBranch %235
%235 = OpLabel
OpBranch %232
%232 = OpLabel
OpBranch %229
%229 = OpLabel
OpBranch %226
%226 = OpLabel
OpBranch %223
%223 = OpLabel
OpBranch %220
%220 = OpLabel
OpBranch %217
%217 = OpLabel
OpBranch %214
%214 = OpLabel
OpBranch %211
%211 = OpLabel
OpBranch %208
%208 = OpLabel
OpBranch %205
%205 = OpLabel
OpBranch %202
%202 = OpLabel
OpBranch %199
%199 = OpLabel
OpBranch %196
%196 = OpLabel
OpBranch %193
%193 = OpLabel
OpBranch %190
%190 = OpLabel
OpBranch %187
%187 = OpLabel
OpBranch %184
%184 = OpLabel
OpBranch %181
%181 = OpLabel
OpBranch %178
%178 = OpLabel
OpBranch %175
%175 = OpLabel
OpBranch %172
%172 = OpLabel
OpBranch %169
%169 = OpLabel
OpBranch %166
%166 = OpLabel
OpBranch %163
%163 = OpLabel
OpBranch %160
%160 = OpLabel
OpBranch %157
%157 = OpLabel
OpBranch %154
%154 = OpLabel
OpBranch %151
%151 = OpLabel
OpBranch %148
%148 = OpLabel
OpBranch %145
%145 = OpLabel
OpBranch %142
%142 = OpLabel
OpBranch %139
%139 = OpLabel
OpBranch %136
%136 = OpLabel
OpBranch %133
%133 = OpLabel
OpBranch %130
%130 = OpLabel
OpBranch %127
%127 = OpLabel
OpBranch %124
%124 = OpLabel
OpBranch %121
%121 = OpLabel
OpBranch %118
%118 = OpLabel
OpBranch %115
%115 = OpLabel
OpBranch %112
%112 = OpLabel
OpBranch %109
%109 = OpLabel
OpBranch %106
%106 = OpLabel
OpBranch %103
%103 = OpLabel
OpBranch %100
%100 = OpLabel
OpBranch %97
%97 = OpLabel
OpBranch %94
%94 = OpLabel
OpBranch %91
%91 = OpLabel
OpBranch %88
%88 = OpLabel
OpBranch %85
%85 = OpLabel
OpBranch %82
%82 = OpLabel
OpBranch %79
%79 = OpLabel
OpBranch %76
%76 = OpLabel
OpBranch %73
%73 = OpLabel
OpBranch %70
%70 = OpLabel
OpBranch %67
%67 = OpLabel
OpBranch %64
%64 = OpLabel
OpBranch %61
%61 = OpLabel
OpBranch %58
%58 = OpLabel
OpBranch %55
%55 = OpLabel
OpBranch %52
%52 = OpLabel
OpBranch %49
%49 = OpLabel
OpBranch %46
%46 = OpLabel
OpBranch %43
%43 = OpLabel
OpBranch %40
%40 = OpLabel
OpBranch %37
%37 = OpLabel
OpBranch %34
%34 = OpLabel
OpBranch %31
%31 = OpLabel
OpBranch %28
%28 = OpLabel
OpBranch %25
%25 = OpLabel
OpBranch %22
%22 = OpLabel
OpBranch %19
%19 = OpLabel
OpBranch %16
%16 = OpLabel
OpBranch %13
%13 = OpLabel
OpBranch %10
%10 = OpLabel
OpBranch %7
%7 = OpLabel
OpReturn
OpFunctionEnd

View File

@ -0,0 +1,129 @@
@compute @workgroup_size(1, 1, 1)
fn main() {
if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
} else if (false) {
}
}