a489ba7fd1
- Replace ostringstream with custom implementation. ~30% performance uplift on vector-shuffle-oom test. Allocations are measurably reduced in Valgrind. - Replace std::vector with SmallVector. Classic malloc optimization, small vectors are backed by inline data. ~ 7-8% gain on vector-shuffle-oom on GCC 8 on Linux. - Use an object pool for IVariant type. We generally allocate a lot of SPIR* objects. We can amortize these allocations neatly by pooling them. - ~15% overall uplift on ./test_shaders.py --iterations 10000 shaders/.
548 lines
15 KiB
C++
548 lines
15 KiB
C++
/*
|
|
* Copyright 2015-2019 Arm Limited
|
|
*
|
|
* Licensed under the Apache License, Version 2.0 (the "License");
|
|
* you may not use this file except in compliance with the License.
|
|
* You may obtain a copy of the License at
|
|
*
|
|
* http://www.apache.org/licenses/LICENSE-2.0
|
|
*
|
|
* Unless required by applicable law or agreed to in writing, software
|
|
* distributed under the License is distributed on an "AS IS" BASIS,
|
|
* WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
|
|
* See the License for the specific language governing permissions and
|
|
* limitations under the License.
|
|
*/
|
|
|
|
#include "spirv_cpp.hpp"
|
|
|
|
using namespace spv;
|
|
using namespace SPIRV_CROSS_NAMESPACE;
|
|
using namespace std;
|
|
|
|
void CompilerCPP::emit_buffer_block(const SPIRVariable &var)
|
|
{
|
|
add_resource_name(var.self);
|
|
|
|
auto &type = get<SPIRType>(var.basetype);
|
|
auto instance_name = to_name(var.self);
|
|
|
|
uint32_t descriptor_set = ir.meta[var.self].decoration.set;
|
|
uint32_t binding = ir.meta[var.self].decoration.binding;
|
|
|
|
emit_block_struct(type);
|
|
auto buffer_name = to_name(type.self);
|
|
|
|
statement("internal::Resource<", buffer_name, type_to_array_glsl(type), "> ", instance_name, "__;");
|
|
statement_no_indent("#define ", instance_name, " __res->", instance_name, "__.get()");
|
|
resource_registrations.push_back(
|
|
join("s.register_resource(", instance_name, "__", ", ", descriptor_set, ", ", binding, ");"));
|
|
statement("");
|
|
}
|
|
|
|
void CompilerCPP::emit_interface_block(const SPIRVariable &var)
|
|
{
|
|
add_resource_name(var.self);
|
|
|
|
auto &type = get<SPIRType>(var.basetype);
|
|
|
|
const char *qual = var.storage == StorageClassInput ? "StageInput" : "StageOutput";
|
|
const char *lowerqual = var.storage == StorageClassInput ? "stage_input" : "stage_output";
|
|
auto instance_name = to_name(var.self);
|
|
uint32_t location = ir.meta[var.self].decoration.location;
|
|
|
|
string buffer_name;
|
|
auto flags = ir.meta[type.self].decoration.decoration_flags;
|
|
if (flags.get(DecorationBlock))
|
|
{
|
|
emit_block_struct(type);
|
|
buffer_name = to_name(type.self);
|
|
}
|
|
else
|
|
buffer_name = type_to_glsl(type);
|
|
|
|
statement("internal::", qual, "<", buffer_name, type_to_array_glsl(type), "> ", instance_name, "__;");
|
|
statement_no_indent("#define ", instance_name, " __res->", instance_name, "__.get()");
|
|
resource_registrations.push_back(join("s.register_", lowerqual, "(", instance_name, "__", ", ", location, ");"));
|
|
statement("");
|
|
}
|
|
|
|
void CompilerCPP::emit_shared(const SPIRVariable &var)
|
|
{
|
|
add_resource_name(var.self);
|
|
|
|
auto instance_name = to_name(var.self);
|
|
statement(CompilerGLSL::variable_decl(var), ";");
|
|
statement_no_indent("#define ", instance_name, " __res->", instance_name);
|
|
}
|
|
|
|
void CompilerCPP::emit_uniform(const SPIRVariable &var)
|
|
{
|
|
add_resource_name(var.self);
|
|
|
|
auto &type = get<SPIRType>(var.basetype);
|
|
auto instance_name = to_name(var.self);
|
|
|
|
uint32_t descriptor_set = ir.meta[var.self].decoration.set;
|
|
uint32_t binding = ir.meta[var.self].decoration.binding;
|
|
uint32_t location = ir.meta[var.self].decoration.location;
|
|
|
|
string type_name = type_to_glsl(type);
|
|
remap_variable_type_name(type, instance_name, type_name);
|
|
|
|
if (type.basetype == SPIRType::Image || type.basetype == SPIRType::SampledImage ||
|
|
type.basetype == SPIRType::AtomicCounter)
|
|
{
|
|
statement("internal::Resource<", type_name, type_to_array_glsl(type), "> ", instance_name, "__;");
|
|
statement_no_indent("#define ", instance_name, " __res->", instance_name, "__.get()");
|
|
resource_registrations.push_back(
|
|
join("s.register_resource(", instance_name, "__", ", ", descriptor_set, ", ", binding, ");"));
|
|
}
|
|
else
|
|
{
|
|
statement("internal::UniformConstant<", type_name, type_to_array_glsl(type), "> ", instance_name, "__;");
|
|
statement_no_indent("#define ", instance_name, " __res->", instance_name, "__.get()");
|
|
resource_registrations.push_back(
|
|
join("s.register_uniform_constant(", instance_name, "__", ", ", location, ");"));
|
|
}
|
|
|
|
statement("");
|
|
}
|
|
|
|
void CompilerCPP::emit_push_constant_block(const SPIRVariable &var)
|
|
{
|
|
add_resource_name(var.self);
|
|
|
|
auto &type = get<SPIRType>(var.basetype);
|
|
auto &flags = ir.meta[var.self].decoration.decoration_flags;
|
|
if (flags.get(DecorationBinding) || flags.get(DecorationDescriptorSet))
|
|
SPIRV_CROSS_THROW("Push constant blocks cannot be compiled to GLSL with Binding or Set syntax. "
|
|
"Remap to location with reflection API first or disable these decorations.");
|
|
|
|
emit_block_struct(type);
|
|
auto buffer_name = to_name(type.self);
|
|
auto instance_name = to_name(var.self);
|
|
|
|
statement("internal::PushConstant<", buffer_name, type_to_array_glsl(type), "> ", instance_name, ";");
|
|
statement_no_indent("#define ", instance_name, " __res->", instance_name, ".get()");
|
|
resource_registrations.push_back(join("s.register_push_constant(", instance_name, "__", ");"));
|
|
statement("");
|
|
}
|
|
|
|
void CompilerCPP::emit_block_struct(SPIRType &type)
|
|
{
|
|
// C++ can't do interface blocks, so we fake it by emitting a separate struct.
|
|
// However, these structs are not allowed to alias anything, so remove it before
|
|
// emitting the struct.
|
|
//
|
|
// The type we have here needs to be resolved to the non-pointer type so we can remove aliases.
|
|
auto &self = get<SPIRType>(type.self);
|
|
self.type_alias = 0;
|
|
emit_struct(self);
|
|
}
|
|
|
|
void CompilerCPP::emit_resources()
|
|
{
|
|
for (auto &id : ir.ids)
|
|
{
|
|
if (id.get_type() == TypeConstant)
|
|
{
|
|
auto &c = id.get<SPIRConstant>();
|
|
|
|
bool needs_declaration = c.specialization || c.is_used_as_lut;
|
|
|
|
if (needs_declaration)
|
|
{
|
|
if (!options.vulkan_semantics && c.specialization)
|
|
{
|
|
c.specialization_constant_macro_name =
|
|
constant_value_macro_name(get_decoration(c.self, DecorationSpecId));
|
|
}
|
|
emit_constant(c);
|
|
}
|
|
}
|
|
else if (id.get_type() == TypeConstantOp)
|
|
{
|
|
emit_specialization_constant_op(id.get<SPIRConstantOp>());
|
|
}
|
|
}
|
|
|
|
// Output all basic struct types which are not Block or BufferBlock as these are declared inplace
|
|
// when such variables are instantiated.
|
|
for (auto &id : ir.ids)
|
|
{
|
|
if (id.get_type() == TypeType)
|
|
{
|
|
auto &type = id.get<SPIRType>();
|
|
if (type.basetype == SPIRType::Struct && type.array.empty() && !type.pointer &&
|
|
(!ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock) &&
|
|
!ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock)))
|
|
{
|
|
emit_struct(type);
|
|
}
|
|
}
|
|
}
|
|
|
|
statement("struct Resources : ", resource_type);
|
|
begin_scope();
|
|
|
|
// Output UBOs and SSBOs
|
|
for (auto &id : ir.ids)
|
|
{
|
|
if (id.get_type() == TypeVariable)
|
|
{
|
|
auto &var = id.get<SPIRVariable>();
|
|
auto &type = get<SPIRType>(var.basetype);
|
|
|
|
if (var.storage != StorageClassFunction && type.pointer && type.storage == StorageClassUniform &&
|
|
!is_hidden_variable(var) &&
|
|
(ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock) ||
|
|
ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock)))
|
|
{
|
|
emit_buffer_block(var);
|
|
}
|
|
}
|
|
}
|
|
|
|
// Output push constant blocks
|
|
for (auto &id : ir.ids)
|
|
{
|
|
if (id.get_type() == TypeVariable)
|
|
{
|
|
auto &var = id.get<SPIRVariable>();
|
|
auto &type = get<SPIRType>(var.basetype);
|
|
if (!is_hidden_variable(var) && var.storage != StorageClassFunction && type.pointer &&
|
|
type.storage == StorageClassPushConstant)
|
|
{
|
|
emit_push_constant_block(var);
|
|
}
|
|
}
|
|
}
|
|
|
|
// Output in/out interfaces.
|
|
for (auto &id : ir.ids)
|
|
{
|
|
if (id.get_type() == TypeVariable)
|
|
{
|
|
auto &var = id.get<SPIRVariable>();
|
|
auto &type = get<SPIRType>(var.basetype);
|
|
|
|
if (var.storage != StorageClassFunction && !is_hidden_variable(var) && type.pointer &&
|
|
(var.storage == StorageClassInput || var.storage == StorageClassOutput) &&
|
|
interface_variable_exists_in_entry_point(var.self))
|
|
{
|
|
emit_interface_block(var);
|
|
}
|
|
}
|
|
}
|
|
|
|
// Output Uniform Constants (values, samplers, images, etc).
|
|
for (auto &id : ir.ids)
|
|
{
|
|
if (id.get_type() == TypeVariable)
|
|
{
|
|
auto &var = id.get<SPIRVariable>();
|
|
auto &type = get<SPIRType>(var.basetype);
|
|
|
|
if (var.storage != StorageClassFunction && !is_hidden_variable(var) && type.pointer &&
|
|
(type.storage == StorageClassUniformConstant || type.storage == StorageClassAtomicCounter))
|
|
{
|
|
emit_uniform(var);
|
|
}
|
|
}
|
|
}
|
|
|
|
// Global variables.
|
|
bool emitted = false;
|
|
for (auto global : global_variables)
|
|
{
|
|
auto &var = get<SPIRVariable>(global);
|
|
if (var.storage == StorageClassWorkgroup)
|
|
{
|
|
emit_shared(var);
|
|
emitted = true;
|
|
}
|
|
}
|
|
|
|
if (emitted)
|
|
statement("");
|
|
|
|
declare_undefined_values();
|
|
|
|
statement("inline void init(spirv_cross_shader& s)");
|
|
begin_scope();
|
|
statement(resource_type, "::init(s);");
|
|
for (auto ® : resource_registrations)
|
|
statement(reg);
|
|
end_scope();
|
|
resource_registrations.clear();
|
|
|
|
end_scope_decl();
|
|
|
|
statement("");
|
|
statement("Resources* __res;");
|
|
if (get_entry_point().model == ExecutionModelGLCompute)
|
|
statement("ComputePrivateResources __priv_res;");
|
|
statement("");
|
|
|
|
// Emit regular globals which are allocated per invocation.
|
|
emitted = false;
|
|
for (auto global : global_variables)
|
|
{
|
|
auto &var = get<SPIRVariable>(global);
|
|
if (var.storage == StorageClassPrivate)
|
|
{
|
|
if (var.storage == StorageClassWorkgroup)
|
|
emit_shared(var);
|
|
else
|
|
statement(CompilerGLSL::variable_decl(var), ";");
|
|
emitted = true;
|
|
}
|
|
}
|
|
|
|
if (emitted)
|
|
statement("");
|
|
}
|
|
|
|
string CompilerCPP::compile()
|
|
{
|
|
// Do not deal with ES-isms like precision, older extensions and such.
|
|
options.es = false;
|
|
options.version = 450;
|
|
backend.float_literal_suffix = true;
|
|
backend.double_literal_suffix = false;
|
|
backend.long_long_literal_suffix = true;
|
|
backend.uint32_t_literal_suffix = true;
|
|
backend.basic_int_type = "int32_t";
|
|
backend.basic_uint_type = "uint32_t";
|
|
backend.swizzle_is_function = true;
|
|
backend.shared_is_implied = true;
|
|
backend.flexible_member_array_supported = false;
|
|
backend.explicit_struct_type = true;
|
|
backend.use_initializer_list = true;
|
|
|
|
build_function_control_flow_graphs_and_analyze();
|
|
update_active_builtins();
|
|
|
|
uint32_t pass_count = 0;
|
|
do
|
|
{
|
|
if (pass_count >= 3)
|
|
SPIRV_CROSS_THROW("Over 3 compilation loops detected. Must be a bug!");
|
|
|
|
resource_registrations.clear();
|
|
reset();
|
|
|
|
// Move constructor for this type is broken on GCC 4.9 ...
|
|
buffer.reset();
|
|
|
|
emit_header();
|
|
emit_resources();
|
|
|
|
emit_function(get<SPIRFunction>(ir.default_entry_point), Bitset());
|
|
|
|
pass_count++;
|
|
} while (is_forcing_recompilation());
|
|
|
|
// Match opening scope of emit_header().
|
|
end_scope_decl();
|
|
// namespace
|
|
end_scope();
|
|
|
|
// Emit C entry points
|
|
emit_c_linkage();
|
|
|
|
// Entry point in CPP is always main() for the time being.
|
|
get_entry_point().name = "main";
|
|
|
|
return buffer.str();
|
|
}
|
|
|
|
void CompilerCPP::emit_c_linkage()
|
|
{
|
|
statement("");
|
|
|
|
statement("spirv_cross_shader_t *spirv_cross_construct(void)");
|
|
begin_scope();
|
|
statement("return new ", impl_type, "();");
|
|
end_scope();
|
|
|
|
statement("");
|
|
statement("void spirv_cross_destruct(spirv_cross_shader_t *shader)");
|
|
begin_scope();
|
|
statement("delete static_cast<", impl_type, "*>(shader);");
|
|
end_scope();
|
|
|
|
statement("");
|
|
statement("void spirv_cross_invoke(spirv_cross_shader_t *shader)");
|
|
begin_scope();
|
|
statement("static_cast<", impl_type, "*>(shader)->invoke();");
|
|
end_scope();
|
|
|
|
statement("");
|
|
statement("static const struct spirv_cross_interface vtable =");
|
|
begin_scope();
|
|
statement("spirv_cross_construct,");
|
|
statement("spirv_cross_destruct,");
|
|
statement("spirv_cross_invoke,");
|
|
end_scope_decl();
|
|
|
|
statement("");
|
|
statement("const struct spirv_cross_interface *",
|
|
interface_name.empty() ? string("spirv_cross_get_interface") : interface_name, "(void)");
|
|
begin_scope();
|
|
statement("return &vtable;");
|
|
end_scope();
|
|
}
|
|
|
|
void CompilerCPP::emit_function_prototype(SPIRFunction &func, const Bitset &)
|
|
{
|
|
if (func.self != ir.default_entry_point)
|
|
add_function_overload(func);
|
|
|
|
local_variable_names = resource_names;
|
|
string decl;
|
|
|
|
auto &type = get<SPIRType>(func.return_type);
|
|
decl += "inline ";
|
|
decl += type_to_glsl(type);
|
|
decl += " ";
|
|
|
|
if (func.self == ir.default_entry_point)
|
|
{
|
|
decl += "main";
|
|
processing_entry_point = true;
|
|
}
|
|
else
|
|
decl += to_name(func.self);
|
|
|
|
decl += "(";
|
|
for (auto &arg : func.arguments)
|
|
{
|
|
add_local_variable_name(arg.id);
|
|
|
|
decl += argument_decl(arg);
|
|
if (&arg != &func.arguments.back())
|
|
decl += ", ";
|
|
|
|
// Hold a pointer to the parameter so we can invalidate the readonly field if needed.
|
|
auto *var = maybe_get<SPIRVariable>(arg.id);
|
|
if (var)
|
|
var->parameter = &arg;
|
|
}
|
|
|
|
decl += ")";
|
|
statement(decl);
|
|
}
|
|
|
|
string CompilerCPP::argument_decl(const SPIRFunction::Parameter &arg)
|
|
{
|
|
auto &type = expression_type(arg.id);
|
|
bool constref = !type.pointer || arg.write_count == 0;
|
|
|
|
auto &var = get<SPIRVariable>(arg.id);
|
|
|
|
string base = type_to_glsl(type);
|
|
string variable_name = to_name(var.self);
|
|
remap_variable_type_name(type, variable_name, base);
|
|
|
|
for (uint32_t i = 0; i < type.array.size(); i++)
|
|
base = join("std::array<", base, ", ", to_array_size(type, i), ">");
|
|
|
|
return join(constref ? "const " : "", base, " &", variable_name);
|
|
}
|
|
|
|
string CompilerCPP::variable_decl(const SPIRType &type, const string &name, uint32_t /* id */)
|
|
{
|
|
string base = type_to_glsl(type);
|
|
remap_variable_type_name(type, name, base);
|
|
bool runtime = false;
|
|
|
|
for (uint32_t i = 0; i < type.array.size(); i++)
|
|
{
|
|
auto &array = type.array[i];
|
|
if (!array && type.array_size_literal[i])
|
|
{
|
|
// Avoid using runtime arrays with std::array since this is undefined.
|
|
// Runtime arrays cannot be passed around as values, so this is fine.
|
|
runtime = true;
|
|
}
|
|
else
|
|
base = join("std::array<", base, ", ", to_array_size(type, i), ">");
|
|
}
|
|
base += ' ';
|
|
return base + name + (runtime ? "[1]" : "");
|
|
}
|
|
|
|
void CompilerCPP::emit_header()
|
|
{
|
|
auto &execution = get_entry_point();
|
|
|
|
statement("// This C++ shader is autogenerated by spirv-cross.");
|
|
statement("#include \"spirv_cross/internal_interface.hpp\"");
|
|
statement("#include \"spirv_cross/external_interface.h\"");
|
|
// Needed to properly implement GLSL-style arrays.
|
|
statement("#include <array>");
|
|
statement("#include <stdint.h>");
|
|
statement("");
|
|
statement("using namespace spirv_cross;");
|
|
statement("using namespace glm;");
|
|
statement("");
|
|
|
|
statement("namespace Impl");
|
|
begin_scope();
|
|
|
|
switch (execution.model)
|
|
{
|
|
case ExecutionModelGeometry:
|
|
case ExecutionModelTessellationControl:
|
|
case ExecutionModelTessellationEvaluation:
|
|
case ExecutionModelGLCompute:
|
|
case ExecutionModelFragment:
|
|
case ExecutionModelVertex:
|
|
statement("struct Shader");
|
|
begin_scope();
|
|
break;
|
|
|
|
default:
|
|
SPIRV_CROSS_THROW("Unsupported execution model.");
|
|
}
|
|
|
|
switch (execution.model)
|
|
{
|
|
case ExecutionModelGeometry:
|
|
impl_type = "GeometryShader<Impl::Shader, Impl::Shader::Resources>";
|
|
resource_type = "GeometryResources";
|
|
break;
|
|
|
|
case ExecutionModelVertex:
|
|
impl_type = "VertexShader<Impl::Shader, Impl::Shader::Resources>";
|
|
resource_type = "VertexResources";
|
|
break;
|
|
|
|
case ExecutionModelFragment:
|
|
impl_type = "FragmentShader<Impl::Shader, Impl::Shader::Resources>";
|
|
resource_type = "FragmentResources";
|
|
break;
|
|
|
|
case ExecutionModelGLCompute:
|
|
impl_type = join("ComputeShader<Impl::Shader, Impl::Shader::Resources, ", execution.workgroup_size.x, ", ",
|
|
execution.workgroup_size.y, ", ", execution.workgroup_size.z, ">");
|
|
resource_type = "ComputeResources";
|
|
break;
|
|
|
|
case ExecutionModelTessellationControl:
|
|
impl_type = "TessControlShader<Impl::Shader, Impl::Shader::Resources>";
|
|
resource_type = "TessControlResources";
|
|
break;
|
|
|
|
case ExecutionModelTessellationEvaluation:
|
|
impl_type = "TessEvaluationShader<Impl::Shader, Impl::Shader::Resources>";
|
|
resource_type = "TessEvaluationResources";
|
|
break;
|
|
|
|
default:
|
|
SPIRV_CROSS_THROW("Unsupported execution model.");
|
|
}
|
|
}
|