AuroraMiddleware/SPIRV-Cross
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Deduce constant LUTs from read-write variables.

Browse Source
This commit is contained in:
Hans-Kristian Arntzen 2018-07-05 13:25:57 +02:00
parent b5ed706860
commit d29f48ef06
20 changed files with 227 additions and 56 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
reference/opt/shaders-hlsl/frag/array-lut-no-loop-variable.frag
View File

@ -15,11 +15,10 @@ struct SPIRV_Cross_Output
void frag_main()
{
float lut[5] = _17;
for (int _46 = 0; _46 < 4; )
{
int _33 = _46 + 1;
FragColor += lut[_33].xxxx;
FragColor += _17[_33].xxxx;
_46 = _33;
continue;
}

3
reference/opt/shaders-msl/frag/array-lut-no-loop-variable.frag
View File

@ -29,11 +29,10 @@ void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
fragment main0_out main0()
{
main0_out out = {};
float lut[5] = {1.0, 2.0, 3.0, 4.0, 5.0};
for (int _46 = 0; _46 < 4; )
{
int _33 = _46 + 1;
out.FragColor += float4(lut[_33]);
out.FragColor += float4(_17[_33]);
_46 = _33;
continue;
}

6
reference/opt/shaders-msl/frag/constant-array.frag
View File

@ -44,10 +44,8 @@ void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
fragment main0_out main0(main0_in in [[stage_in]])
{
main0_out out = {};
float4 indexable[3] = {float4(1.0), float4(2.0), float4(3.0)};
float4 indexable_1[2][2] = {{float4(1.0), float4(2.0)}, {float4(8.0), float4(10.0)}};
Foobar indexable_2[2] = {{10.0, 40.0}, {90.0, 70.0}};
out.FragColor = ((indexable[in.index] + (indexable_1[in.index][in.index + 1])) + float4(30.0)) + float4(indexable_2[in.index].a + indexable_2[in.index].b);
Foobar indexable[2] = {{10.0, 40.0}, {90.0, 70.0}};
out.FragColor = ((_37[in.index] + (_55[in.index][in.index + 1])) + float4(30.0)) + float4(indexable[in.index].a + indexable[in.index].b);
return out;
}

5
reference/opt/shaders/frag/array-lut-no-loop-variable.frag
View File

@ -2,15 +2,16 @@
precision mediump float;
precision highp int;
const float _17[5] = float[](1.0, 2.0, 3.0, 4.0, 5.0);
layout(location = 0) out vec4 FragColor;
void main()
{
float lut[5] = float[](1.0, 2.0, 3.0, 4.0, 5.0);
for (int _46 = 0; _46 < 4; )
{
mediump int _33 = _46 + 1;
FragColor += vec4(lut[_33]);
FragColor += vec4(_17[_33]);
_46 = _33;
continue;
}

9
reference/opt/shaders/frag/constant-array.frag
View File

@ -2,6 +2,9 @@
precision mediump float;
precision highp int;
const vec4 _37[3] = vec4[](vec4(1.0), vec4(2.0), vec4(3.0));
const vec4 _55[2][2] = vec4[][](vec4[](vec4(1.0), vec4(2.0)), vec4[](vec4(8.0), vec4(10.0)));
struct Foobar
{
float a;
@ -13,9 +16,7 @@ layout(location = 0) flat in mediump int index;
void main()
{
highp vec4 indexable[3] = vec4[](vec4(1.0), vec4(2.0), vec4(3.0));
highp vec4 indexable_1[2][2] = vec4[][](vec4[](vec4(1.0), vec4(2.0)), vec4[](vec4(8.0), vec4(10.0)));
Foobar indexable_2[2] = Foobar[](Foobar(10.0, 40.0), Foobar(90.0, 70.0));
FragColor = ((indexable[index] + (indexable_1[index][index + 1])) + vec4(30.0)) + vec4(indexable_2[index].a + indexable_2[index].b);
Foobar indexable[2] = Foobar[](Foobar(10.0, 40.0), Foobar(90.0, 70.0));
FragColor = ((_37[index] + (_55[index][index + 1])) + vec4(30.0)) + vec4(indexable[index].a + indexable[index].b);
}

3
reference/shaders-hlsl/frag/array-lut-no-loop-variable.frag
View File

@ -15,8 +15,7 @@ struct SPIRV_Cross_Output
void frag_main()
{
float lut[5] = _17;
for (int i = 0; i < 4; i++, FragColor += lut[i].xxxx)
for (int i = 0; i < 4; i++, FragColor += _17[i].xxxx)
{
}
}

1
reference/shaders-msl/asm/frag/op-constant-null.asm.frag
View File

@ -39,7 +39,6 @@ fragment main0_out main0()
float4 b = float4(0.0);
float2x3 c = float2x3(float3(0.0), float3(0.0));
D d = {float4(0.0), 0.0};
float4 e[4] = {float4(0.0), float4(0.0), float4(0.0), float4(0.0)};
out.FragColor = a;
return out;
}

3
reference/shaders-msl/comp/composite-construct.comp
View File

@ -42,8 +42,7 @@ kernel void main0(device SSBO0& _16 [[buffer(0)]], device SSBO1& _32 [[buffer(1)
float4 _37[2] = { _16.as[gl_GlobalInvocationID.x], _32.bs[gl_GlobalInvocationID.x] };
float4 values[2];
spvArrayCopy(values, _37);
float4 copy_values[2] = {float4(20.0), float4(40.0)};
Composite c = Composite{ values[0], copy_values[1] };
Composite c = Composite{ values[0], _43[1] };
_16.as[0] = values[gl_LocalInvocationIndex];
_32.bs[1] = c.b;
}

3
reference/shaders-msl/frag/array-lut-no-loop-variable.frag
View File

@ -29,8 +29,7 @@ void spvArrayCopyConstant(thread T (&dst)[N], constant T (&src)[N])
fragment main0_out main0()
{
main0_out out = {};
float lut[5] = {1.0, 2.0, 3.0, 4.0, 5.0};
for (int i = 0; i < 4; i++, out.FragColor += float4(lut[i]))
for (int i = 0; i < 4; i++, out.FragColor += float4(_17[i]))
{
}
return out;

8
reference/shaders-msl/frag/constant-array.frag
View File

@ -49,12 +49,10 @@ float4 resolve(thread const Foobar& f)
fragment main0_out main0(main0_in in [[stage_in]])
{
main0_out out = {};
float4 indexable[3] = {float4(1.0), float4(2.0), float4(3.0)};
float4 indexable_1[2][2] = {{float4(1.0), float4(2.0)}, {float4(8.0), float4(10.0)}};
Foobar param = {10.0, 20.0};
Foobar indexable_2[2] = {{10.0, 40.0}, {90.0, 70.0}};
Foobar param_1 = indexable_2[in.index];
out.FragColor = ((indexable[in.index] + (indexable_1[in.index][in.index + 1])) + resolve(param)) + resolve(param_1);
Foobar indexable[2] = {{10.0, 40.0}, {90.0, 70.0}};
Foobar param_1 = indexable[in.index];
out.FragColor = ((_37[in.index] + (_55[in.index][in.index + 1])) + resolve(param)) + resolve(param_1);
return out;
}

3
reference/shaders/asm/frag/op-constant-null.asm.frag
View File

@ -2,6 +2,8 @@
precision mediump float;
precision highp int;
const vec4 _14[4] = vec4[](vec4(0.0), vec4(0.0), vec4(0.0), vec4(0.0));
struct D
{
vec4 a;
@ -16,7 +18,6 @@ void main()
vec4 b = vec4(0.0);
mat2x3 c = mat2x3(vec3(0.0), vec3(0.0));
D d = D(vec4(0.0), 0.0);
vec4 e[4] = vec4[](vec4(0.0), vec4(0.0), vec4(0.0), vec4(0.0));
FragColor = a;
}

7
reference/shaders/comp/composite-construct.comp
View File

@ -1,6 +1,9 @@
#version 310 es
layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
const vec4 _66[2] = vec4[](vec4(10.0), vec4(30.0));
const float _94[2][3] = float[][](float[](1.0, 1.0, 1.0), float[](2.0, 2.0, 2.0));
struct Composite
{
vec4 a[2];
@ -25,13 +28,11 @@ vec4 summe(vec4 values[3][2])
void main()
{
vec4 values[2] = vec4[](_41.as[gl_GlobalInvocationID.x], _55.bs[gl_GlobalInvocationID.x]);
vec4 const_values[2] = vec4[](vec4(10.0), vec4(30.0));
vec4 copy_values[2] = const_values;
vec4 copy_values[2] = _66;
vec4 copy_values2[2] = values;
vec4 param[3][2] = vec4[][](values, copy_values, copy_values2);
_41.as[gl_GlobalInvocationID.x] = summe(param);
Composite c = Composite(values, copy_values);
float arrayofarray[2][3] = float[][](float[](1.0, 1.0, 1.0), float[](2.0, 2.0, 2.0));
float b = 10.0;
float values_scalar[4] = float[](b, b, b, b);
}

5
reference/shaders/frag/array-lut-no-loop-variable.frag
View File

@ -2,12 +2,13 @@
precision mediump float;
precision highp int;
const float _17[5] = float[](1.0, 2.0, 3.0, 4.0, 5.0);
layout(location = 0) out vec4 FragColor;
void main()
{
float lut[5] = float[](1.0, 2.0, 3.0, 4.0, 5.0);
for (mediump int i = 0; i < 4; i++, FragColor += vec4(lut[i]))
for (mediump int i = 0; i < 4; i++, FragColor += vec4(_17[i]))
{
}
}

11
reference/shaders/frag/constant-array.frag
View File

@ -2,6 +2,9 @@
precision mediump float;
precision highp int;
const vec4 _37[3] = vec4[](vec4(1.0), vec4(2.0), vec4(3.0));
const vec4 _55[2][2] = vec4[][](vec4[](vec4(1.0), vec4(2.0)), vec4[](vec4(8.0), vec4(10.0)));
struct Foobar
{
float a;
@ -18,11 +21,9 @@ vec4 resolve(Foobar f)
void main()
{
highp vec4 indexable[3] = vec4[](vec4(1.0), vec4(2.0), vec4(3.0));
highp vec4 indexable_1[2][2] = vec4[][](vec4[](vec4(1.0), vec4(2.0)), vec4[](vec4(8.0), vec4(10.0)));
Foobar param = Foobar(10.0, 20.0);
Foobar indexable_2[2] = Foobar[](Foobar(10.0, 40.0), Foobar(90.0, 70.0));
Foobar param_1 = indexable_2[index];
FragColor = ((indexable[index] + (indexable_1[index][index + 1])) + resolve(param)) + resolve(param_1);
Foobar indexable[2] = Foobar[](Foobar(10.0, 40.0), Foobar(90.0, 70.0));
Foobar param_1 = indexable[index];
FragColor = ((_37[index] + (_55[index][index + 1])) + resolve(param)) + resolve(param_1);
}

3
spirv_common.hpp
View File

@ -1101,6 +1101,9 @@ struct SPIRConstant : IVariant
// If this constant is used as an array length which creates specialization restrictions on some backends.
bool is_used_as_array_length = false;
// If true, this is a LUT, and should always be declared in the outer scope.
bool is_used_as_lut = false;
// For composites which are constant arrays, etc.
std::vector<uint32_t> subconstants;
};

148
spirv_cross.cpp
View File

@ -3918,10 +3918,150 @@ bool Compiler::AnalyzeVariableScopeAccessHandler::handle(spv::Op op, const uint3
return true;
}
void Compiler::analyze_variable_scope(SPIRFunction &entry)
Compiler::StaticExpressionAccessHandler::StaticExpressionAccessHandler(Compiler &compiler_,
uint32_t variable_id_)
: compiler(compiler_), variable_id(variable_id_)
{
AnalyzeVariableScopeAccessHandler handler(*this, entry);
}
bool Compiler::StaticExpressionAccessHandler::follow_function_call(const SPIRFunction &)
{
return false;
}
bool Compiler::StaticExpressionAccessHandler::handle(spv::Op op, const uint32_t *args, uint32_t length)
{
switch (op)
{
case OpStore:
if (length < 2)
return false;
if (args[0] == variable_id)
{
static_expression = args[1];
write_count++;
}
break;
case OpLoad:
if (length < 3)
return false;
if (args[2] == variable_id && static_expression == 0) // Tried to read from variable before it was initialized.
return false;
break;
case OpAccessChain:
case OpInBoundsAccessChain:
if (length < 3)
return false;
if (args[2] == variable_id) // If we try to access chain our candidate variable before we store to it, bail.
return false;
break;
default:
break;
}
return true;
}
void Compiler::find_function_local_luts(SPIRFunction &entry,
const AnalyzeVariableScopeAccessHandler &handler)
{
auto &cfg = *function_cfgs.find(entry.self)->second;
// For each variable which is statically accessed.
for (auto &accessed_var : handler.accessed_variables_to_block)
{
auto &blocks = accessed_var.second;
auto &var = get<SPIRVariable>(accessed_var.first);
auto &type = expression_type(accessed_var.first);
// Only consider function local variables here.
if (var.storage != StorageClassFunction)
continue;
// We cannot be a phi variable.
if (var.phi_variable)
continue;
// Only consider arrays here.
if (type.array.empty())
continue;
// HACK: Do not consider structs. This is a quirk with how types are currently being emitted.
// Structs are emitted after specialization constants and composite constants.
// FIXME: Fix declaration order so declared constants can have struct types.
if (type.basetype == SPIRType::Struct)
continue;
// If the variable has an initializer, make sure it is a constant expression.
uint32_t static_constant_expression = 0;
if (var.initializer)
{
if (ids[var.initializer].get_type() != TypeConstant)
continue;
static_constant_expression = var.initializer;
// There can be no stores to this variable, we have now proved we have a LUT.
if (handler.complete_write_variables_to_block.count(var.self) != 0 ||
handler.partial_write_variables_to_block.count(var.self) != 0)
continue;
}
else
{
// We can have one, and only one write to the variable, and that write needs to be a constant.
// No partial writes allowed.
if (handler.partial_write_variables_to_block.count(var.self) != 0)
continue;
auto itr = handler.complete_write_variables_to_block.find(var.self);
// No writes?
if (itr == end(handler.complete_write_variables_to_block))
continue;
// We write to the variable in more than one block.
auto &write_blocks = itr->second;
if (write_blocks.size() != 1)
continue;
// The write needs to happen in the dominating block.
DominatorBuilder builder(cfg);
for (auto &block : blocks)
builder.add_block(block);
uint32_t dominator = builder.get_dominator();
// The complete write happened in a branch or similar, cannot deduce static expression.
if (write_blocks.count(dominator) == 0)
continue;
// Find the static expression for this variable.
StaticExpressionAccessHandler static_expression_handler(*this, var.self);
traverse_all_reachable_opcodes(get<SPIRBlock>(dominator), static_expression_handler);
// We want one, and exactly one write
if (static_expression_handler.write_count != 1 || static_expression_handler.static_expression == 0)
continue;
// Is it a constant expression?
if (ids[static_expression_handler.static_expression].get_type() != TypeConstant)
continue;
// We found a LUT!
static_constant_expression = static_expression_handler.static_expression;
}
get<SPIRConstant>(static_constant_expression).is_used_as_lut = true;
var.static_expression = static_constant_expression;
var.statically_assigned = true;
var.remapped_variable = true;
}
}
void Compiler::analyze_variable_scope(SPIRFunction &entry, AnalyzeVariableScopeAccessHandler &handler)
{
// First, we map out all variable access within a function.
// Essentially a map of block -> { variables accessed in the basic block }
traverse_all_reachable_opcodes(entry, handler);
@ -4423,7 +4563,9 @@ void Compiler::build_function_control_flow_graphs_and_analyze()
for (auto &f : function_cfgs)
{
auto &func = get<SPIRFunction>(f.first);
analyze_variable_scope(func);
AnalyzeVariableScopeAccessHandler scope_handler(*this, func);
analyze_variable_scope(func, scope_handler);
find_function_local_luts(func, scope_handler);
// Check if we can actually use the loop variables we found in analyze_variable_scope.
// To use multiple initializers, we need the same type and qualifiers.

17
spirv_cross.hpp
View File

@ -676,8 +676,6 @@ protected:
variable_remap_callback(type, var_name, type_name);
}
void analyze_variable_scope(SPIRFunction &function);
void parse();
void parse(const Instruction &i);
@ -903,6 +901,21 @@ protected:
const SPIRBlock *current_block = nullptr;
};
struct StaticExpressionAccessHandler : OpcodeHandler
{
StaticExpressionAccessHandler(Compiler &compiler_, uint32_t variable_id_);
bool follow_function_call(const SPIRFunction &) override;
bool handle(spv::Op op, const uint32_t *args, uint32_t length) override;
Compiler &compiler;
uint32_t variable_id;
uint32_t static_expression = 0;
uint32_t write_count = 0;
};
void analyze_variable_scope(SPIRFunction &function, AnalyzeVariableScopeAccessHandler &handler);
void find_function_local_luts(SPIRFunction &function, const AnalyzeVariableScopeAccessHandler &handler);
void make_constant_null(uint32_t id, uint32_t type);
std::vector<spv::Capability> declared_capabilities;

38
spirv_glsl.cpp
View File

@ -1651,7 +1651,7 @@ void CompilerGLSL::emit_specialization_constant_op(const SPIRConstantOp &constan
statement("const ", variable_decl(type, name), " = ", constant_op_expression(constant), ";");
}
void CompilerGLSL::emit_specialization_constant(const SPIRConstant &constant)
void CompilerGLSL::emit_constant(const SPIRConstant &constant)
{
auto &type = get<SPIRType>(constant.constant_type);
auto name = to_name(constant.self);
@ -2072,25 +2072,26 @@ void CompilerGLSL::emit_resources()
//
// TODO: If we have the fringe case that we create a spec constant which depends on a struct type,
// we'll have to deal with that, but there's currently no known way to express that.
if (options.vulkan_semantics)
for (auto &id : ids)
{
for (auto &id : ids)
if (id.get_type() == TypeConstant)
{
if (id.get_type() == TypeConstant)
{
auto &c = id.get<SPIRConstant>();
if (!c.specialization)
continue;
auto &c = id.get<SPIRConstant>();
emit_specialization_constant(c);
emitted = true;
}
else if (id.get_type() == TypeConstantOp)
bool needs_declaration = (c.specialization && options.vulkan_semantics) ||
c.is_used_as_lut;
if (needs_declaration)
{
emit_specialization_constant_op(id.get<SPIRConstantOp>());
emit_constant(c);
emitted = true;
}
}
else if (options.vulkan_semantics && id.get_type() == TypeConstantOp)
{
emit_specialization_constant_op(id.get<SPIRConstantOp>());
emitted = true;
}
}
if (emitted)
@ -2425,6 +2426,8 @@ string CompilerGLSL::to_expression(uint32_t id)
return builtin_to_glsl(dec.builtin_type, StorageClassGeneric);
else if (c.specialization && options.vulkan_semantics)
return to_name(id);
else if (c.is_used_as_lut)
return to_name(id);
else if (type.basetype == SPIRType::Struct && !backend.can_declare_struct_inline)
return to_name(id);
else if (!type.array.empty() && !backend.can_declare_arrays_inline)
@ -6191,6 +6194,10 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction)
var->static_expression = ops[1];
else if (var && var->loop_variable && !var->loop_variable_enable)
var->static_expression = ops[1];
else if (var && var->remapped_variable)
{
// Skip the write.
}
else if (var && flattened_structs.count(ops[0]))
{
store_flattened_struct(*var, ops[1]);
@ -8935,6 +8942,11 @@ void CompilerGLSL::emit_function(SPIRFunction &func, const Bitset &return_flags)
entry_block.dominated_variables.push_back(var.self);
var.deferred_declaration = true;
}
else if (var.storage == StorageClassFunction && var.remapped_variable && var.static_expression)
{
// No need to declare this variable, it has a static expression.
var.deferred_declaration = false;
}
else if (expression_is_lvalue(v))
{
add_local_variable_name(var.self);

2
spirv_glsl.hpp
View File

@ -374,7 +374,7 @@ protected:
void emit_flattened_io_block(const SPIRVariable &var, const char *qual);
void emit_block_chain(SPIRBlock &block);
void emit_hoisted_temporaries(std::vector<std::pair<uint32_t, uint32_t>> &temporaries);
void emit_specialization_constant(const SPIRConstant &constant);
void emit_constant(const SPIRConstant &constant);
void emit_specialization_constant_op(const SPIRConstantOp &constant);
std::string emit_continue_block(uint32_t continue_block);
bool attempt_emit_loop_header(SPIRBlock &block, SPIRBlock::Method method);

5
spirv_msl.cpp
View File

@ -2110,6 +2110,11 @@ bool CompilerMSL::maybe_emit_array_assignment(uint32_t id_lhs, uint32_t id_rhs)
return false;
auto *var = maybe_get<SPIRVariable>(id_lhs);
// Is this a remapped, static constant? Don't do anything.
if (var->remapped_variable && var->statically_assigned)
return true;
if (ids[id_rhs].get_type() == TypeConstant && var && var->deferred_declaration)
{
// Special case, if we end up declaring a variable when assigning the constant array,