diff --git a/reference/opt/shaders-msl/comp/torture-loop.comp b/reference/opt/shaders-msl/comp/torture-loop.comp index 5e14eb9f..c9e38513 100644 --- a/reference/opt/shaders-msl/comp/torture-loop.comp +++ b/reference/opt/shaders-msl/comp/torture-loop.comp @@ -47,7 +47,8 @@ kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], devic _96 = _40; float4 _100; uint _101; - for (uint _97 = 0u, _99 = _98; _97 < 16u; _95 = _100, _96++, _97++, _99 = _101) + uint _99; + for (uint _97 = 0u; _97 < 16u; _95 = _100, _96++, _97++, _99 = _101) { _100 = _95; _101 = 0u; diff --git a/reference/opt/shaders-msl/frag/false-loop-init.frag b/reference/opt/shaders-msl/frag/false-loop-init.frag new file mode 100644 index 00000000..c67bb9d3 --- /dev/null +++ b/reference/opt/shaders-msl/frag/false-loop-init.frag @@ -0,0 +1,38 @@ +#include +#include + +using namespace metal; + +constant uint _49 = {}; + +struct main0_in +{ + float4 accum [[user(locn0)]]; +}; + +struct main0_out +{ + float4 result [[color(0)]]; +}; + +fragment main0_out main0(main0_in in [[stage_in]]) +{ + main0_out out = {}; + out.result = float4(0.0); + uint _51; + uint _50; + for (int _48 = 0; _48 < 4; _48 += int(_51), _50 = _51) + { + if (in.accum.y > 10.0) + { + _51 = 40u; + } + else + { + _51 = 30u; + } + out.result += in.accum; + } + return out; +} + diff --git a/reference/opt/shaders/comp/torture-loop.comp b/reference/opt/shaders/comp/torture-loop.comp index 1ba8a739..4e5c402e 100644 --- a/reference/opt/shaders/comp/torture-loop.comp +++ b/reference/opt/shaders/comp/torture-loop.comp @@ -45,7 +45,8 @@ void main() _96 = _40; vec4 _100; uint _101; - for (uint _97 = 0u, _99 = _98; _97 < 16u; _95 = _100, _96++, _97++, _99 = _101) + uint _99; + for (uint _97 = 0u; _97 < 16u; _95 = _100, _96++, _97++, _99 = _101) { _100 = _95; _101 = 0u; diff --git a/reference/opt/shaders/flatten/multi-dimensional.desktop.flatten_dim.frag b/reference/opt/shaders/flatten/multi-dimensional.desktop.flatten_dim.frag index df33ef73..6ccede21 100644 --- a/reference/opt/shaders/flatten/multi-dimensional.desktop.flatten_dim.frag +++ b/reference/opt/shaders/flatten/multi-dimensional.desktop.flatten_dim.frag @@ -13,7 +13,9 @@ void main() vec4 values3[2 * 3 * 1]; int _96; int _97; - for (int _92 = 0, _94 = _93, _95 = _93; _92 < 2; _92++, _94 = _96, _95 = _97) + int _94; + int _95; + for (int _92 = 0; _92 < 2; _92++, _94 = _96, _95 = _97) { _96 = 0; _97 = _95; diff --git a/reference/opt/shaders/frag/false-loop-init.frag b/reference/opt/shaders/frag/false-loop-init.frag new file mode 100644 index 00000000..1db46c1b --- /dev/null +++ b/reference/opt/shaders/frag/false-loop-init.frag @@ -0,0 +1,28 @@ +#version 310 es +precision mediump float; +precision highp int; + +layout(location = 0) out vec4 result; +layout(location = 0) in vec4 accum; + +uint _49; + +void main() +{ + result = vec4(0.0); + uint _51; + uint _50; + for (int _48 = 0; _48 < 4; _48 += int(_51), _50 = _51) + { + if (accum.y > 10.0) + { + _51 = 40u; + } + else + { + _51 = 30u; + } + result += accum; + } +} + diff --git a/reference/shaders-msl/frag/false-loop-init.frag b/reference/shaders-msl/frag/false-loop-init.frag new file mode 100644 index 00000000..e0792474 --- /dev/null +++ b/reference/shaders-msl/frag/false-loop-init.frag @@ -0,0 +1,35 @@ +#include +#include + +using namespace metal; + +struct main0_in +{ + float4 accum [[user(locn0)]]; +}; + +struct main0_out +{ + float4 result [[color(0)]]; +}; + +fragment main0_out main0(main0_in in [[stage_in]]) +{ + main0_out out = {}; + out.result = float4(0.0); + uint j; + for (int i = 0; i < 4; i += int(j)) + { + if (in.accum.y > 10.0) + { + j = 40u; + } + else + { + j = 30u; + } + out.result += in.accum; + } + return out; +} + diff --git a/reference/shaders/frag/false-loop-init.frag b/reference/shaders/frag/false-loop-init.frag new file mode 100644 index 00000000..b0ed5577 --- /dev/null +++ b/reference/shaders/frag/false-loop-init.frag @@ -0,0 +1,25 @@ +#version 310 es +precision mediump float; +precision highp int; + +layout(location = 0) out vec4 result; +layout(location = 0) in vec4 accum; + +void main() +{ + result = vec4(0.0); + mediump uint j; + for (mediump int i = 0; i < 4; i += int(j)) + { + if (accum.y > 10.0) + { + j = 40u; + } + else + { + j = 30u; + } + result += accum; + } +} + diff --git a/shaders-msl/frag/false-loop-init.frag b/shaders-msl/frag/false-loop-init.frag new file mode 100644 index 00000000..7ce5b52b --- /dev/null +++ b/shaders-msl/frag/false-loop-init.frag @@ -0,0 +1,19 @@ +#version 310 es +precision mediump float; + +layout(location = 0) in vec4 accum; +layout(location = 0) out vec4 result; + +void main() +{ + result = vec4(0.0); + uint j; + for (int i = 0; i < 4; i += int(j)) + { + if (accum.y > 10.0) + j = 40u; + else + j = 30u; + result += accum; + } +} diff --git a/shaders/frag/false-loop-init.frag b/shaders/frag/false-loop-init.frag new file mode 100644 index 00000000..7ce5b52b --- /dev/null +++ b/shaders/frag/false-loop-init.frag @@ -0,0 +1,19 @@ +#version 310 es +precision mediump float; + +layout(location = 0) in vec4 accum; +layout(location = 0) out vec4 result; + +void main() +{ + result = vec4(0.0); + uint j; + for (int i = 0; i < 4; i += int(j)) + { + if (accum.y > 10.0) + j = 40u; + else + j = 30u; + result += accum; + } +} diff --git a/spirv_cross.cpp b/spirv_cross.cpp index c2f9396e..26983a43 100644 --- a/spirv_cross.cpp +++ b/spirv_cross.cpp @@ -3503,6 +3503,10 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry) assert(header); auto &header_block = this->get(header); + auto &blocks = handler.accessed_variables_to_block[loop_variable.first]; + + // If a loop variable is not used before the loop, it's probably not a loop variable. + bool has_accessed_variable = blocks.count(header) != 0; // Now, there are two conditions we need to meet for the variable to be a loop variable. // 1. The dominating block must have a branch-free path to the loop header, @@ -3513,6 +3517,9 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry) bool static_loop_init = true; while (dominator != header) { + if (blocks.count(dominator) != 0) + has_accessed_variable = true; + auto &succ = cfg.get_succeeding_edges(dominator); if (succ.size() != 1) { @@ -3530,12 +3537,11 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry) dominator = succ.front(); } - if (!static_loop_init) + if (!static_loop_init || !has_accessed_variable) continue; // The second condition we need to meet is that no access after the loop // merge can occur. Walk the CFG to see if we find anything. - auto &blocks = handler.accessed_variables_to_block[loop_variable.first]; seen_blocks.clear(); cfg.walk_from(seen_blocks, header_block.merge_block, [&](uint32_t walk_block) { diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index eebcd2db..426cee06 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -7180,10 +7180,18 @@ string CompilerGLSL::variable_decl(const SPIRVariable &variable) auto res = join(to_qualifiers_glsl(variable.self), variable_decl(type, to_name(variable.self), variable.self)); - if (variable.loop_variable) - res += join(" = ", to_expression(variable.static_expression)); + if (variable.loop_variable && variable.static_expression) + { + uint32_t expr = variable.static_expression; + if (ids[expr].get_type() != TypeUndef) + res += join(" = ", to_expression(variable.static_expression)); + } else if (variable.initializer) - res += join(" = ", to_expression(variable.initializer)); + { + uint32_t expr = variable.initializer; + if (ids[expr].get_type() != TypeUndef) + res += join(" = ", to_expression(variable.initializer)); + } return res; } @@ -8022,11 +8030,23 @@ string CompilerGLSL::emit_for_loop_initializers(const SPIRBlock &block) // We can only declare for loop initializers if all variables are of same type. // If we cannot do this, declare individual variables before the loop header. - if (block.loop_variables.size() == 1) + // We might have a loop variable candidate which was not assigned to for some reason. + uint32_t missing_initializers = 0; + for (auto &variable : block.loop_variables) + { + uint32_t expr = get(variable).static_expression; + + // Sometimes loop variables are initialized with OpUndef, but we can just declare + // a plain variable without initializer in this case. + if (expr == 0 || ids[expr].get_type() == TypeUndef) + missing_initializers++; + } + + if (block.loop_variables.size() == 1 && missing_initializers == 0) { return variable_decl(get(block.loop_variables.front())); } - else if (!same_types) + else if (!same_types || missing_initializers == uint32_t(block.loop_variables.size())) { for (auto &loop_var : block.loop_variables) statement(variable_decl(get(loop_var)), ";"); @@ -8034,20 +8054,32 @@ string CompilerGLSL::emit_for_loop_initializers(const SPIRBlock &block) } else { - auto &var = get(block.loop_variables.front()); - auto &type = get(var.basetype); - - // Don't remap the type here as we have multiple names, - // doesn't make sense to remap types for loop variables anyways. - // It is assumed here that all relevant qualifiers are equal for all loop variables. - string expr = join(to_qualifiers_glsl(var.self), type_to_glsl(type), " "); + // We have a mix of loop variables, either ones with a clear initializer, or ones without. + // Separate the two streams. + string expr; for (auto &loop_var : block.loop_variables) { - auto &v = get(loop_var); - expr += join(to_name(loop_var), " = ", to_expression(v.static_expression)); - if (&loop_var != &block.loop_variables.back()) - expr += ", "; + uint32_t static_expr = get(loop_var).static_expression; + if (static_expr == 0 || ids[static_expr].get_type() == TypeUndef) + { + statement(variable_decl(get(loop_var)), ";"); + } + else + { + if (expr.empty()) + { + // For loop initializers are of the form (block.loop_variables.front()); + auto &type = get(var.basetype); + expr = join(to_qualifiers_glsl(var.self), type_to_glsl(type), " "); + } + else + expr += ", "; + + auto &v = get(loop_var); + expr += join(to_name(loop_var), " = ", to_expression(v.static_expression)); + } } return expr; } @@ -8058,11 +8090,21 @@ bool CompilerGLSL::for_loop_initializers_are_same_type(const SPIRBlock &block) if (block.loop_variables.size() <= 1) return true; - uint32_t expected = get(block.loop_variables[0]).basetype; - uint64_t expected_flags = get_decoration_mask(block.loop_variables[0]); + uint32_t expected = 0; + uint64_t expected_flags = 0; for (auto &var : block.loop_variables) { - if (expected != get(var).basetype) + // Don't care about uninitialized variables as they will not be part of the initializers. + uint32_t expr = get(var).static_expression; + if (expr == 0 || ids[expr].get_type() == TypeUndef) + continue; + + if (expected == 0) + { + expected = get(var).basetype; + expected_flags = get_decoration_mask(var); + } + else if (expected != get(var).basetype) return false; // Precision flags and things like that must also match.