From ada560b289b757ac8289fd2d6a54adf203beb753 Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Thu, 20 May 2021 21:31:37 +0000 Subject: [PATCH] writer/hlsl: Fix continuing block emission Inline the `continuing` block in the places where `continue` is called. Simplifies the emission, and fixes emission of `let` statements in the loop. Also fix random indenting of intrinsic functions. Fixed: tint:744 Fixed: tint:818 Change-Id: I06994dbc724bc646e0435a1035b00760eaf5f5ab Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51784 Kokoro: Kokoro Reviewed-by: David Neto Commit-Queue: Ben Clayton --- src/writer/hlsl/generator_impl.cc | 180 +- src/writer/hlsl/generator_impl.h | 29 +- .../hlsl/generator_impl_continue_test.cc | 11 +- .../hlsl/generator_impl_intrinsic_test.cc | 2 +- src/writer/hlsl/generator_impl_loop_test.cc | 56 +- test/bug/tint/744.wgsl | 29 + test/bug/tint/744.wgsl.expected.hlsl | 41 + test/bug/tint/744.wgsl.expected.msl | 47 + test/bug/tint/744.wgsl.expected.spvasm | 127 ++ test/bug/tint/744.wgsl.expected.wgsl | 44 + test/bug/tint/749.spvasm.expected.hlsl | 1551 ++++++++++++++++- .../local/ptr_function.wgsl.expected.hlsl | 8 +- .../load/local/ptr_private.wgsl.expected.hlsl | 9 +- .../load/local/ptr_uniform.wgsl.expected.hlsl | 13 +- .../local/ptr_workgroup.wgsl.expected.hlsl | 10 +- test/samples/compute_boids.wgsl.expected.hlsl | 54 +- 16 files changed, 2009 insertions(+), 202 deletions(-) create mode 100644 test/bug/tint/744.wgsl create mode 100644 test/bug/tint/744.wgsl.expected.hlsl create mode 100644 test/bug/tint/744.wgsl.expected.msl create mode 100644 test/bug/tint/744.wgsl.expected.spvasm create mode 100644 test/bug/tint/744.wgsl.expected.wgsl diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc index 8d2deaac60..626cc77968 100644 --- a/src/writer/hlsl/generator_impl.cc +++ b/src/writer/hlsl/generator_impl.cc @@ -34,6 +34,7 @@ #include "src/sem/variable.h" #include "src/transform/calculate_array_length.h" #include "src/transform/decompose_storage_access.h" +#include "src/utils/scoped_assignment.h" #include "src/writer/append_vector.h" #include "src/writer/float_to_string.h" @@ -422,20 +423,14 @@ bool GeneratorImpl::EmitBinary(std::ostream& pre, bool GeneratorImpl::EmitBlock(std::ostream& out, const ast::BlockStatement* stmt) { - out << "{" << std::endl; - increment_indent(); - - for (auto* s : *stmt) { - if (!EmitStatement(out, s)) { - return false; + return EmitBlockBraces(out, [&] { + for (auto* s : *stmt) { + if (!EmitStatement(out, s)) { + return false; + } } - } - - decrement_indent(); - make_indent(out); - out << "}"; - - return true; + return true; + }); } bool GeneratorImpl::EmitBlockAndNewline(std::ostream& out, @@ -614,7 +609,6 @@ bool GeneratorImpl::EmitCall(std::ostream& pre, return false; } - make_indent(out); out << name << "("; bool first = true; @@ -1338,6 +1332,9 @@ bool GeneratorImpl::EmitTypeConstructor(std::ostream& pre, } bool GeneratorImpl::EmitContinue(std::ostream& out, ast::ContinueStatement*) { + if (!emit_continuing_(out)) { + return false; + } make_indent(out); out << "continue;" << std::endl; return true; @@ -2192,100 +2189,30 @@ bool GeneratorImpl::EmitZeroValue(std::ostream& out, const sem::Type* type) { } bool GeneratorImpl::EmitLoop(std::ostream& out, ast::LoopStatement* stmt) { - loop_emission_counter_++; + make_indent(out); - std::string guard = generate_name("tint_hlsl_is_first_" + - std::to_string(loop_emission_counter_)); + auto emit_continuing = [this, stmt](std::ostream& o) { + if (stmt->has_continuing()) { + make_indent(o); + if (!EmitBlock(o, stmt->continuing())) { + return false; + } + o << std::endl; + } + return true; + }; - if (stmt->has_continuing()) { - make_indent(out); - - // Continuing variables get their own scope. - out << "{" << std::endl; - increment_indent(); - - make_indent(out); - out << "bool " << guard << " = true;" << std::endl; - - // A continuing block may use variables declared in the method body. As a - // first pass, if we have a continuing, we pull all declarations outside - // the for loop into the continuing scope. Then, the variable declarations - // will be turned into assignments. - for (auto* s : *stmt->body()) { - if (auto* v = s->As()) { - if (!EmitVariable(out, v->variable(), true)) { - return false; - } + TINT_SCOPED_ASSIGNMENT(emit_continuing_, emit_continuing); + bool ok = EmitBlockBraces(out, "while (true)", [&] { + for (auto* s : stmt->body()->statements()) { + if (!EmitStatement(out, s)) { + return false; } } - } - - make_indent(out); - out << "for(;;) {" << std::endl; - increment_indent(); - - if (stmt->has_continuing()) { - make_indent(out); - out << "if (!" << guard << ") "; - - if (!EmitBlockAndNewline(out, stmt->continuing())) { - return false; - } - - make_indent(out); - out << guard << " = false;" << std::endl; - out << std::endl; - } - - for (auto* s : *(stmt->body())) { - // If we have a continuing block we've already emitted the variable - // declaration before the loop, so treat it as an assignment. - if (auto* decl = s->As()) { - if (stmt->has_continuing()) { - make_indent(out); - - auto* var = decl->variable(); - - std::ostringstream pre; - std::ostringstream constructor_out; - if (var->constructor() != nullptr) { - if (!EmitExpression(pre, constructor_out, var->constructor())) { - return false; - } - } - out << pre.str(); - - out << builder_.Symbols().NameFor(var->symbol()) << " = "; - if (var->constructor() != nullptr) { - out << constructor_out.str(); - } else { - auto* type = builder_.Sem().Get(var)->Type()->UnwrapRef(); - if (!EmitZeroValue(out, type)) { - return false; - } - } - out << ";" << std::endl; - continue; - } - } - - if (!EmitStatement(out, s)) { - return false; - } - } - - decrement_indent(); - make_indent(out); - out << "}" << std::endl; - - // Close the scope for any continuing variables. - if (stmt->has_continuing()) { - decrement_indent(); - make_indent(out); - out << "}" << std::endl; - } - - return true; + return emit_continuing(out); + }); + out << std::endl; + return ok; } bool GeneratorImpl::EmitMemberAccessor(std::ostream& pre, @@ -2379,7 +2306,7 @@ bool GeneratorImpl::EmitStatement(std::ostream& out, ast::Statement* stmt) { return EmitSwitch(out, s); } if (auto* v = stmt->As()) { - return EmitVariable(out, v->variable(), false); + return EmitVariable(out, v->variable()); } diagnostics_.add_error("unknown statement type: " + builder_.str(stmt)); @@ -2662,9 +2589,7 @@ bool GeneratorImpl::EmitUnaryOp(std::ostream& pre, return true; } -bool GeneratorImpl::EmitVariable(std::ostream& out, - ast::Variable* var, - bool skip_constructor) { +bool GeneratorImpl::EmitVariable(std::ostream& out, ast::Variable* var) { make_indent(out); auto* sem = builder_.Sem().Get(var); @@ -2677,19 +2602,17 @@ bool GeneratorImpl::EmitVariable(std::ostream& out, } std::ostringstream constructor_out; - if (!skip_constructor) { - constructor_out << " = "; + constructor_out << " = "; - if (var->constructor()) { - std::ostringstream pre; - if (!EmitExpression(pre, constructor_out, var->constructor())) { - return false; - } - out << pre.str(); - } else { - if (!EmitZeroValue(constructor_out, type)) { - return false; - } + if (var->constructor()) { + std::ostringstream pre; + if (!EmitExpression(pre, constructor_out, var->constructor())) { + return false; + } + out << pre.str(); + } else { + if (!EmitZeroValue(constructor_out, type)) { + return false; } } @@ -2789,6 +2712,23 @@ std::string GeneratorImpl::get_buffer_name(ast::Expression* expr) { return ""; } +template +bool GeneratorImpl::EmitBlockBraces(std::ostream& out, + const std::string& prefix, + F&& cb) { + out << prefix << (prefix.empty() ? "{" : " {") << std::endl; + increment_indent(); + + if (!cb()) { + return false; + } + + decrement_indent(); + make_indent(out); + out << "}"; + return true; +} + } // namespace hlsl } // namespace writer } // namespace tint diff --git a/src/writer/hlsl/generator_impl.h b/src/writer/hlsl/generator_impl.h index 250782b426..9fee7845d8 100644 --- a/src/writer/hlsl/generator_impl.h +++ b/src/writer/hlsl/generator_impl.h @@ -18,6 +18,7 @@ #include #include #include +#include #include "src/ast/assignment_statement.h" #include "src/ast/bitcast_expression.h" @@ -319,11 +320,8 @@ class GeneratorImpl : public TextGenerator { /// Handles generating a variable /// @param out the output stream /// @param var the variable to generate - /// @param skip_constructor set true if the constructor should be skipped /// @returns true if the variable was emitted - bool EmitVariable(std::ostream& out, - ast::Variable* var, - bool skip_constructor); + bool EmitVariable(std::ostream& out, ast::Variable* var); /// Handles generating a program scope constant variable /// @param out the output stream /// @param var the variable to emit @@ -387,10 +385,31 @@ class GeneratorImpl : public TextGenerator { return builder_.TypeOf(type); } + /// Emits `prefix`, followed by an opening brace `{`, then calls `cb` to emit + /// the block body, then finally emits the closing brace `}`. + /// @param out the output stream + /// @param prefix the string to emit before the opening brace + /// @param cb a function or function-like object with the signature `bool()` + /// that emits the block body. + /// @returns the return value of `cb`. + template + bool EmitBlockBraces(std::ostream& out, const std::string& prefix, F&& cb); + + /// Emits an opening brace `{`, then calls `cb` to emit the block body, then + /// finally emits the closing brace `}`. + /// @param out the output stream + /// @param cb a function or function-like object with the signature `bool()` + /// that emits the block body. + /// @returns the return value of `cb`. + template + bool EmitBlockBraces(std::ostream& out, F&& cb) { + return EmitBlockBraces(out, "", std::forward(cb)); + } + ProgramBuilder builder_; Symbol current_ep_sym_; bool generating_entry_point_ = false; - uint32_t loop_emission_counter_ = 0; + std::function emit_continuing_; ScopeStack global_variables_; std::unordered_map ep_sym_to_in_data_; std::unordered_map ep_sym_to_out_data_; diff --git a/src/writer/hlsl/generator_impl_continue_test.cc b/src/writer/hlsl/generator_impl_continue_test.cc index 07d6b5775d..6a2482b6ae 100644 --- a/src/writer/hlsl/generator_impl_continue_test.cc +++ b/src/writer/hlsl/generator_impl_continue_test.cc @@ -22,15 +22,18 @@ namespace { using HlslGeneratorImplTest_Continue = TestHelper; TEST_F(HlslGeneratorImplTest_Continue, Emit_Continue) { - auto* c = create(); - WrapInFunction(Loop(Block(c))); + auto* loop = Loop(Block(create())); + WrapInFunction(loop); GeneratorImpl& gen = Build(); gen.increment_indent(); - ASSERT_TRUE(gen.EmitStatement(out, c)) << gen.error(); - EXPECT_EQ(result(), " continue;\n"); + ASSERT_TRUE(gen.EmitStatement(out, loop)) << gen.error(); + EXPECT_EQ(result(), R"( while (true) { + continue; + } +)"); } } // namespace diff --git a/src/writer/hlsl/generator_impl_intrinsic_test.cc b/src/writer/hlsl/generator_impl_intrinsic_test.cc index e6f3966a7f..866032651f 100644 --- a/src/writer/hlsl/generator_impl_intrinsic_test.cc +++ b/src/writer/hlsl/generator_impl_intrinsic_test.cc @@ -268,7 +268,7 @@ TEST_F(HlslGeneratorImplTest_Intrinsic, Intrinsic_Call) { gen.increment_indent(); ASSERT_TRUE(gen.EmitExpression(pre, out, call)) << gen.error(); - EXPECT_EQ(result(), " dot(param1, param2)"); + EXPECT_EQ(result(), "dot(param1, param2)"); } TEST_F(HlslGeneratorImplTest_Intrinsic, Pack4x8Snorm) { diff --git a/src/writer/hlsl/generator_impl_loop_test.cc b/src/writer/hlsl/generator_impl_loop_test.cc index 75630015b3..906d56eda4 100644 --- a/src/writer/hlsl/generator_impl_loop_test.cc +++ b/src/writer/hlsl/generator_impl_loop_test.cc @@ -34,7 +34,7 @@ TEST_F(HlslGeneratorImplTest_Loop, Emit_Loop) { gen.increment_indent(); ASSERT_TRUE(gen.EmitStatement(out, l)) << gen.error(); - EXPECT_EQ(result(), R"( for(;;) { + EXPECT_EQ(result(), R"( while (true) { discard; } )"); @@ -52,15 +52,10 @@ TEST_F(HlslGeneratorImplTest_Loop, Emit_LoopWithContinuing) { gen.increment_indent(); ASSERT_TRUE(gen.EmitStatement(out, l)) << gen.error(); - EXPECT_EQ(result(), R"( { - bool tint_hlsl_is_first_1 = true; - for(;;) { - if (!tint_hlsl_is_first_1) { - return; - } - tint_hlsl_is_first_1 = false; - - discard; + EXPECT_EQ(result(), R"( while (true) { + discard; + { + return; } } )"); @@ -89,26 +84,16 @@ TEST_F(HlslGeneratorImplTest_Loop, Emit_LoopNestedWithContinuing) { gen.increment_indent(); ASSERT_TRUE(gen.EmitStatement(out, outer)) << gen.error(); - EXPECT_EQ(result(), R"( { - bool tint_hlsl_is_first_1 = true; - for(;;) { - if (!tint_hlsl_is_first_1) { - lhs = rhs; - } - tint_hlsl_is_first_1 = false; - + EXPECT_EQ(result(), R"( while (true) { + while (true) { + discard; { - bool tint_hlsl_is_first_2 = true; - for(;;) { - if (!tint_hlsl_is_first_2) { - return; - } - tint_hlsl_is_first_2 = false; - - discard; - } + return; } } + { + lhs = rhs; + } } )"); } @@ -153,18 +138,11 @@ TEST_F(HlslGeneratorImplTest_Loop, Emit_LoopWithVarUsedInContinuing) { gen.increment_indent(); ASSERT_TRUE(gen.EmitStatement(out, outer)) << gen.error(); - EXPECT_EQ(result(), R"( { - bool tint_hlsl_is_first_1 = true; - float lhs; - float other; - for(;;) { - if (!tint_hlsl_is_first_1) { - lhs = rhs; - } - tint_hlsl_is_first_1 = false; - - lhs = 2.400000095f; - other = 0.0f; + EXPECT_EQ(result(), R"( while (true) { + float lhs = 2.400000095f; + float other = 0.0f; + { + lhs = rhs; } } )"); diff --git a/test/bug/tint/744.wgsl b/test/bug/tint/744.wgsl new file mode 100644 index 0000000000..8950ed5d27 --- /dev/null +++ b/test/bug/tint/744.wgsl @@ -0,0 +1,29 @@ +[[block]] struct Uniforms { + aShape : vec2; + bShape : vec2; + outShape : vec2; +}; +[[block]] struct Matrix { + numbers: array; +}; + +[[group(0), binding(0)]] var firstMatrix : [[access(read)]] Matrix; +[[group(0), binding(1)]] var secondMatrix : [[access(read)]] Matrix; +[[group(0), binding(2)]] var resultMatrix : [[access(write)]] Matrix; +[[group(0), binding(3)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(2,2,1)]] +fn main([[builtin(global_invocation_id)]] global_id : vec3) { + let resultCell : vec2 = vec2(global_id.y, global_id.x); + let dimInner : u32 = uniforms.aShape.y; + let dimOutter: u32 = uniforms.outShape.y; + var result : u32 = 0u; + for (var i : u32 = 0u; i < dimInner; i = i + 1u) { + let a : u32 = i + resultCell.x * dimInner; + let b : u32 = resultCell.y + i * dimOutter; + result = result + firstMatrix.numbers[a] * secondMatrix.numbers[b]; + } + + let index : u32 = resultCell.y + resultCell.x * dimOutter; + resultMatrix.numbers[index] = result; +} diff --git a/test/bug/tint/744.wgsl.expected.hlsl b/test/bug/tint/744.wgsl.expected.hlsl new file mode 100644 index 0000000000..d026848268 --- /dev/null +++ b/test/bug/tint/744.wgsl.expected.hlsl @@ -0,0 +1,41 @@ +struct Uniforms { + uint2 aShape; + uint2 bShape; + uint2 outShape; +}; +struct tint_symbol_1 { + uint3 global_id : SV_DispatchThreadID; +}; + +ConstantBuffer uniforms : register(b3, space0); + +ByteAddressBuffer firstMatrix : register(t0, space0); +ByteAddressBuffer secondMatrix : register(t1, space0); +RWByteAddressBuffer resultMatrix : register(u2, space0); + +[numthreads(2, 2, 1)] +void main(tint_symbol_1 tint_symbol) { + const uint3 global_id = tint_symbol.global_id; + const uint2 resultCell = uint2(global_id.y, global_id.x); + const uint dimInner = uniforms.aShape.y; + const uint dimOutter = uniforms.outShape.y; + uint result = 0u; + { + uint i = 0u; + while (true) { + if (!((i < dimInner))) { + break; + } + const uint a = (i + (resultCell.x * dimInner)); + const uint b = (resultCell.y + (i * dimOutter)); + result = (result + (firstMatrix.Load((4u * a)) * secondMatrix.Load((4u * b)))); + { + i = (i + 1u); + } + } + } + const uint index = (resultCell.y + (resultCell.x * dimOutter)); + resultMatrix.Store((4u * index), asuint(result)); + return; +} + diff --git a/test/bug/tint/744.wgsl.expected.msl b/test/bug/tint/744.wgsl.expected.msl new file mode 100644 index 0000000000..a99c53a55e --- /dev/null +++ b/test/bug/tint/744.wgsl.expected.msl @@ -0,0 +1,47 @@ +#include + +using namespace metal; +struct Uniforms { + /* 0x0000 */ packed_uint2 aShape; + /* 0x0008 */ packed_uint2 bShape; + /* 0x0010 */ packed_uint2 outShape; +}; +struct Matrix { + /* 0x0000 */ uint numbers[1]; +}; +struct tint_symbol_2 { + uint3 global_id [[thread_position_in_grid]]; +}; + +kernel void tint_symbol(tint_symbol_2 tint_symbol_1 [[stage_in]], constant Uniforms& uniforms [[buffer(3)]], const device Matrix& firstMatrix [[buffer(0)]], const device Matrix& secondMatrix [[buffer(1)]], device Matrix& resultMatrix [[buffer(2)]]) { + uint3 const global_id = tint_symbol_1.global_id; + uint2 const resultCell = uint2(global_id.y, global_id.x); + uint const dimInner = uniforms.aShape.y; + uint const dimOutter = uniforms.outShape.y; + uint result = 0u; + { + uint i = 0u; + { + bool tint_msl_is_first_1 = true; + uint const a; + uint const b; + for(;;) { + if (!tint_msl_is_first_1) { + i = (i + 1u); + } + tint_msl_is_first_1 = false; + + if (!((i < dimInner))) { + break; + } + a = (i + (resultCell.x * dimInner)); + b = (resultCell.y + (i * dimOutter)); + result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b])); + } + } + } + uint const index = (resultCell.y + (resultCell.x * dimOutter)); + resultMatrix.numbers[index] = result; + return; +} + diff --git a/test/bug/tint/744.wgsl.expected.spvasm b/test/bug/tint/744.wgsl.expected.spvasm new file mode 100644 index 0000000000..b7d0993740 --- /dev/null +++ b/test/bug/tint/744.wgsl.expected.spvasm @@ -0,0 +1,127 @@ +; SPIR-V +; Version: 1.3 +; Generator: Google Tint Compiler; 0 +; Bound: 71 +; Schema: 0 + OpCapability Shader + OpMemoryModel Logical GLSL450 + OpEntryPoint GLCompute %main "main" %tint_symbol + OpExecutionMode %main LocalSize 2 2 1 + OpName %Matrix "Matrix" + OpMemberName %Matrix 0 "numbers" + OpName %firstMatrix "firstMatrix" + OpName %secondMatrix "secondMatrix" + OpName %resultMatrix "resultMatrix" + OpName %Uniforms "Uniforms" + OpMemberName %Uniforms 0 "aShape" + OpMemberName %Uniforms 1 "bShape" + OpMemberName %Uniforms 2 "outShape" + OpName %uniforms "uniforms" + OpName %tint_symbol "tint_symbol" + OpName %main "main" + OpName %result "result" + OpName %i "i" + OpDecorate %Matrix Block + OpMemberDecorate %Matrix 0 Offset 0 + OpDecorate %_runtimearr_uint ArrayStride 4 + OpDecorate %firstMatrix NonWritable + OpDecorate %firstMatrix DescriptorSet 0 + OpDecorate %firstMatrix Binding 0 + OpDecorate %secondMatrix NonWritable + OpDecorate %secondMatrix DescriptorSet 0 + OpDecorate %secondMatrix Binding 1 + OpDecorate %resultMatrix NonReadable + OpDecorate %resultMatrix DescriptorSet 0 + OpDecorate %resultMatrix Binding 2 + OpDecorate %Uniforms Block + OpMemberDecorate %Uniforms 0 Offset 0 + OpMemberDecorate %Uniforms 1 Offset 8 + OpMemberDecorate %Uniforms 2 Offset 16 + OpDecorate %uniforms DescriptorSet 0 + OpDecorate %uniforms Binding 3 + OpDecorate %tint_symbol BuiltIn GlobalInvocationId + %uint = OpTypeInt 32 0 +%_runtimearr_uint = OpTypeRuntimeArray %uint + %Matrix = OpTypeStruct %_runtimearr_uint +%_ptr_StorageBuffer_Matrix = OpTypePointer StorageBuffer %Matrix +%firstMatrix = OpVariable %_ptr_StorageBuffer_Matrix StorageBuffer +%secondMatrix = OpVariable %_ptr_StorageBuffer_Matrix StorageBuffer +%resultMatrix = OpVariable %_ptr_StorageBuffer_Matrix StorageBuffer + %v2uint = OpTypeVector %uint 2 + %Uniforms = OpTypeStruct %v2uint %v2uint %v2uint +%_ptr_Uniform_Uniforms = OpTypePointer Uniform %Uniforms + %uniforms = OpVariable %_ptr_Uniform_Uniforms Uniform + %v3uint = OpTypeVector %uint 3 +%_ptr_Input_v3uint = OpTypePointer Input %v3uint +%tint_symbol = OpVariable %_ptr_Input_v3uint Input + %void = OpTypeVoid + %15 = OpTypeFunction %void + %uint_1 = OpConstant %uint 1 +%_ptr_Input_uint = OpTypePointer Input %uint + %uint_0 = OpConstant %uint 0 +%_ptr_Uniform_uint = OpTypePointer Uniform %uint + %uint_2 = OpConstant %uint 2 +%_ptr_Function_uint = OpTypePointer Function %uint + %35 = OpConstantNull %uint + %bool = OpTypeBool +%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint + %main = OpFunction %void None %15 + %18 = OpLabel + %result = OpVariable %_ptr_Function_uint Function %35 + %i = OpVariable %_ptr_Function_uint Function %35 + %21 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_1 + %22 = OpLoad %uint %21 + %24 = OpAccessChain %_ptr_Input_uint %tint_symbol %uint_0 + %25 = OpLoad %uint %24 + %26 = OpCompositeConstruct %v2uint %22 %25 + %28 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_0 %uint_1 + %29 = OpLoad %uint %28 + %31 = OpAccessChain %_ptr_Uniform_uint %uniforms %uint_2 %uint_1 + %32 = OpLoad %uint %31 + OpStore %result %uint_0 + OpStore %i %uint_0 + OpBranch %37 + %37 = OpLabel + OpLoopMerge %38 %39 None + OpBranch %40 + %40 = OpLabel + %42 = OpLoad %uint %i + %43 = OpULessThan %bool %42 %29 + %41 = OpLogicalNot %bool %43 + OpSelectionMerge %45 None + OpBranchConditional %41 %46 %45 + %46 = OpLabel + OpBranch %38 + %45 = OpLabel + %47 = OpLoad %uint %i + %48 = OpCompositeExtract %uint %26 0 + %49 = OpIMul %uint %48 %29 + %50 = OpIAdd %uint %47 %49 + %51 = OpCompositeExtract %uint %26 1 + %52 = OpLoad %uint %i + %53 = OpIMul %uint %52 %32 + %54 = OpIAdd %uint %51 %53 + %55 = OpLoad %uint %result + %57 = OpAccessChain %_ptr_StorageBuffer_uint %firstMatrix %uint_0 %50 + %58 = OpLoad %uint %57 + %59 = OpAccessChain %_ptr_StorageBuffer_uint %secondMatrix %uint_0 %54 + %60 = OpLoad %uint %59 + %61 = OpIMul %uint %58 %60 + %62 = OpIAdd %uint %55 %61 + OpStore %result %62 + OpBranch %39 + %39 = OpLabel + %63 = OpLoad %uint %i + %64 = OpIAdd %uint %63 %uint_1 + OpStore %i %64 + OpBranch %37 + %38 = OpLabel + %65 = OpCompositeExtract %uint %26 1 + %66 = OpCompositeExtract %uint %26 0 + %67 = OpIMul %uint %66 %32 + %68 = OpIAdd %uint %65 %67 + %69 = OpAccessChain %_ptr_StorageBuffer_uint %resultMatrix %uint_0 %68 + %70 = OpLoad %uint %result + OpStore %69 %70 + OpReturn + OpFunctionEnd diff --git a/test/bug/tint/744.wgsl.expected.wgsl b/test/bug/tint/744.wgsl.expected.wgsl new file mode 100644 index 0000000000..cc4accafe2 --- /dev/null +++ b/test/bug/tint/744.wgsl.expected.wgsl @@ -0,0 +1,44 @@ +[[block]] +struct Uniforms { + aShape : vec2; + bShape : vec2; + outShape : vec2; +}; + +[[block]] +struct Matrix { + numbers : array; +}; + +[[group(0), binding(0)]] var firstMatrix : [[access(read)]] Matrix; + +[[group(0), binding(1)]] var secondMatrix : [[access(read)]] Matrix; + +[[group(0), binding(2)]] var resultMatrix : [[access(write)]] Matrix; + +[[group(0), binding(3)]] var uniforms : Uniforms; + +[[stage(compute), workgroup_size(2, 2, 1)]] +fn main([[builtin(global_invocation_id)]] global_id : vec3) { + let resultCell : vec2 = vec2(global_id.y, global_id.x); + let dimInner : u32 = uniforms.aShape.y; + let dimOutter : u32 = uniforms.outShape.y; + var result : u32 = 0u; + { + var i : u32 = 0u; + loop { + if (!((i < dimInner))) { + break; + } + let a : u32 = (i + (resultCell.x * dimInner)); + let b : u32 = (resultCell.y + (i * dimOutter)); + result = (result + (firstMatrix.numbers[a] * secondMatrix.numbers[b])); + + continuing { + i = (i + 1u); + } + } + } + let index : u32 = (resultCell.y + (resultCell.x * dimOutter)); + resultMatrix.numbers[index] = result; +} diff --git a/test/bug/tint/749.spvasm.expected.hlsl b/test/bug/tint/749.spvasm.expected.hlsl index 76a11e3a73..e6674e672b 100644 --- a/test/bug/tint/749.spvasm.expected.hlsl +++ b/test/bug/tint/749.spvasm.expected.hlsl @@ -1 +1,1550 @@ -SKIP: crbug.com/tint/818: [HLSL] Loop emission broken for `let` function declarations +struct QuicksortObject { + int numbers[10]; +}; +struct buf0 { + float2 resolution; +}; + +ConstantBuffer x_188 : register(b0, space0); + +struct main_in { + float4 gl_FragCoord : SV_Position; +}; + +struct main_out { + float4 x_GLF_color : SV_Target0; +}; + +static QuicksortObject obj; + +void swap_i1_i1_(inout int i, inout int j) { + int temp = 0; + const int x_932 = temp; + temp = 0; + temp = x_932; + const float3 x_523 = float3(float3(1.0f, 2.0f, 3.0f).z, float3(1.0f, 2.0f, 3.0f).y, float3(1.0f, 2.0f, 3.0f).z); + const int x_933 = i; + i = 0; + i = x_933; + const int x_28 = i; + const int x_934 = j; + j = 0; + j = x_934; + const float3 x_524 = float3(x_523.y, x_523.x, x_523.y); + const int x_935 = temp; + temp = 0; + temp = x_935; + const int x_30_save = x_28; + const int x_936 = obj.numbers[x_30_save]; + obj.numbers[x_30_save] = 0; + obj.numbers[x_30_save] = x_936; + const int x_31 = obj.numbers[x_30_save]; + const int x_937 = temp; + temp = 0; + temp = x_937; + temp = x_31; + const int x_938 = j; + j = 0; + j = x_938; + const float3 x_525 = float3(x_523.z, float3(1.0f, 2.0f, 3.0f).x, x_523.y); + const int x_939 = i; + i = 0; + i = x_939; + const int x_32 = i; + const int x_940 = obj.numbers[x_30_save]; + obj.numbers[x_30_save] = 0; + obj.numbers[x_30_save] = x_940; + const int x_33 = j; + const int x_941 = i; + i = 0; + i = x_941; + const float3 x_526 = float3(x_525.x, x_525.z, x_525.z); + const int x_942 = obj.numbers[x_30_save]; + obj.numbers[x_30_save] = 0; + obj.numbers[x_30_save] = x_942; + const int x_34_save = x_33; + const int x_35 = obj.numbers[x_34_save]; + const QuicksortObject x_943 = obj; + const int tint_symbol[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_1 = {tint_symbol}; + obj = tint_symbol_1; + obj = x_943; + const float2 x_527 = float2(x_526.x, x_526.x); + const int x_36_save = x_32; + const float3 x_528 = float3(x_524.x, x_524.z, x_524.x); + obj.numbers[x_36_save] = x_35; + const QuicksortObject x_944 = obj; + const int tint_symbol_2[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_3 = {tint_symbol_2}; + obj = tint_symbol_3; + obj = x_944; + const float3 x_529 = float3(x_526.y, x_526.z, x_526.x); + const int x_945 = i; + i = 0; + i = x_945; + const int x_37 = j; + const int x_946 = temp; + temp = 0; + temp = x_946; + const float2 x_530 = float2(x_529.z, x_529.y); + const int x_947 = obj.numbers[x_34_save]; + obj.numbers[x_34_save] = 0; + obj.numbers[x_34_save] = x_947; + const int x_38 = temp; + const int x_948 = j; + j = 0; + j = x_948; + const float3 x_531 = float3(x_527.x, x_526.y, x_526.x); + const int x_949 = obj.numbers[x_36_save]; + obj.numbers[x_36_save] = 0; + obj.numbers[x_36_save] = x_949; + const QuicksortObject x_950 = obj; + const int tint_symbol_4[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_5 = {tint_symbol_4}; + obj = tint_symbol_5; + obj = x_950; + const float3 x_532 = float3(x_528.x, x_528.y, x_528.x); + const int x_951 = obj.numbers[x_34_save]; + obj.numbers[x_34_save] = 0; + obj.numbers[x_34_save] = x_951; + obj.numbers[x_37] = x_38; + return; +} + +int performPartition_i1_i1_(inout int l, inout int h) { + int param_3 = 0; + int i_1 = 0; + int j_1 = 0; + int param_2 = 0; + int param_1 = 0; + int param = 0; + int pivot = 0; + float2 x_537 = float2(0.0f, 0.0f); + float3 x_538 = float3(0.0f, 0.0f, 0.0f); + const int x_952 = h; + h = 0; + h = x_952; + const int x_41 = h; + const int x_953 = l; + l = 0; + l = x_953; + const int x_42_save = x_41; + const int x_954 = obj.numbers[x_42_save]; + obj.numbers[x_42_save] = 0; + obj.numbers[x_42_save] = x_954; + const int x_43 = obj.numbers[x_42_save]; + const int x_955 = param_3; + param_3 = 0; + param_3 = x_955; + const float3 x_534 = float3(float3(1.0f, 2.0f, 3.0f).z, float3(1.0f, 2.0f, 3.0f).x, float3(1.0f, 2.0f, 3.0f).z); + const int x_956 = param_1; + param_1 = 0; + param_1 = x_956; + pivot = x_43; + const int x_45 = l; + const int x_957 = h; + h = 0; + h = x_957; + const int x_958 = j_1; + j_1 = 0; + j_1 = x_958; + const float3 x_535 = float3(x_534.y, x_534.z, x_534.y); + const int x_959 = l; + l = 0; + l = x_959; + i_1 = (x_45 - asint(1u)); + const int x_49 = l; + const float3 x_536 = float3(x_534.x, x_534.z, x_535.x); + j_1 = 10; + const QuicksortObject x_960 = obj; + const int tint_symbol_6[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_7 = {tint_symbol_6}; + obj = tint_symbol_7; + obj = x_960; + while (true) { + const int x_961 = pivot; + pivot = 0; + pivot = x_961; + const int x_962 = param_1; + param_1 = 0; + param_1 = x_962; + const int x_55 = j_1; + const int x_963 = pivot; + pivot = 0; + pivot = x_963; + x_537 = float2(float3(1.0f, 2.0f, 3.0f).y, float3(1.0f, 2.0f, 3.0f).z); + const QuicksortObject x_964 = obj; + const int tint_symbol_8[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_9 = {tint_symbol_8}; + obj = tint_symbol_9; + obj = x_964; + const int x_56 = h; + const int x_965 = h; + h = 0; + h = x_965; + const int x_966 = param; + param = 0; + param = x_966; + const int x_967 = j_1; + j_1 = 0; + j_1 = x_967; + x_538 = float3(x_534.x, x_537.y, x_534.z); + const int x_968 = param; + param = 0; + param = x_968; + if ((x_55 <= (x_56 - asint(1u)))) { + } else { + break; + } + const int x_60 = j_1; + const int x_969 = obj.numbers[x_42_save]; + obj.numbers[x_42_save] = 0; + obj.numbers[x_42_save] = x_969; + const int x_61_save = x_60; + const int x_970 = h; + h = 0; + h = x_970; + const float3 x_539 = float3(x_537.x, x_535.z, x_537.x); + const int x_971 = param_1; + param_1 = 0; + param_1 = x_971; + const int x_62 = obj.numbers[x_61_save]; + const QuicksortObject x_972 = obj; + const int tint_symbol_10[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_11 = {tint_symbol_10}; + obj = tint_symbol_11; + obj = x_972; + const int x_63 = pivot; + const float2 x_540 = float2(float3(1.0f, 2.0f, 3.0f).y, x_534.z); + const int x_973 = i_1; + i_1 = 0; + i_1 = x_973; + const int x_974 = l; + l = 0; + l = x_974; + const float3 x_541 = float3(x_534.y, x_534.x, x_534.y); + const int x_975 = pivot; + pivot = 0; + pivot = x_975; + if ((x_62 <= x_63)) { + const float3 x_542 = float3(x_541.z, x_541.x, x_541.x); + const int x_976 = param_3; + param_3 = 0; + param_3 = x_976; + const int x_67 = i_1; + const int x_977 = pivot; + pivot = 0; + pivot = x_977; + const float2 x_543 = float2(x_539.x, x_541.y); + const int x_978 = i_1; + i_1 = 0; + i_1 = x_978; + const int x_979 = param; + param = 0; + param = x_979; + i_1 = (x_67 + asint(1u)); + const int x_980 = l; + l = 0; + l = x_980; + const float3 x_544 = float3(float3(1.0f, 2.0f, 3.0f).z, float3(1.0f, 2.0f, 3.0f).y, x_540.x); + const int x_70 = i_1; + const float2 x_545 = float2(x_537.y, x_538.x); + const int x_981 = param; + param = 0; + param = x_981; + param = x_70; + const int x_982 = param; + param = 0; + param = x_982; + const float2 x_546 = float2(x_545.x, x_545.x); + const int x_983 = i_1; + i_1 = 0; + i_1 = x_983; + const int x_72 = j_1; + param_1 = x_72; + const int x_984 = param_3; + param_3 = 0; + param_3 = x_984; + swap_i1_i1_(param, param_1); + const int x_985 = param_1; + param_1 = 0; + param_1 = x_985; + } + const QuicksortObject x_986 = obj; + const int tint_symbol_12[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_13 = {tint_symbol_12}; + obj = tint_symbol_13; + obj = x_986; + { + const int x_987 = h; + h = 0; + h = x_987; + const int x_74 = j_1; + const int x_988 = h; + h = 0; + h = x_988; + const float3 x_547 = float3(x_539.x, x_541.z, x_541.z); + const int x_989 = obj.numbers[x_61_save]; + obj.numbers[x_61_save] = 0; + obj.numbers[x_61_save] = x_989; + const int x_990 = param; + param = 0; + param = x_990; + j_1 = (1 + x_74); + const int x_991 = param_1; + param_1 = 0; + param_1 = x_991; + const float3 x_548 = float3(x_541.y, x_541.z, x_541.x); + const int x_992 = obj.numbers[x_61_save]; + obj.numbers[x_61_save] = 0; + obj.numbers[x_61_save] = x_992; + } + } + const int x_76 = i_1; + const int x_993 = obj.numbers[x_42_save]; + obj.numbers[x_42_save] = 0; + obj.numbers[x_42_save] = x_993; + const float2 x_549 = float2(x_534.x, x_534.y); + const QuicksortObject x_994 = obj; + const int tint_symbol_14[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_15 = {tint_symbol_14}; + obj = tint_symbol_15; + obj = x_994; + const int x_995 = h; + h = 0; + h = x_995; + i_1 = (1 + x_76); + const int x_996 = param_1; + param_1 = 0; + param_1 = x_996; + const int x_79 = i_1; + const int x_997 = j_1; + j_1 = 0; + j_1 = x_997; + const float2 x_550 = float2(x_534.x, x_534.x); + const int x_998 = param_1; + param_1 = 0; + param_1 = x_998; + param_2 = x_79; + const float2 x_551 = float2(x_534.y, x_536.x); + const int x_999 = pivot; + pivot = 0; + pivot = x_999; + const int x_81 = h; + const float2 x_552 = float2(x_550.x, x_549.y); + const int x_1000 = h; + h = 0; + h = x_1000; + param_3 = x_81; + const int x_1001 = i_1; + i_1 = 0; + i_1 = x_1001; + const float2 x_553 = float2(x_549.y, x_552.x); + const int x_1002 = h; + h = 0; + h = x_1002; + swap_i1_i1_(param_2, param_3); + const int x_1003 = l; + l = 0; + l = x_1003; + const float2 x_554 = float2(x_536.z, float3(1.0f, 2.0f, 3.0f).y); + const int x_1004 = param_1; + param_1 = 0; + param_1 = x_1004; + const int x_83 = i_1; + const int x_1005 = param; + param = 0; + param = x_1005; + const float2 x_555 = float2(x_534.y, x_534.x); + const int x_1006 = j_1; + j_1 = 0; + j_1 = x_1006; + return x_83; +} + +void quicksort_() { + int param_4 = 0; + int h_1 = 0; + int p = 0; + int l_1 = 0; + int top = 0; + int stack[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + int param_5 = 0; + l_1 = 0; + const int x_1007 = param_5; + param_5 = 0; + param_5 = x_1007; + h_1 = 9; + const int x_1008[10] = stack; + const int tint_symbol_16[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_16; + stack = x_1008; + const float2 x_556 = float2(float3(1.0f, 2.0f, 3.0f).y, float3(1.0f, 2.0f, 3.0f).y); + const int x_1009 = param_5; + param_5 = 0; + param_5 = x_1009; + top = -1; + const int x_1010 = p; + p = 0; + p = x_1010; + const int x_93 = top; + const float2 x_557 = float2(float3(1.0f, 2.0f, 3.0f).x, float3(1.0f, 2.0f, 3.0f).x); + const int x_1011 = p; + p = 0; + p = x_1011; + const int x_94 = (x_93 + asint(1u)); + const int x_1012 = top; + top = 0; + top = x_1012; + const float2 x_558 = float2(x_556.y, x_557.y); + const int x_1013 = param_4; + param_4 = 0; + param_4 = x_1013; + top = x_94; + const int x_1014 = h_1; + h_1 = 0; + h_1 = x_1014; + const float3 x_559 = float3(x_557.y, x_557.x, x_557.x); + const int x_1015 = param_4; + param_4 = 0; + param_4 = x_1015; + const int x_95 = l_1; + const QuicksortObject x_1016 = obj; + const int tint_symbol_17[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_18 = {tint_symbol_17}; + obj = tint_symbol_18; + obj = x_1016; + const float3 x_560 = float3(x_559.y, x_559.x, x_557.x); + const int x_96_save = x_94; + const int x_1017[10] = stack; + const int tint_symbol_19[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_19; + stack = x_1017; + const float3 x_561 = float3(x_556.y, x_556.y, x_556.y); + const int x_1018 = l_1; + l_1 = 0; + l_1 = 0; + stack[x_96_save] = x_95; + const int x_1019 = param_5; + param_5 = 0; + param_5 = x_1019; + const int x_97 = top; + const int x_1020 = param_4; + param_4 = 0; + param_4 = x_1020; + const float3 x_562 = float3(float3(1.0f, 2.0f, 3.0f).z, x_558.y, float3(1.0f, 2.0f, 3.0f).y); + const int x_1021 = stack[x_96_save]; + stack[x_96_save] = 0; + stack[x_96_save] = x_1021; + const int x_98 = (x_97 + 1); + const int x_1022 = stack[x_96_save]; + stack[x_96_save] = 0; + stack[x_96_save] = x_1022; + const float3 x_563 = float3(x_559.x, x_559.z, x_556.y); + top = x_98; + const int x_1023 = param_4; + param_4 = 0; + param_4 = x_1023; + const int x_99 = h_1; + const int x_1024 = param_4; + param_4 = 0; + param_4 = x_1024; + const float3 x_564 = float3(x_558.x, x_561.x, x_558.y); + const int x_1025 = l_1; + l_1 = 0; + l_1 = x_1025; + const int x_100_save = x_98; + const int x_1026 = param_5; + param_5 = 0; + param_5 = x_1026; + const float2 x_565 = float2(x_564.z, x_564.z); + const int x_1027 = p; + p = 0; + p = x_1027; + stack[x_100_save] = x_99; + while (true) { + const float3 x_566 = float3(x_563.x, x_563.x, x_563.x); + const int x_1028 = h_1; + h_1 = 0; + h_1 = x_1028; + const int x_1029[10] = stack; + const int tint_symbol_20[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_20; + stack = x_1029; + const int x_106 = top; + const int x_1030[10] = stack; + const int tint_symbol_21[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_21; + stack = x_1030; + const float2 x_567 = float2(x_558.x, x_564.z); + const int x_1031 = param_4; + param_4 = 0; + param_4 = x_1031; + if ((x_106 >= asint(0u))) { + } else { + break; + } + const QuicksortObject x_1032 = obj; + const int tint_symbol_22[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_23 = {tint_symbol_22}; + obj = tint_symbol_23; + obj = x_1032; + const float3 x_568 = float3(x_559.y, x_559.x, x_563.y); + const int x_1033 = param_4; + param_4 = 0; + param_4 = x_1033; + const int x_108 = top; + const float3 x_569 = float3(x_565.x, x_567.y, x_565.x); + const int x_1034 = h_1; + h_1 = 0; + h_1 = x_1034; + const float2 x_570 = float2(x_556.x, x_556.x); + const int x_1035 = p; + p = 0; + p = x_1035; + top = (x_108 - asint(1u)); + const int x_1036 = p; + p = 0; + p = x_1036; + const int x_110_save = x_108; + const int x_1037 = stack[x_96_save]; + stack[x_96_save] = 0; + stack[x_96_save] = x_1037; + const int x_111 = stack[x_110_save]; + const int x_1038[10] = stack; + const int tint_symbol_24[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_24; + stack = x_1038; + const float3 x_571 = float3(x_559.y, x_559.x, x_564.y); + const int x_1039 = l_1; + l_1 = 0; + l_1 = x_1039; + h_1 = x_111; + const int x_1040[10] = stack; + const int tint_symbol_25[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_25; + stack = x_1040; + const float2 x_572 = float2(x_562.y, x_561.y); + const int x_1041 = p; + p = 0; + p = x_1041; + const int x_112 = top; + const int x_1042 = param_4; + param_4 = 0; + param_4 = x_1042; + const int x_1043 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1043; + const float2 x_573 = float2(float3(1.0f, 2.0f, 3.0f).y, float3(1.0f, 2.0f, 3.0f).z); + top = (x_112 - 1); + const int x_1044 = param_5; + param_5 = 0; + param_5 = x_1044; + const float3 x_574 = float3(x_570.y, x_565.x, x_570.y); + const int x_1045 = h_1; + h_1 = 0; + h_1 = x_1045; + const int x_114_save = x_112; + const float2 x_575 = float2(x_564.y, x_564.z); + const int x_1046 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1046; + const int x_115 = stack[x_114_save]; + const int x_1047 = p; + p = 0; + p = x_1047; + const float3 x_576 = float3(x_573.y, x_573.y, x_565.x); + const int x_1048 = param_5; + param_5 = 0; + param_5 = x_1048; + l_1 = x_115; + const int x_1049 = top; + top = 0; + top = x_1049; + const int x_118 = l_1; + param_4 = x_118; + const int x_1050 = stack[x_110_save]; + stack[x_110_save] = 0; + stack[x_110_save] = x_1050; + const float2 x_577 = float2(x_569.y, x_569.z); + const int x_120 = h_1; + const float2 x_578 = float2(x_558.x, float3(1.0f, 2.0f, 3.0f).y); + param_5 = x_120; + const int x_1051 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1051; + const int x_121 = performPartition_i1_i1_(param_4, param_5); + const float2 x_579 = float2(x_567.x, x_568.x); + const int x_1052 = param_5; + param_5 = 0; + param_5 = x_1052; + p = x_121; + const int x_1053 = param_4; + param_4 = 0; + param_4 = x_1053; + const int x_122 = p; + const int x_1054 = h_1; + h_1 = 0; + h_1 = x_1054; + const float2 x_580 = float2(x_568.y, x_568.y); + const int x_1055 = l_1; + l_1 = 0; + l_1 = x_1055; + const int x_1056 = h_1; + h_1 = 0; + h_1 = x_1056; + const int x_124 = l_1; + const int x_1057 = stack[x_110_save]; + stack[x_110_save] = 0; + stack[x_110_save] = x_1057; + const int x_1058 = h_1; + h_1 = 0; + h_1 = x_1058; + const float2 x_582 = float2(x_567.y, x_573.x); + const int x_1059 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1059; + if (((x_122 - asint(1u)) > x_124)) { + const int x_1060 = param_4; + param_4 = 0; + param_4 = x_1060; + const int x_128 = top; + const float2 x_583 = float2(x_571.y, x_556.y); + const int x_1061 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1061; + const int x_1062[10] = stack; + const int tint_symbol_26[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_26; + stack = x_1062; + const float2 x_584 = float2(x_569.z, x_569.y); + const float3 x_585 = float3(x_580.y, x_577.x, x_577.x); + const int x_130 = l_1; + const int x_1063 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1063; + const float2 x_586 = float2(x_564.x, x_585.x); + const int x_1064 = param_5; + param_5 = 0; + param_5 = x_1064; + const int x_131_save = (1 + x_128); + const int x_1065 = stack[x_110_save]; + stack[x_110_save] = 0; + stack[x_110_save] = x_1065; + const float3 x_587 = float3(x_566.y, x_566.y, x_563.x); + const int x_1066 = param_5; + param_5 = 0; + param_5 = x_1066; + stack[x_131_save] = x_130; + const int x_132 = top; + const int x_1067 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1067; + const float2 x_588 = float2(x_575.y, x_575.x); + const int x_1068 = stack[x_131_save]; + stack[x_131_save] = 0; + stack[x_131_save] = x_1068; + const int x_133 = asint((1u + asuint(x_132))); + const int x_1069 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1069; + const float3 x_589 = float3(x_576.z, x_588.y, x_576.z); + const int x_1070 = h_1; + h_1 = 0; + h_1 = x_1070; + top = x_133; + const int x_1071[10] = stack; + const int tint_symbol_27[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_27; + stack = x_1071; + const int x_134 = p; + const float2 x_590 = float2(x_576.x, x_573.y); + const int x_1072 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1072; + const int x_136_save = x_133; + const int x_1073 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1073; + stack[x_136_save] = (x_134 - asint(1u)); + const int x_1074 = stack[x_96_save]; + stack[x_96_save] = 0; + stack[x_96_save] = x_1074; + const float2 x_591 = float2(x_569.z, x_569.y); + const int x_1075 = stack[x_136_save]; + stack[x_136_save] = 0; + stack[x_136_save] = x_1075; + } + const int x_1076 = stack[x_96_save]; + stack[x_96_save] = 0; + stack[x_96_save] = x_1076; + const float2 x_592 = float2(float3(1.0f, 2.0f, 3.0f).x, float3(1.0f, 2.0f, 3.0f).y); + const QuicksortObject x_1077 = obj; + const int tint_symbol_28[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_29 = {tint_symbol_28}; + obj = tint_symbol_29; + obj = x_1077; + const int x_137 = p; + const int x_1078 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1078; + const float3 x_593 = float3(x_571.z, x_556.x, x_556.y); + const int x_1079 = p; + p = 0; + p = x_1079; + const float3 x_594 = float3(x_563.z, x_563.x, x_575.x); + const int x_1080 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1080; + const int x_139 = h_1; + const int x_1081 = top; + top = 0; + top = x_1081; + const float3 x_595 = float3(x_560.z, x_568.x, x_560.x); + const int x_1082 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1082; + const int x_1083 = p; + p = 0; + p = x_1083; + if ((asint((1u + asuint(x_137))) < x_139)) { + const int x_1084 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1084; + const float2 x_596 = float2(x_592.y, x_582.x); + const int x_1085 = l_1; + l_1 = 0; + l_1 = x_1085; + const int x_143 = top; + const int x_1086 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1086; + const float3 x_597 = float3(x_562.y, x_560.y, x_560.y); + const int x_144 = (x_143 + 1); + const int x_1087 = param_5; + param_5 = 0; + param_5 = x_1087; + top = x_144; + const int x_1088 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1088; + const int x_145 = p; + const int x_1089 = param_5; + param_5 = 0; + param_5 = x_1089; + const float3 x_599 = float3(x_560.z, x_560.x, x_568.x); + const int x_1090 = p; + p = 0; + p = x_1090; + const float3 x_600 = float3(x_556.x, x_580.x, x_580.x); + const int x_1091 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1091; + const int x_147_save = x_144; + const int x_1092 = stack[x_110_save]; + stack[x_110_save] = 0; + stack[x_110_save] = x_1092; + const float2 x_601 = float2(x_563.x, x_563.y); + stack[x_147_save] = asint((1u + asuint(x_145))); + const int x_1093[10] = stack; + const int tint_symbol_30[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_30; + stack = x_1093; + const int x_148 = top; + const int x_1094 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1094; + const float2 x_602 = float2(x_565.y, x_599.y); + const int x_1095[10] = stack; + const int tint_symbol_31[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + stack = tint_symbol_31; + stack = x_1095; + const int x_149 = (x_148 + asint(1u)); + const int x_1096 = stack[x_147_save]; + stack[x_147_save] = 0; + stack[x_147_save] = x_1096; + top = x_149; + const int x_1097 = param_4; + param_4 = 0; + param_4 = x_1097; + const int x_150 = h_1; + const int x_1098 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1098; + const int x_1099 = stack[x_96_save]; + stack[x_96_save] = 0; + stack[x_96_save] = x_1099; + stack[x_149] = x_150; + const int x_1100 = stack[x_114_save]; + stack[x_114_save] = 0; + stack[x_114_save] = x_1100; + const float3 x_603 = float3(x_568.y, x_564.x, x_564.x); + const int x_1101 = l_1; + l_1 = 0; + l_1 = x_1101; + } + const int x_1102 = stack[x_100_save]; + stack[x_100_save] = 0; + stack[x_100_save] = x_1102; + { + const int x_1103 = l_1; + l_1 = 0; + l_1 = x_1103; + const float2 x_604 = float2(x_563.z, x_564.x); + const QuicksortObject x_1104 = obj; + const int tint_symbol_32[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_33 = {tint_symbol_32}; + obj = tint_symbol_33; + obj = x_1104; + } + } + const int x_1105 = h_1; + h_1 = 0; + h_1 = x_1105; + return; +} + +main_out main(main_in tint_in) { + main_out tint_out = (main_out)0; + float3 color = float3(0.0f, 0.0f, 0.0f); + int i_2 = 0; + float2 uv = float2(0.0f, 0.0f); + const float2 x_717 = uv; + uv = float2(0.0f, 0.0f); + uv = x_717; + i_2 = 0; + const QuicksortObject x_721 = obj; + const int tint_symbol_34[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_35 = {tint_symbol_34}; + obj = tint_symbol_35; + obj = x_721; + if (true) { + const QuicksortObject x_722 = obj; + const int tint_symbol_36[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_37 = {tint_symbol_36}; + obj = tint_symbol_37; + obj = x_722; + const float2 x_431 = float2(float3(1.0f, 2.0f, 3.0f).x, float3(1.0f, 2.0f, 3.0f).x); + const int x_158 = i_2; + const float2 x_723 = uv; + uv = float2(0.0f, 0.0f); + uv = x_723; + const float3 x_725 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_725; + const float2 x_432 = float2(x_431.y, x_431.y); + const QuicksortObject x_726 = obj; + const int tint_symbol_38[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_39 = {tint_symbol_38}; + obj = tint_symbol_39; + obj = x_726; + } + const QuicksortObject x_756 = obj; + const int tint_symbol_40[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_41 = {tint_symbol_40}; + obj = tint_symbol_41; + obj = x_756; + const float2 x_446 = float2(float2(0.0f, 0.0f).x, float2(0.0f, 0.0f).x); + const int x_757 = i_2; + i_2 = 0; + i_2 = x_757; + quicksort_(); + const QuicksortObject x_758 = obj; + const int tint_symbol_42[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_43 = {tint_symbol_42}; + obj = tint_symbol_43; + obj = x_758; + const float4 x_184 = tint_in.gl_FragCoord; + const float2 x_759 = uv; + uv = float2(0.0f, 0.0f); + uv = x_759; + const float2 x_447 = float2(float2(0.0f, 0.0f).y, float2(0.0f, 0.0f).y); + const float2 x_760 = uv; + uv = float2(0.0f, 0.0f); + uv = x_760; + const float2 x_185 = float2(x_184.x, x_184.y); + const float3 x_448 = float3(x_185.y, x_446.y, x_446.y); + const QuicksortObject x_761 = obj; + const int tint_symbol_44[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_45 = {tint_symbol_44}; + obj = tint_symbol_45; + obj = x_761; + const float2 x_762 = uv; + uv = float2(0.0f, 0.0f); + uv = x_762; + const float2 x_191 = x_188.resolution; + const QuicksortObject x_763 = obj; + const int tint_symbol_46[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_47 = {tint_symbol_46}; + obj = tint_symbol_47; + obj = x_763; + const float3 x_449 = float3(x_184.y, float3(1.0f, 2.0f, 3.0f).z, x_184.w); + const float3 x_764 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_764; + const float2 x_192 = (x_185 / x_191); + const QuicksortObject x_765 = obj; + const int tint_symbol_48[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_49 = {tint_symbol_48}; + obj = tint_symbol_49; + obj = x_765; + const float2 x_450 = float2(x_447.x, x_185.y); + const float3 x_766 = color; + color = float3(0.0f, 0.0f, 0.0f); + const float3 x_767 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_767; + color = x_766; + uv = x_192; + color = float3(1.0f, 2.0f, 3.0f); + const float3 x_768 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_768; + const float3 x_451 = float3(x_185.x, x_185.y, x_446.y); + const QuicksortObject x_769 = obj; + const int tint_symbol_50[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_51 = {tint_symbol_50}; + obj = tint_symbol_51; + obj = x_769; + const int x_770 = obj.numbers[0u]; + obj.numbers[0u] = 0; + obj.numbers[0u] = x_770; + const int x_201 = obj.numbers[0u]; + const QuicksortObject x_771 = obj; + const int tint_symbol_52[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_53 = {tint_symbol_52}; + obj = tint_symbol_53; + obj = x_771; + const int x_772 = obj.numbers[0u]; + obj.numbers[0u] = 0; + obj.numbers[0u] = x_772; + const float x_206 = color.x; + const float x_773 = color.x; + color.x = 0.0f; + color.x = x_773; + const float2 x_452 = float2(float3(1.0f, 2.0f, 3.0f).z, float3(1.0f, 2.0f, 3.0f).y); + const int x_774 = i_2; + i_2 = 0; + i_2 = x_774; + const QuicksortObject x_775 = obj; + const int tint_symbol_54[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_55 = {tint_symbol_54}; + obj = tint_symbol_55; + obj = x_775; + const float3 x_453 = float3(x_451.x, x_450.x, x_450.y); + color.x = (x_206 + float(x_201)); + const float2 x_776 = uv; + uv = float2(0.0f, 0.0f); + uv = x_776; + const float2 x_777 = uv; + uv = float2(0.0f, 0.0f); + uv = x_777; + const float2 x_454 = float2(x_184.y, x_184.y); + const float x_210 = uv.x; + const float2 x_455 = float2(x_192.y, x_192.x); + const float x_778 = uv.x; + uv.x = 0.0f; + uv.x = x_778; + const QuicksortObject x_779 = obj; + const int tint_symbol_56[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_57 = {tint_symbol_56}; + obj = tint_symbol_57; + obj = x_779; + if ((x_210 > 0.25f)) { + const int x_780 = i_2; + i_2 = 0; + i_2 = x_780; + const int x_781 = obj.numbers[0u]; + obj.numbers[0u] = 0; + obj.numbers[0u] = x_781; + const float3 x_456 = float3(float2(0.0f, 0.0f).y, x_448.y, x_448.y); + const float x_782 = uv.x; + uv.x = 0.0f; + uv.x = x_782; + const int x_216 = obj.numbers[1]; + const QuicksortObject x_783 = obj; + const int tint_symbol_58[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_59 = {tint_symbol_58}; + obj = tint_symbol_59; + obj = x_783; + const float2 x_457 = float2(x_454.x, x_454.x); + const float2 x_784 = uv; + uv = float2(0.0f, 0.0f); + uv = x_784; + const QuicksortObject x_785 = obj; + const int tint_symbol_60[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_61 = {tint_symbol_60}; + obj = tint_symbol_61; + obj = x_785; + const float2 x_458 = float2(float3(1.0f, 2.0f, 3.0f).z, float2(0.0f, 0.0f).y); + const int x_786 = i_2; + i_2 = 0; + i_2 = x_786; + const float x_219 = color[0]; + const float x_787 = color[0]; + color[0] = 0.0f; + color[0] = x_787; + const float3 x_788 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_788; + const float3 x_789 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_789; + const float3 x_459 = float3(x_454.y, x_454.y, x_447.y); + const float x_790 = color[0]; + color[0] = 0.0f; + color[0] = x_790; + color.x = (float(x_216) + x_219); + const int x_791 = obj.numbers[0u]; + obj.numbers[0u] = 0; + obj.numbers[0u] = x_791; + } + const float x_792 = uv.x; + uv.x = 0.0f; + uv.x = x_792; + const float x_793 = uv.x; + uv.x = 0.0f; + uv.x = x_793; + const float x_223 = uv.x; + const float x_794 = uv.x; + uv.x = 0.0f; + uv.x = x_794; + const float3 x_460 = float3(x_453.z, x_453.y, x_453.y); + const float2 x_795 = uv; + uv = float2(0.0f, 0.0f); + uv = x_795; + const float x_796 = uv.x; + uv.x = 0.0f; + uv.x = x_796; + const float2 x_461 = float2(float2(0.0f, 0.0f).y, float2(0.0f, 0.0f).y); + const float x_797 = uv.x; + uv.x = 0.0f; + uv.x = x_797; + if ((x_223 > 0.5f)) { + const float x_798 = uv.x; + uv.x = 0.0f; + uv.x = x_798; + const float2 x_462 = float2(x_446.x, x_446.x); + const float x_799 = color.x; + color.x = 0.0f; + color.x = x_799; + const float x_800 = color.x; + color.x = 0.0f; + color.x = x_800; + const float3 x_463 = float3(x_453.x, x_453.z, x_461.y); + const float x_801 = color.x; + color.x = 0.0f; + color.x = x_801; + const int x_230 = obj.numbers[2u]; + const float x_802 = uv.x; + uv.x = 0.0f; + uv.x = x_802; + const float x_803 = color.x; + color.x = 0.0f; + color.x = x_803; + const int x_804 = obj.numbers[2u]; + obj.numbers[2u] = 0; + obj.numbers[2u] = x_804; + const float2 x_464 = float2(x_450.y, x_191.x); + const float x_805 = color.y; + color.y = 0.0f; + color.y = x_805; + const float x_234 = color.y; + const int x_806 = obj.numbers[2u]; + obj.numbers[2u] = 0; + obj.numbers[2u] = x_806; + const float2 x_465 = float2(x_463.x, x_185.x); + const float x_807 = color.x; + color.x = 0.0f; + color.x = x_807; + const int x_808 = i_2; + i_2 = 0; + i_2 = x_808; + const float2 x_466 = float2(x_455.y, float2(0.0f, 0.0f).y); + const int x_809 = i_2; + i_2 = 0; + i_2 = x_809; + color.y = (float(x_230) + x_234); + const float x_810 = uv.x; + uv.x = 0.0f; + uv.x = x_810; + } + const int x_811 = i_2; + i_2 = 0; + i_2 = x_811; + const float2 x_467 = float2(x_191.x, x_191.x); + const float x_812 = uv.x; + uv.x = 0.0f; + uv.x = x_812; + const float x_238 = uv[0]; + const float3 x_813 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_813; + const float x_814 = color.x; + color.x = 0.0f; + color.x = x_814; + if ((x_238 > 0.75f)) { + const float x_815 = color.x; + color.x = 0.0f; + color.x = x_815; + const int x_245 = obj.numbers[3]; + const float x_816 = color.x; + color.x = 0.0f; + color.x = x_816; + const QuicksortObject x_817 = obj; + const int tint_symbol_62[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_63 = {tint_symbol_62}; + obj = tint_symbol_63; + obj = x_817; + const float3 x_468 = float3(x_467.x, x_467.x, x_467.x); + const float x_818 = uv[0]; + uv[0] = 0.0f; + uv[0] = x_818; + const float x_819 = uv.x; + uv.x = 0.0f; + uv.x = x_819; + const float x_249 = color.z; + const float3 x_820 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_820; + const float3 x_469 = float3(x_467.x, x_191.y, x_467.y); + const float x_821 = color.z; + color.z = 0.0f; + color.z = x_821; + const int x_822 = obj.numbers[0u]; + obj.numbers[0u] = 0; + obj.numbers[0u] = x_822; + const float2 x_470 = float2(float2(0.0f, 0.0f).x, float2(0.0f, 0.0f).y); + const float x_823 = color.z; + color.z = 0.0f; + color.z = x_823; + color.z = (x_249 + float(x_245)); + const float2 x_824 = uv; + uv = float2(0.0f, 0.0f); + uv = x_824; + const float2 x_471 = float2(x_470.y, x_470.y); + } + const float x_825 = uv[0]; + uv[0] = 0.0f; + uv[0] = x_825; + const float3 x_472 = float3(x_454.x, x_454.y, x_454.y); + const int x_254 = obj.numbers[4]; + const float x_826 = uv[0]; + uv[0] = 0.0f; + uv[0] = x_826; + const float3 x_827 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_827; + const float3 x_473 = float3(x_446.y, x_453.x, x_453.x); + const int x_828 = obj.numbers[4]; + obj.numbers[4] = 0; + obj.numbers[4] = x_828; + const float2 x_474 = float2(x_191.x, x_184.z); + const float x_829 = uv.x; + uv.x = 0.0f; + uv.x = x_829; + const float x_257 = color.y; + const float x_830 = color.y; + color.y = 0.0f; + color.y = x_830; + const float2 x_475 = float2(x_467.x, x_450.x); + const float x_831 = uv.x; + uv.x = 0.0f; + uv.x = x_831; + const float x_832 = color.x; + color.x = 0.0f; + color.x = x_832; + const float2 x_476 = float2(x_451.z, x_460.y); + color.y = (x_257 + float(x_254)); + const float3 x_477 = float3(float2(0.0f, 0.0f).x, x_472.x, float2(0.0f, 0.0f).y); + const float x_833 = uv.x; + uv.x = 0.0f; + uv.x = x_833; + const float x_834 = color.x; + color.x = 0.0f; + color.x = x_834; + const float2 x_478 = float2(x_472.x, x_472.y); + const float x_835 = uv.y; + uv.y = 0.0f; + uv.y = x_835; + const float x_261 = uv.y; + const int x_836 = i_2; + i_2 = 0; + i_2 = x_836; + const float3 x_479 = float3(float2(0.0f, 0.0f).y, x_454.y, float2(0.0f, 0.0f).x); + const int x_837 = obj.numbers[0u]; + obj.numbers[0u] = 0; + obj.numbers[0u] = x_837; + const float x_838 = color.y; + color.y = 0.0f; + color.y = x_838; + const float3 x_480 = float3(x_446.x, x_446.x, float2(0.0f, 0.0f).y); + const float x_839 = uv.x; + uv.x = 0.0f; + uv.x = x_839; + if ((x_261 > 0.25f)) { + const float2 x_481 = float2(x_447.x, x_480.z); + const float3 x_840 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_840; + const int x_267 = obj.numbers[5u]; + const float x_841 = color.x; + color.x = 0.0f; + color.x = x_841; + const int x_842 = i_2; + i_2 = 0; + i_2 = x_842; + const int x_843 = i_2; + i_2 = 0; + i_2 = x_843; + const float x_270 = color.x; + const float x_844 = uv[0]; + uv[0] = 0.0f; + uv[0] = x_844; + const float3 x_482 = float3(x_455.x, x_475.y, x_455.y); + const QuicksortObject x_845 = obj; + const int tint_symbol_64[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_65 = {tint_symbol_64}; + obj = tint_symbol_65; + obj = x_845; + const float x_846 = uv.y; + uv.y = 0.0f; + uv.y = x_846; + const int x_847 = i_2; + i_2 = 0; + i_2 = x_847; + const float3 x_483 = float3(x_184.w, x_184.w, x_192.x); + const float x_848 = uv.x; + uv.x = 0.0f; + uv.x = x_848; + color.x = (float(x_267) + x_270); + const float3 x_484 = float3(x_454.y, x_450.x, x_454.y); + const float x_849 = uv.x; + uv.x = 0.0f; + uv.x = x_849; + } + const float x_850 = color.x; + color.x = 0.0f; + color.x = x_850; + const float3 x_485 = float3(x_467.x, x_450.y, x_450.x); + const float x_851 = uv.y; + uv.y = 0.0f; + uv.y = x_851; + const int x_852 = obj.numbers[4]; + obj.numbers[4] = 0; + obj.numbers[4] = x_852; + const float x_274 = uv.y; + const int x_853 = obj.numbers[0u]; + obj.numbers[0u] = 0; + obj.numbers[0u] = x_853; + if ((x_274 > 0.5f)) { + const float x_854 = uv.x; + uv.x = 0.0f; + uv.x = x_854; + const float2 x_486 = float2(x_480.y, x_455.y); + const float x_855 = color.y; + color.y = 0.0f; + color.y = x_855; + const float2 x_487 = float2(x_449.z, x_449.y); + const float x_856 = uv.y; + uv.y = 0.0f; + uv.y = x_856; + const int x_280 = obj.numbers[6u]; + const float x_857 = uv.y; + uv.y = 0.0f; + uv.y = x_857; + const int x_858 = i_2; + i_2 = 0; + i_2 = x_858; + const int x_859 = obj.numbers[4]; + obj.numbers[4] = 0; + obj.numbers[4] = x_859; + const float2 x_488 = float2(x_473.z, x_473.y); + const float x_283 = color.y; + const float2 x_860 = uv; + uv = float2(0.0f, 0.0f); + uv = x_860; + const float x_861 = color.x; + color.x = 0.0f; + color.x = x_861; + const float2 x_489 = float2(x_475.y, x_475.x); + const int x_862 = obj.numbers[6u]; + obj.numbers[6u] = 0; + obj.numbers[6u] = x_862; + const int x_863 = obj.numbers[6u]; + obj.numbers[6u] = 0; + obj.numbers[6u] = x_863; + const float2 x_490 = float2(x_480.z, x_480.z); + const QuicksortObject x_864 = obj; + const int tint_symbol_66[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_67 = {tint_symbol_66}; + obj = tint_symbol_67; + obj = x_864; + color.y = (float(x_280) + x_283); + const float x_865 = color.x; + color.x = 0.0f; + color.x = x_865; + const float2 x_491 = float2(float3(1.0f, 2.0f, 3.0f).y, x_454.x); + const float x_866 = color.y; + color.y = 0.0f; + color.y = x_866; + } + const float2 x_492 = float2(x_455.y, x_455.y); + const float x_867 = color.x; + color.x = 0.0f; + color.x = x_867; + const float x_287 = uv.y; + const QuicksortObject x_868 = obj; + const int tint_symbol_68[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_69 = {tint_symbol_68}; + obj = tint_symbol_69; + obj = x_868; + const float2 x_493 = float2(x_475.x, x_475.y); + const float x_869 = uv[0]; + uv[0] = 0.0f; + uv[0] = x_869; + const float x_870 = color.y; + color.y = 0.0f; + color.y = x_870; + const float3 x_494 = float3(x_191.x, x_191.y, x_191.y); + const int x_871 = obj.numbers[4]; + obj.numbers[4] = 0; + obj.numbers[4] = x_871; + if ((x_287 > 0.75f)) { + const float3 x_872 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_872; + const float x_873 = color.x; + color.x = 0.0f; + color.x = x_873; + const float3 x_495 = float3(x_192.y, x_192.x, x_192.y); + const float3 x_874 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_874; + const int x_293 = obj.numbers[7]; + const float x_875 = uv.x; + uv.x = 0.0f; + uv.x = x_875; + const float3 x_496 = float3(x_475.x, x_467.y, x_467.x); + const float x_876 = color.y; + color.y = 0.0f; + color.y = x_876; + const float2 x_497 = float2(x_477.x, x_461.y); + const int x_877 = obj.numbers[0u]; + obj.numbers[0u] = 0; + obj.numbers[0u] = x_877; + const float x_878 = color.y; + color.y = 0.0f; + color.y = x_878; + const float3 x_498 = float3(x_478.x, x_478.y, x_478.x); + const float x_879 = color.x; + color.x = 0.0f; + color.x = x_879; + const float x_296 = color.z; + const float x_880 = uv.y; + uv.y = 0.0f; + uv.y = x_880; + const float2 x_499 = float2(x_184.x, x_184.y); + const float x_881 = uv.x; + uv.x = 0.0f; + uv.x = x_881; + const float x_882 = uv.y; + uv.y = 0.0f; + uv.y = x_882; + const float x_883 = uv.y; + uv.y = 0.0f; + uv.y = x_883; + const float3 x_500 = float3(x_499.y, x_499.y, x_494.z); + const float x_884 = color.z; + color.z = 0.0f; + color.z = x_884; + color.z = (float(x_293) + x_296); + const float x_885 = color.y; + color.y = 0.0f; + color.y = x_885; + const float2 x_501 = float2(x_453.x, x_453.z); + const float x_886 = color.x; + color.x = 0.0f; + color.x = x_886; + } + const int x_887 = i_2; + i_2 = 0; + i_2 = x_887; + const float2 x_502 = float2(x_451.y, x_192.y); + const float2 x_888 = uv; + uv = float2(0.0f, 0.0f); + uv = x_888; + const int x_301 = obj.numbers[8]; + const int x_889 = i_2; + i_2 = 0; + i_2 = x_889; + const float2 x_503 = float2(x_185.x, x_451.z); + const int x_890 = obj.numbers[8]; + obj.numbers[8] = 0; + obj.numbers[8] = x_890; + const float x_891 = color.y; + color.y = 0.0f; + color.y = x_891; + const float2 x_504 = float2(x_453.y, float2(0.0f, 0.0f).x); + const float x_892 = color.x; + color.x = 0.0f; + color.x = x_892; + const float3 x_505 = float3(x_504.x, x_504.y, x_504.x); + const float x_893 = color.z; + color.z = 0.0f; + color.z = x_893; + const float x_304 = color.z; + const float x_894 = color.x; + color.x = 0.0f; + color.x = x_894; + const float2 x_506 = float2(x_493.x, x_492.x); + const int x_895 = obj.numbers[4]; + obj.numbers[4] = 0; + obj.numbers[4] = x_895; + const float x_896 = uv.y; + uv.y = 0.0f; + uv.y = x_896; + const float2 x_507 = float2(x_461.x, x_447.x); + const float x_897 = color.y; + color.y = 0.0f; + color.y = x_897; + color.z = (x_304 + float(x_301)); + const float2 x_898 = uv; + uv = float2(0.0f, 0.0f); + uv = x_898; + const float x_899 = uv.x; + uv.x = 0.0f; + uv.x = x_899; + const float3 x_508 = float3(x_461.y, x_461.x, x_506.y); + const float x_900 = uv.x; + uv.x = 0.0f; + uv.x = x_900; + const float x_308 = uv.x; + const float x_901 = color.y; + color.y = 0.0f; + color.y = x_901; + const float3 x_509 = float3(x_503.y, x_503.x, x_448.z); + const float x_902 = uv.y; + uv.y = 0.0f; + uv.y = x_902; + const float x_310 = uv.y; + const float x_903 = uv.y; + uv.y = 0.0f; + uv.y = x_903; + const float x_904 = color.z; + color.z = 0.0f; + color.z = x_904; + const float3 x_510 = float3(float3(1.0f, 2.0f, 3.0f).y, x_485.y, x_485.z); + const float x_905 = color.z; + color.z = 0.0f; + color.z = x_905; + const int x_906 = i_2; + i_2 = 0; + i_2 = x_906; + const float2 x_511 = float2(x_485.z, x_485.y); + const float3 x_907 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_907; + const float x_908 = uv.y; + uv.y = 0.0f; + uv.y = x_908; + const float3 x_512 = float3(x_455.y, x_455.y, x_455.y); + const int x_909 = obj.numbers[4]; + obj.numbers[4] = 0; + obj.numbers[4] = x_909; + if ((abs((x_308 - x_310)) < 0.25f)) { + const float x_910 = uv.x; + uv.x = 0.0f; + uv.x = x_910; + const QuicksortObject x_911 = obj; + const int tint_symbol_70[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_71 = {tint_symbol_70}; + obj = tint_symbol_71; + obj = x_911; + const float3 x_513 = float3(x_505.z, x_505.x, x_448.x); + const int x_912 = obj.numbers[8]; + obj.numbers[8] = 0; + obj.numbers[8] = x_912; + const int x_317 = obj.numbers[9u]; + const float3 x_514 = float3(x_474.y, x_474.y, x_474.y); + const float x_913 = uv.y; + uv.y = 0.0f; + uv.y = x_913; + const float x_320 = color.x; + const float x_914 = uv.y; + uv.y = 0.0f; + uv.y = x_914; + const float2 x_515 = float2(x_502.x, x_502.y); + const float x_915 = color.x; + color.x = 0.0f; + color.x = x_915; + const float3 x_916 = color; + color = float3(0.0f, 0.0f, 0.0f); + color = x_916; + const float2 x_516 = float2(x_452.x, x_452.x); + const float2 x_917 = uv; + uv = float2(0.0f, 0.0f); + uv = x_917; + const float x_918 = uv.x; + uv.x = 0.0f; + uv.x = x_918; + const float3 x_517 = float3(float2(0.0f, 0.0f).x, float2(0.0f, 0.0f).x, float2(0.0f, 0.0f).y); + color.x = (float(x_317) + x_320); + const float x_919 = color.x; + color.x = 0.0f; + color.x = x_919; + const float3 x_518 = float3(x_480.y, x_508.x, x_480.x); + const float x_920 = color.x; + color.x = 0.0f; + color.x = x_920; + } + const float x_921 = uv.y; + uv.y = 0.0f; + uv.y = x_921; + const float3 x_325 = color; + const float x_922 = uv[0]; + uv[0] = 0.0f; + uv[0] = x_922; + const float3 x_519 = float3(x_447.x, x_446.x, x_446.y); + const float3 x_326 = normalize(x_325); + const float x_923 = uv.x; + uv.x = 0.0f; + uv.x = x_923; + const QuicksortObject x_924 = obj; + const int tint_symbol_72[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_73 = {tint_symbol_72}; + obj = tint_symbol_73; + obj = x_924; + const QuicksortObject x_925 = obj; + const int tint_symbol_74[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_75 = {tint_symbol_74}; + obj = tint_symbol_75; + obj = x_925; + const float x_926 = color.y; + color.y = 0.0f; + color.y = x_926; + const float2 x_520 = float2(x_506.y, x_519.y); + const float x_927 = color.y; + color.y = 0.0f; + color.y = x_927; + const float4 x_330 = float4(x_326.x, x_326.y, x_326.z, 1.0f); + const float x_928 = uv.y; + uv.y = 0.0f; + uv.y = x_928; + const float3 x_521 = float3(float3(1.0f, 2.0f, 3.0f).y, float3(1.0f, 2.0f, 3.0f).y, x_520.y); + const float x_929 = uv.x; + uv.x = 0.0f; + uv.x = x_929; + tint_out.x_GLF_color = x_330; + const QuicksortObject x_930 = obj; + const int tint_symbol_76[10] = {0, 0, 0, 0, 0, 0, 0, 0, 0, 0}; + const QuicksortObject tint_symbol_77 = {tint_symbol_76}; + obj = tint_symbol_77; + obj = x_930; + const float3 x_522 = float3(x_330.w, x_330.y, x_493.x); + const float x_931 = color.x; + color.x = 0.0f; + color.x = x_931; + return tint_out; +} + diff --git a/test/ptr_ref/load/local/ptr_function.wgsl.expected.hlsl b/test/ptr_ref/load/local/ptr_function.wgsl.expected.hlsl index b3db42f63d..fd572806fc 100644 --- a/test/ptr_ref/load/local/ptr_function.wgsl.expected.hlsl +++ b/test/ptr_ref/load/local/ptr_function.wgsl.expected.hlsl @@ -1 +1,7 @@ -SKIP: Failed to generate: error: pointers not supported in HLSL +[numthreads(1, 1, 1)] +void main() { + int i = 123; + const int use = (i + 1); + return; +} + diff --git a/test/ptr_ref/load/local/ptr_private.wgsl.expected.hlsl b/test/ptr_ref/load/local/ptr_private.wgsl.expected.hlsl index b3db42f63d..c28a1f3b1f 100644 --- a/test/ptr_ref/load/local/ptr_private.wgsl.expected.hlsl +++ b/test/ptr_ref/load/local/ptr_private.wgsl.expected.hlsl @@ -1 +1,8 @@ -SKIP: Failed to generate: error: pointers not supported in HLSL +static int i = 123; + +[numthreads(1, 1, 1)] +void main() { + const int use = (i + 1); + return; +} + diff --git a/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl index b3db42f63d..dad67de454 100644 --- a/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl +++ b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl @@ -1 +1,12 @@ -SKIP: Failed to generate: error: pointers not supported in HLSL +struct S { + int a; +}; + +ConstantBuffer v : register(b0, space0); + +[numthreads(1, 1, 1)] +void main() { + const int use = (v.a + 1); + return; +} + diff --git a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.hlsl b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.hlsl index b3db42f63d..f282b82855 100644 --- a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.hlsl +++ b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.hlsl @@ -1 +1,9 @@ -SKIP: Failed to generate: error: pointers not supported in HLSL +groupshared int i; + +[numthreads(1, 1, 1)] +void main() { + i = 123; + const int use = (i + 1); + return; +} + diff --git a/test/samples/compute_boids.wgsl.expected.hlsl b/test/samples/compute_boids.wgsl.expected.hlsl index e29d25bac0..080bcf3ae7 100644 --- a/test/samples/compute_boids.wgsl.expected.hlsl +++ b/test/samples/compute_boids.wgsl.expected.hlsl @@ -31,8 +31,8 @@ tint_symbol_2 vert_main(tint_symbol_1 tint_symbol) { const float2 a_particlePos = tint_symbol.a_particlePos; const float2 a_particleVel = tint_symbol.a_particleVel; const float2 a_pos = tint_symbol.a_pos; - float angle = -( atan2(a_particleVel.x, a_particleVel.y)); - float2 pos = float2(((a_pos.x * cos(angle)) - (a_pos.y * sin(angle))), ((a_pos.x * sin(angle)) + (a_pos.y * cos(angle)))); + float angle = -(atan2(a_particleVel.x, a_particleVel.y)); + float2 pos = float2(((a_pos.x * cos(angle)) - (a_pos.y * sin(angle))), ((a_pos.x * sin(angle)) + (a_pos.y * cos(angle)))); const tint_symbol_2 tint_symbol_8 = {float4((pos + a_particlePos), 0.0f, 1.0f)}; return tint_symbol_8; } @@ -60,33 +60,31 @@ void comp_main(tint_symbol_5 tint_symbol_4) { float2 vel = float2(0.0f, 0.0f); { uint i = 0u; - { - bool tint_hlsl_is_first_1 = true; - for(;;) { - if (!tint_hlsl_is_first_1) { + while (true) { + if (!((i < 5u))) { + break; + } + if ((i == index)) { + { i = (i + 1u); } - tint_hlsl_is_first_1 = false; - - if (!((i < 5u))) { - break; - } - if ((i == index)) { - continue; - } - pos = asfloat(particlesA.Load2((16u * i))).xy; - vel = asfloat(particlesA.Load2(((16u * i) + 8u))).xy; - if (( distance(pos, vPos) < params.rule1Distance)) { - cMass = (cMass + pos); - cMassCount = (cMassCount + 1); - } - if (( distance(pos, vPos) < params.rule2Distance)) { - colVel = (colVel - (pos - vPos)); - } - if (( distance(pos, vPos) < params.rule3Distance)) { - cVel = (cVel + vel); - cVelCount = (cVelCount + 1); - } + continue; + } + pos = asfloat(particlesA.Load2((16u * i))).xy; + vel = asfloat(particlesA.Load2(((16u * i) + 8u))).xy; + if ((distance(pos, vPos) < params.rule1Distance)) { + cMass = (cMass + pos); + cMassCount = (cMassCount + 1); + } + if ((distance(pos, vPos) < params.rule2Distance)) { + colVel = (colVel - (pos - vPos)); + } + if ((distance(pos, vPos) < params.rule3Distance)) { + cVel = (cVel + vel); + cVelCount = (cVelCount + 1); + } + { + i = (i + 1u); } } } @@ -97,7 +95,7 @@ void comp_main(tint_symbol_5 tint_symbol_4) { cVel = (cVel / float2(float(cVelCount), float(cVelCount))); } vVel = (((vVel + (cMass * params.rule1Scale)) + (colVel * params.rule2Scale)) + (cVel * params.rule3Scale)); - vVel = ( normalize(vVel) * clamp( length(vVel), 0.0f, 0.100000001f)); + vVel = (normalize(vVel) * clamp(length(vVel), 0.0f, 0.100000001f)); vPos = (vPos + (vVel * params.deltaT)); if ((vPos.x < -1.0f)) { vPos.x = 1.0f;