Merge pull request #408 from KhronosGroup/fix-407
Check if a loop variable candidate is actually used.
This commit is contained in:
commit
3fc2561734
@ -47,7 +47,8 @@ kernel void main0(uint3 gl_GlobalInvocationID [[thread_position_in_grid]], devic
|
|||||||
_96 = _40;
|
_96 = _40;
|
||||||
float4 _100;
|
float4 _100;
|
||||||
uint _101;
|
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;
|
_100 = _95;
|
||||||
_101 = 0u;
|
_101 = 0u;
|
||||||
|
38
reference/opt/shaders-msl/frag/false-loop-init.frag
Normal file
38
reference/opt/shaders-msl/frag/false-loop-init.frag
Normal file
@ -0,0 +1,38 @@
|
|||||||
|
#include <metal_stdlib>
|
||||||
|
#include <simd/simd.h>
|
||||||
|
|
||||||
|
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;
|
||||||
|
}
|
||||||
|
|
@ -45,7 +45,8 @@ void main()
|
|||||||
_96 = _40;
|
_96 = _40;
|
||||||
vec4 _100;
|
vec4 _100;
|
||||||
uint _101;
|
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;
|
_100 = _95;
|
||||||
_101 = 0u;
|
_101 = 0u;
|
||||||
|
@ -13,7 +13,9 @@ void main()
|
|||||||
vec4 values3[2 * 3 * 1];
|
vec4 values3[2 * 3 * 1];
|
||||||
int _96;
|
int _96;
|
||||||
int _97;
|
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;
|
_96 = 0;
|
||||||
_97 = _95;
|
_97 = _95;
|
||||||
|
28
reference/opt/shaders/frag/false-loop-init.frag
Normal file
28
reference/opt/shaders/frag/false-loop-init.frag
Normal file
@ -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;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
35
reference/shaders-msl/frag/false-loop-init.frag
Normal file
35
reference/shaders-msl/frag/false-loop-init.frag
Normal file
@ -0,0 +1,35 @@
|
|||||||
|
#include <metal_stdlib>
|
||||||
|
#include <simd/simd.h>
|
||||||
|
|
||||||
|
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;
|
||||||
|
}
|
||||||
|
|
25
reference/shaders/frag/false-loop-init.frag
Normal file
25
reference/shaders/frag/false-loop-init.frag
Normal file
@ -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;
|
||||||
|
}
|
||||||
|
}
|
||||||
|
|
19
shaders-msl/frag/false-loop-init.frag
Normal file
19
shaders-msl/frag/false-loop-init.frag
Normal file
@ -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;
|
||||||
|
}
|
||||||
|
}
|
19
shaders/frag/false-loop-init.frag
Normal file
19
shaders/frag/false-loop-init.frag
Normal file
@ -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;
|
||||||
|
}
|
||||||
|
}
|
@ -3503,6 +3503,10 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry)
|
|||||||
|
|
||||||
assert(header);
|
assert(header);
|
||||||
auto &header_block = this->get<SPIRBlock>(header);
|
auto &header_block = this->get<SPIRBlock>(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.
|
// 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,
|
// 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;
|
bool static_loop_init = true;
|
||||||
while (dominator != header)
|
while (dominator != header)
|
||||||
{
|
{
|
||||||
|
if (blocks.count(dominator) != 0)
|
||||||
|
has_accessed_variable = true;
|
||||||
|
|
||||||
auto &succ = cfg.get_succeeding_edges(dominator);
|
auto &succ = cfg.get_succeeding_edges(dominator);
|
||||||
if (succ.size() != 1)
|
if (succ.size() != 1)
|
||||||
{
|
{
|
||||||
@ -3530,12 +3537,11 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry)
|
|||||||
dominator = succ.front();
|
dominator = succ.front();
|
||||||
}
|
}
|
||||||
|
|
||||||
if (!static_loop_init)
|
if (!static_loop_init || !has_accessed_variable)
|
||||||
continue;
|
continue;
|
||||||
|
|
||||||
// The second condition we need to meet is that no access after the loop
|
// 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.
|
// 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();
|
seen_blocks.clear();
|
||||||
cfg.walk_from(seen_blocks, header_block.merge_block, [&](uint32_t walk_block) {
|
cfg.walk_from(seen_blocks, header_block.merge_block, [&](uint32_t walk_block) {
|
||||||
|
@ -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));
|
auto res = join(to_qualifiers_glsl(variable.self), variable_decl(type, to_name(variable.self), variable.self));
|
||||||
|
|
||||||
if (variable.loop_variable)
|
if (variable.loop_variable && variable.static_expression)
|
||||||
res += join(" = ", to_expression(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)
|
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;
|
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.
|
// 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 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<SPIRVariable>(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<SPIRVariable>(block.loop_variables.front()));
|
return variable_decl(get<SPIRVariable>(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)
|
for (auto &loop_var : block.loop_variables)
|
||||||
statement(variable_decl(get<SPIRVariable>(loop_var)), ";");
|
statement(variable_decl(get<SPIRVariable>(loop_var)), ";");
|
||||||
@ -8034,20 +8054,32 @@ string CompilerGLSL::emit_for_loop_initializers(const SPIRBlock &block)
|
|||||||
}
|
}
|
||||||
else
|
else
|
||||||
{
|
{
|
||||||
auto &var = get<SPIRVariable>(block.loop_variables.front());
|
// We have a mix of loop variables, either ones with a clear initializer, or ones without.
|
||||||
auto &type = get<SPIRType>(var.basetype);
|
// Separate the two streams.
|
||||||
|
string expr;
|
||||||
// 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), " ");
|
|
||||||
|
|
||||||
for (auto &loop_var : block.loop_variables)
|
for (auto &loop_var : block.loop_variables)
|
||||||
{
|
{
|
||||||
auto &v = get<SPIRVariable>(loop_var);
|
uint32_t static_expr = get<SPIRVariable>(loop_var).static_expression;
|
||||||
expr += join(to_name(loop_var), " = ", to_expression(v.static_expression));
|
if (static_expr == 0 || ids[static_expr].get_type() == TypeUndef)
|
||||||
if (&loop_var != &block.loop_variables.back())
|
{
|
||||||
expr += ", ";
|
statement(variable_decl(get<SPIRVariable>(loop_var)), ";");
|
||||||
|
}
|
||||||
|
else
|
||||||
|
{
|
||||||
|
if (expr.empty())
|
||||||
|
{
|
||||||
|
// For loop initializers are of the form <type id = value, id = value, id = value, etc ...
|
||||||
|
auto &var = get<SPIRVariable>(block.loop_variables.front());
|
||||||
|
auto &type = get<SPIRType>(var.basetype);
|
||||||
|
expr = join(to_qualifiers_glsl(var.self), type_to_glsl(type), " ");
|
||||||
|
}
|
||||||
|
else
|
||||||
|
expr += ", ";
|
||||||
|
|
||||||
|
auto &v = get<SPIRVariable>(loop_var);
|
||||||
|
expr += join(to_name(loop_var), " = ", to_expression(v.static_expression));
|
||||||
|
}
|
||||||
}
|
}
|
||||||
return expr;
|
return expr;
|
||||||
}
|
}
|
||||||
@ -8058,11 +8090,21 @@ bool CompilerGLSL::for_loop_initializers_are_same_type(const SPIRBlock &block)
|
|||||||
if (block.loop_variables.size() <= 1)
|
if (block.loop_variables.size() <= 1)
|
||||||
return true;
|
return true;
|
||||||
|
|
||||||
uint32_t expected = get<SPIRVariable>(block.loop_variables[0]).basetype;
|
uint32_t expected = 0;
|
||||||
uint64_t expected_flags = get_decoration_mask(block.loop_variables[0]);
|
uint64_t expected_flags = 0;
|
||||||
for (auto &var : block.loop_variables)
|
for (auto &var : block.loop_variables)
|
||||||
{
|
{
|
||||||
if (expected != get<SPIRVariable>(var).basetype)
|
// Don't care about uninitialized variables as they will not be part of the initializers.
|
||||||
|
uint32_t expr = get<SPIRVariable>(var).static_expression;
|
||||||
|
if (expr == 0 || ids[expr].get_type() == TypeUndef)
|
||||||
|
continue;
|
||||||
|
|
||||||
|
if (expected == 0)
|
||||||
|
{
|
||||||
|
expected = get<SPIRVariable>(var).basetype;
|
||||||
|
expected_flags = get_decoration_mask(var);
|
||||||
|
}
|
||||||
|
else if (expected != get<SPIRVariable>(var).basetype)
|
||||||
return false;
|
return false;
|
||||||
|
|
||||||
// Precision flags and things like that must also match.
|
// Precision flags and things like that must also match.
|
||||||
|
Loading…
Reference in New Issue
Block a user