5547b25afe
Undef values may be of struct type and may be used in constants. Therefore, they must be interleaved with constants and types. Fixes the rest of the Vulkan CTS test `dEQP-VK.spirv_assembly.instruction.compute.opundef.undefined_spec_constant_composite`. (Please excuse the churn in the reference output; it's an inevitable result of this change.)
554 lines
16 KiB
C++
554 lines
16 KiB
C++
/*
|
|
* Copyright 2015-2021 Arm Limited
|
|
* SPDX-License-Identifier: Apache-2.0 OR MIT
|
|
*
|
|
* 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.
|
|
*/
|
|
|
|
/*
|
|
* At your option, you may choose to accept this material under either:
|
|
* 1. The Apache License, Version 2.0, found at <http://www.apache.org/licenses/LICENSE-2.0>, or
|
|
* 2. The MIT License, found at <http://opensource.org/licenses/MIT>.
|
|
*/
|
|
|
|
#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("");
|
|
|
|
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()
|
|
{
|
|
ir.fixup_reserved_names();
|
|
|
|
// 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.unsized_array_supported = false;
|
|
backend.explicit_struct_type = true;
|
|
backend.use_initializer_list = true;
|
|
|
|
fixup_type_alias();
|
|
reorder_type_alias();
|
|
build_function_control_flow_graphs_and_analyze();
|
|
update_active_builtins();
|
|
|
|
uint32_t pass_count = 0;
|
|
do
|
|
{
|
|
resource_registrations.clear();
|
|
reset(pass_count);
|
|
|
|
// 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.");
|
|
}
|
|
}
|