From 042475e88e7486f40e6b1b87d623b36610f526e6 Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Thu, 28 Jul 2016 11:16:02 +0200 Subject: [PATCH] Add support for multiple entry points. - Only consider I/O variables if part of OpEntryPoint. - Keep a safe fallback if #entry-points is 1 to avoid potentially breaking previously working shaders. --- main.cpp | 34 ++++++- .../shaders/asm/comp/multiple-entry.asm.comp | 27 +++++ shaders/asm/comp/multiple-entry.asm.comp | 97 ++++++++++++++++++ spirv_common.hpp | 26 +++++ spirv_cpp.cpp | 11 ++- spirv_cross.cpp | 98 ++++++++++++++++--- spirv_cross.hpp | 36 ++++--- spirv_glsl.cpp | 17 +++- spirv_msl.cpp | 36 ++++--- test_shaders.py | 4 +- 10 files changed, 337 insertions(+), 49 deletions(-) create mode 100644 reference/shaders/asm/comp/multiple-entry.asm.comp create mode 100644 shaders/asm/comp/multiple-entry.asm.comp diff --git a/main.cpp b/main.cpp index 473ca846..e41c251e 100644 --- a/main.cpp +++ b/main.cpp @@ -218,10 +218,37 @@ static void print_resources(const Compiler &compiler, const char *tag, const vec fprintf(stderr, "=============\n\n"); } +static const char *execution_model_to_str(spv::ExecutionModel model) +{ + switch (model) + { + case spv::ExecutionModelVertex: + return "vertex"; + case spv::ExecutionModelTessellationControl: + return "tessellation control"; + case ExecutionModelTessellationEvaluation: + return "tessellation evaluation"; + case ExecutionModelGeometry: + return "geometry"; + case ExecutionModelFragment: + return "fragment"; + case ExecutionModelGLCompute: + return "compute"; + default: + return "???"; + } +} + static void print_resources(const Compiler &compiler, const ShaderResources &res) { uint64_t modes = compiler.get_execution_mode_mask(); + fprintf(stderr, "Entry points:\n"); + auto entry_points = compiler.get_entry_points(); + for (auto &e : entry_points) + fprintf(stderr, " %s (%s)\n", e.c_str(), execution_model_to_str(compiler.get_entry_point(e).model)); + fprintf(stderr, "\n"); + fprintf(stderr, "Execution modes:\n"); for (unsigned i = 0; i < 64; i++) { @@ -348,6 +375,7 @@ struct CLIArguments vector pls_out; vector remaps; vector extensions; + string entry; uint32_t iterations = 1; bool cpp = false; @@ -361,7 +389,7 @@ static void print_help() "version>] [--dump-resources] [--help] [--force-temporary] [--cpp] [--cpp-interface-name ] " "[--metal] [--vulkan-semantics] [--flatten-ubo] [--fixup-clipspace] [--iterations iter] [--pls-in " "format input-name] [--pls-out format output-name] [--remap source_name target_name components] " - "[--extension ext]\n"); + "[--extension ext] [--entry name]\n"); } static bool remap_generic(Compiler &compiler, const vector &resources, const Remap &remap) @@ -480,6 +508,7 @@ int main(int argc, char *argv[]) cbs.add("--metal", [&args](CLIParser &) { args.metal = true; }); cbs.add("--vulkan-semantics", [&args](CLIParser &) { args.vulkan_semantics = true; }); cbs.add("--extension", [&args](CLIParser &parser) { args.extensions.push_back(parser.next_string()); }); + cbs.add("--entry", [&args](CLIParser &parser) { args.entry = parser.next_string(); }); cbs.add("--remap", [&args](CLIParser &parser) { string src = parser.next_string(); string dst = parser.next_string(); @@ -531,6 +560,9 @@ int main(int argc, char *argv[]) else compiler = unique_ptr(new CompilerGLSL(read_spirv_file(args.input))); + if (!args.entry.empty()) + compiler->set_entry_point(args.entry); + if (!args.set_version && !compiler->get_options().version) { fprintf(stderr, "Didn't specify GLSL version and SPIR-V did not specify language.\n"); diff --git a/reference/shaders/asm/comp/multiple-entry.asm.comp b/reference/shaders/asm/comp/multiple-entry.asm.comp new file mode 100644 index 00000000..b30a164f --- /dev/null +++ b/reference/shaders/asm/comp/multiple-entry.asm.comp @@ -0,0 +1,27 @@ +#version 310 es +layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in; + +layout(binding = 0, std430) restrict buffer _6 +{ + ivec4 _0; + uvec4 _1; +} _8; + +layout(binding = 1, std430) restrict buffer _7 +{ + uvec4 _0; + ivec4 _1; +} _9; + +void main() +{ + _9._0 = (_8._1 + uvec4(_8._0)); + _9._0 = (uvec4(_8._0) + _8._1); + _9._0 = (_8._1 + _8._1); + _9._0 = uvec4(_8._0 + _8._0); + _9._1 = ivec4(_8._1 + _8._1); + _9._1 = (_8._0 + _8._0); + _9._1 = (ivec4(_8._1) + _8._0); + _9._1 = (_8._0 + ivec4(_8._1)); +} + diff --git a/shaders/asm/comp/multiple-entry.asm.comp b/shaders/asm/comp/multiple-entry.asm.comp new file mode 100644 index 00000000..0cfb5543 --- /dev/null +++ b/shaders/asm/comp/multiple-entry.asm.comp @@ -0,0 +1,97 @@ +; SPIR-V +; Version: 1.0 +; Generator: Khronos Glslang Reference Front End; 1 +; Bound: 30 +; Schema: 0 + OpCapability Shader + %1 = OpExtInstImport "GLSL.std.450" + OpMemoryModel Logical GLSL450 + OpEntryPoint Fragment %func_alt "main2" %frag_in %frag_out + OpEntryPoint GLCompute %func "main" + OpExecutionMode %func LocalSize 1 1 1 + OpSource ESSL 310 + OpSourceExtension "GL_GOOGLE_cpp_style_line_directive" + OpSourceExtension "GL_GOOGLE_include_directive" + OpMemberDecorate %input_struct 0 Offset 0 + OpMemberDecorate %input_struct 1 Offset 16 + OpMemberDecorate %output_struct 0 Offset 0 + OpMemberDecorate %output_struct 1 Offset 16 + OpDecorate %input_struct BufferBlock + OpDecorate %inputs DescriptorSet 0 + OpDecorate %inputs Binding 0 + OpDecorate %inputs Restrict + OpDecorate %output_struct BufferBlock + OpDecorate %outputs DescriptorSet 0 + OpDecorate %outputs Binding 1 + OpDecorate %outputs Restrict + OpDecorate %frag_in Location 0 + OpDecorate %frag_out Location 0 + + %void = OpTypeVoid + %main_func = OpTypeFunction %void + + %uint = OpTypeInt 32 0 + %uvec4 = OpTypeVector %uint 4 + + %int = OpTypeInt 32 1 + %ivec4 = OpTypeVector %int 4 + + %ivec4_ptr = OpTypePointer Uniform %ivec4 + %uvec4_ptr = OpTypePointer Uniform %uvec4 + + %float = OpTypeFloat 32 + %vec4 = OpTypeVector %float 4 + %vec4_input_ptr = OpTypePointer Input %vec4 + %vec4_output_ptr = OpTypePointer Output %vec4 + + %zero = OpConstant %int 0 + %one = OpConstant %int 1 + + %input_struct = OpTypeStruct %ivec4 %uvec4 + %input_struct_ptr = OpTypePointer Uniform %input_struct + %inputs = OpVariable %input_struct_ptr Uniform + %output_struct = OpTypeStruct %uvec4 %ivec4 + %output_struct_ptr = OpTypePointer Uniform %output_struct + %outputs = OpVariable %output_struct_ptr Uniform + + %frag_in = OpVariable %vec4_input_ptr Input + %frag_out = OpVariable %vec4_output_ptr Output + + %func = OpFunction %void None %main_func + %block = OpLabel + + %input1_ptr = OpAccessChain %ivec4_ptr %inputs %zero + %input0_ptr = OpAccessChain %uvec4_ptr %inputs %one + %input1 = OpLoad %ivec4 %input1_ptr + %input0 = OpLoad %uvec4 %input0_ptr + + %output_ptr_uvec4 = OpAccessChain %uvec4_ptr %outputs %zero + %output_ptr_ivec4 = OpAccessChain %ivec4_ptr %outputs %one + +; Test all variants of IAdd + %result_iadd_0 = OpIAdd %uvec4 %input0 %input1 + %result_iadd_1 = OpIAdd %uvec4 %input1 %input0 + %result_iadd_2 = OpIAdd %uvec4 %input0 %input0 + %result_iadd_3 = OpIAdd %uvec4 %input1 %input1 + %result_iadd_4 = OpIAdd %ivec4 %input0 %input0 + %result_iadd_5 = OpIAdd %ivec4 %input1 %input1 + %result_iadd_6 = OpIAdd %ivec4 %input0 %input1 + %result_iadd_7 = OpIAdd %ivec4 %input1 %input0 + OpStore %output_ptr_uvec4 %result_iadd_0 + OpStore %output_ptr_uvec4 %result_iadd_1 + OpStore %output_ptr_uvec4 %result_iadd_2 + OpStore %output_ptr_uvec4 %result_iadd_3 + OpStore %output_ptr_ivec4 %result_iadd_4 + OpStore %output_ptr_ivec4 %result_iadd_5 + OpStore %output_ptr_ivec4 %result_iadd_6 + OpStore %output_ptr_ivec4 %result_iadd_7 + + OpReturn + OpFunctionEnd + + %func_alt = OpFunction %void None %main_func + %block_alt = OpLabel + %frag_input_value = OpLoad %vec4 %frag_in + OpStore %frag_out %frag_input_value + OpReturn + OpFunctionEnd diff --git a/spirv_common.hpp b/spirv_common.hpp index 6d020b05..e1a3f849 100644 --- a/spirv_common.hpp +++ b/spirv_common.hpp @@ -232,6 +232,32 @@ struct SPIRExtension : IVariant Extension ext; }; +// SPIREntryPoint is not a variant since its IDs are used to decorate OpFunction, +// so in order to avoid conflicts, we can't stick them in the ids array. +struct SPIREntryPoint +{ + SPIREntryPoint(uint32_t self_, spv::ExecutionModel execution_model, std::string entry_name) + : self(self_) + , name(std::move(entry_name)) + , model(execution_model) + { + } + SPIREntryPoint() = default; + + uint32_t self = 0; + std::string name; + std::vector interface_variables; + + uint64_t flags = 0; + struct + { + uint32_t x = 0, y = 0, z = 0; + } workgroup_size; + uint32_t invocations = 0; + uint32_t output_vertices = 0; + spv::ExecutionModel model = {}; +}; + struct SPIRExpression : IVariant { enum diff --git a/spirv_cpp.cpp b/spirv_cpp.cpp index b22e7e3e..588ef7c7 100644 --- a/spirv_cpp.cpp +++ b/spirv_cpp.cpp @@ -193,7 +193,8 @@ void CompilerCPP::emit_resources() auto &type = get(var.basetype); if (var.storage != StorageClassFunction && !is_builtin_variable(var) && !var.remapped_variable && - type.pointer && (var.storage == StorageClassInput || var.storage == StorageClassOutput)) + type.pointer && (var.storage == StorageClassInput || var.storage == StorageClassOutput) && + interface_variable_exists_in_entry_point(var.self)) { emit_interface_block(var); } @@ -244,7 +245,7 @@ void CompilerCPP::emit_resources() statement(""); statement("Resources* __res;"); - if (execution.model == ExecutionModelGLCompute) + if (get_entry_point().model == ExecutionModelGLCompute) statement("ComputePrivateResources __priv_res;"); statement(""); @@ -299,7 +300,7 @@ string CompilerCPP::compile() emit_header(); emit_resources(); - emit_function(get(execution.entry_point), 0); + emit_function(get(entry_point), 0); pass_count++; } while (force_recompile); @@ -362,7 +363,7 @@ void CompilerCPP::emit_function_prototype(SPIRFunction &func, uint64_t) decl += type_to_glsl(type); decl += " "; - if (func.self == execution.entry_point) + if (func.self == entry_point) { decl += "main"; processing_entry_point = true; @@ -424,6 +425,8 @@ string CompilerCPP::variable_decl(const SPIRType &type, const string &name) 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\""); diff --git a/spirv_cross.cpp b/spirv_cross.cpp index d8b84eca..e010a07d 100644 --- a/spirv_cross.cpp +++ b/spirv_cross.cpp @@ -419,7 +419,7 @@ ShaderResources Compiler::get_shader_resources() const continue; // Input - if (var.storage == StorageClassInput) + if (var.storage == StorageClassInput && interface_variable_exists_in_entry_point(var.self)) { if (meta[type.self].decoration.decoration_flags & (1ull << DecorationBlock)) res.stage_inputs.push_back({ var.self, var.basetype, type.self, meta[type.self].decoration.alias }); @@ -432,7 +432,7 @@ ShaderResources Compiler::get_shader_resources() const res.subpass_inputs.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias }); } // Outputs - else if (var.storage == StorageClassOutput) + else if (var.storage == StorageClassOutput && interface_variable_exists_in_entry_point(var.self)) { if (meta[type.self].decoration.decoration_flags & (1ull << DecorationBlock)) res.stage_outputs.push_back({ var.self, var.basetype, type.self, meta[type.self].decoration.alias }); @@ -929,20 +929,23 @@ void Compiler::parse(const Instruction &instruction) case OpEntryPoint: { - if (execution.entry_point) - throw CompilerError("More than one entry point not supported."); + auto itr = entry_points.emplace(ops[1], SPIREntryPoint(ops[1], static_cast(ops[0]), + extract_string(spirv, instruction.offset + 2))); + auto &e = itr.first->second; - execution.model = static_cast(ops[0]); - execution.entry_point = ops[1]; + // Strings need nul-terminator and consume the whole word. + uint32_t strlen_words = (e.name.size() + 1 + 3) >> 2; + e.interface_variables.insert(end(e.interface_variables), ops + strlen_words + 2, ops + instruction.length); + + // If we don't have an entry, make the first one our "default". + if (!entry_point) + entry_point = ops[1]; break; } case OpExecutionMode: { - uint32_t entry = ops[0]; - if (entry != execution.entry_point) - throw CompilerError("Cannot set execution mode to non-existing entry point."); - + auto &execution = entry_points[ops[0]]; auto mode = static_cast(ops[1]); execution.flags |= 1ull << mode; @@ -1921,7 +1924,7 @@ std::vector Compiler::get_active_buffer_ranges(uint32_t id) const { std::vector ranges; BufferAccessHandler handler(*this, ranges, id); - traverse_all_reachable_opcodes(get(execution.entry_point), handler); + traverse_all_reachable_opcodes(get(entry_point), handler); return ranges; } @@ -1974,11 +1977,13 @@ bool Compiler::types_are_logically_equivalent(const SPIRType &a, const SPIRType uint64_t Compiler::get_execution_mode_mask() const { - return execution.flags; + return get_entry_point().flags; } void Compiler::set_execution_mode(ExecutionMode mode, uint32_t arg0, uint32_t arg1, uint32_t arg2) { + auto &execution = get_entry_point(); + execution.flags |= 1ull << mode; switch (mode) { @@ -2003,11 +2008,13 @@ void Compiler::set_execution_mode(ExecutionMode mode, uint32_t arg0, uint32_t ar void Compiler::unset_execution_mode(ExecutionMode mode) { + auto &execution = get_entry_point(); execution.flags &= ~(1ull << mode); } uint32_t Compiler::get_execution_mode_argument(spv::ExecutionMode mode, uint32_t index) const { + auto &execution = get_entry_point(); switch (mode) { case ExecutionModeLocalSize: @@ -2036,6 +2043,7 @@ uint32_t Compiler::get_execution_mode_argument(spv::ExecutionMode mode, uint32_t ExecutionModel Compiler::get_execution_model() const { + auto &execution = get_entry_point(); return execution.model; } @@ -2076,3 +2084,69 @@ void Compiler::inherit_expression_dependencies(uint32_t dst, uint32_t source_exp // Eliminate duplicated dependencies. e_deps.erase(unique(begin(e_deps), end(e_deps)), end(e_deps)); } + +vector Compiler::get_entry_points() const +{ + vector entries; + for (auto &entry : entry_points) + entries.push_back(entry.second.name); + return entries; +} + +void Compiler::set_entry_point(const std::string &name) +{ + auto &entry = get_entry_point(name); + entry_point = entry.self; +} + +SPIREntryPoint &Compiler::get_entry_point(const std::string &name) +{ + auto itr = + find_if(begin(entry_points), end(entry_points), + [&](const std::pair &entry) -> bool { return entry.second.name == name; }); + + if (itr == end(entry_points)) + throw CompilerError("Entry point does not exist."); + + return itr->second; +} + +const SPIREntryPoint &Compiler::get_entry_point(const std::string &name) const +{ + auto itr = + find_if(begin(entry_points), end(entry_points), + [&](const std::pair &entry) -> bool { return entry.second.name == name; }); + + if (itr == end(entry_points)) + throw CompilerError("Entry point does not exist."); + + return itr->second; +} + +const SPIREntryPoint &Compiler::get_entry_point() const +{ + return entry_points.find(entry_point)->second; +} + +SPIREntryPoint &Compiler::get_entry_point() +{ + return entry_points.find(entry_point)->second; +} + +bool Compiler::interface_variable_exists_in_entry_point(uint32_t id) const +{ + auto &var = get(id); + if (var.storage != StorageClassInput && var.storage != StorageClassOutput) + throw CompilerError("Only Input and Output variables are part of a shader linking interface."); + + // This is to avoid potential problems with very old glslang versions which did + // not emit input/output interfaces properly. + // We can assume they only had a single entry point, and single entry point + // shaders could easily be assumed to use every interface variable anyways. + if (entry_points.size() <= 1) + return true; + + auto &execution = get_entry_point(); + return find(begin(execution.interface_variables), end(execution.interface_variables), id) != + end(execution.interface_variables); +} diff --git a/spirv_cross.hpp b/spirv_cross.hpp index e1df27ee..d1ca4c10 100644 --- a/spirv_cross.hpp +++ b/spirv_cross.hpp @@ -191,6 +191,18 @@ public: void set_subpass_input_remapped_components(uint32_t id, uint32_t components); uint32_t get_subpass_input_remapped_components(uint32_t id) const; + // All operations work on the current entry point. + // Entry points can be swapped out with set_entry_point(). + // Entry points should be set right after the constructor completes as some reflection functions traverse the graph from the entry point. + // Resource reflection also depends on the entry point. + // By default, the current entry point is set to the first OpEntryPoint which appears in the SPIR-V module. + std::vector get_entry_points() const; + void set_entry_point(const std::string &name); + + // Returns the internal data structure for entry points to allow poking around. + const SPIREntryPoint &get_entry_point(const std::string &name) const; + SPIREntryPoint &get_entry_point(const std::string &name); + // Query and modify OpExecutionMode. uint64_t get_execution_mode_mask() const; void unset_execution_mode(spv::ExecutionMode mode); @@ -266,20 +278,12 @@ protected: return nullptr; } - struct Execution - { - uint64_t flags = 0; - spv::ExecutionModel model; - uint32_t entry_point = 0; - struct - { - uint32_t x = 0, y = 0, z = 0; - } workgroup_size; - uint32_t invocations = 0; - uint32_t output_vertices = 0; - - Execution() = default; - } execution; + uint32_t entry_point = 0; + // Normally, we'd stick SPIREntryPoint in ids array, but it conflicts with SPIRFunction. + // Entry points can therefore be seen as some sort of meta structure. + std::unordered_map entry_points; + const SPIREntryPoint &get_entry_point() const; + SPIREntryPoint &get_entry_point(); struct Source { @@ -359,6 +363,10 @@ protected: bool types_are_logically_equivalent(const SPIRType &a, const SPIRType &b) const; void inherit_expression_dependencies(uint32_t dst, uint32_t source); + // For proper multiple entry point support, allow querying if an Input or Output + // variable is part of that entry points interface. + bool interface_variable_exists_in_entry_point(uint32_t id) const; + private: void parse(); void parse(const Instruction &i); diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index 8930f1a8..cfb652b2 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -221,7 +221,7 @@ string CompilerGLSL::compile() emit_header(); emit_resources(); - emit_function(get(execution.entry_point), 0); + emit_function(get(entry_point), 0); pass_count++; } while (force_recompile); @@ -231,6 +231,7 @@ string CompilerGLSL::compile() void CompilerGLSL::emit_header() { + auto &execution = get_entry_point(); statement("#version ", options.version, options.es && options.version > 100 ? " es" : ""); for (auto &header : header_lines) @@ -948,6 +949,7 @@ void CompilerGLSL::emit_buffer_block(const SPIRVariable &var) void CompilerGLSL::emit_interface_block(const SPIRVariable &var) { + auto &execution = get_entry_point(); auto &type = get(var.basetype); // Either make it plain in/out or in/out blocks depending on what shader is doing ... @@ -1064,6 +1066,7 @@ string CompilerGLSL::remap_swizzle(uint32_t result_type, uint32_t input_componen void CompilerGLSL::emit_pls() { + auto &execution = get_entry_point(); if (execution.model != ExecutionModelFragment) throw CompilerError("Pixel local storage only supported in fragment shaders."); @@ -1096,6 +1099,8 @@ void CompilerGLSL::emit_pls() void CompilerGLSL::emit_resources() { + auto &execution = get_entry_point(); + // Legacy GL uses gl_FragData[], redeclare all fragment outputs // with builtins. if (execution.model == ExecutionModelFragment && is_legacy()) @@ -1183,7 +1188,8 @@ void CompilerGLSL::emit_resources() auto &type = get(var.basetype); if (var.storage != StorageClassFunction && !is_builtin_variable(var) && !var.remapped_variable && - type.pointer && (var.storage == StorageClassInput || var.storage == StorageClassOutput)) + type.pointer && (var.storage == StorageClassInput || var.storage == StorageClassOutput) && + interface_variable_exists_in_entry_point(var.self)) { emit_interface_block(var); emitted = true; @@ -4006,7 +4012,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) case OpControlBarrier: { // Ignore execution and memory scope. - if (execution.model == ExecutionModelGLCompute) + if (get_entry_point().model == ExecutionModelGLCompute) { uint32_t mem = get(ops[2]).scalar(); if (mem == MemorySemanticsWorkgroupMemoryMask) @@ -4101,6 +4107,8 @@ const char *CompilerGLSL::flags_to_precision_qualifiers_glsl(const SPIRType &typ { if (options.es) { + auto &execution = get_entry_point(); + // Structs do not have precision qualifiers, neither do doubles (desktop only anyways, so no mediump/highp). if (type.basetype != SPIRType::Float && type.basetype != SPIRType::Int && type.basetype != SPIRType::UInt && type.basetype != SPIRType::Image && type.basetype != SPIRType::SampledImage && @@ -4525,7 +4533,7 @@ void CompilerGLSL::emit_function_prototype(SPIRFunction &func, uint64_t return_f decl += type_to_glsl(type); decl += " "; - if (func.self == execution.entry_point) + if (func.self == entry_point) { decl += "main"; processing_entry_point = true; @@ -4625,6 +4633,7 @@ void CompilerGLSL::emit_function(SPIRFunction &func, uint64_t return_flags) void CompilerGLSL::emit_fixup() { + auto &execution = get_entry_point(); if (execution.model == ExecutionModelVertex && options.vertex.fixup_clipspace) { const char *suffix = backend.float_literal_suffix ? "f" : ""; diff --git a/spirv_msl.cpp b/spirv_msl.cpp index dd9dfc16..19c6be2b 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -84,7 +84,7 @@ string CompilerMSL::compile(MSLConfiguration &msl_cfg, vector *p_ emit_header(); emit_resources(); emit_function_declarations(); - emit_function(get(execution.entry_point), 0); + emit_function(get(entry_point), 0); pass_count++; } while (force_recompile); @@ -120,7 +120,7 @@ void CompilerMSL::extract_builtins() // Non-constant variables cannot have global scope in Metal. void CompilerMSL::localize_global_variables() { - auto &entry_func = get(execution.entry_point); + auto &entry_func = get(entry_point); auto iter = global_variables.begin(); while (iter != global_variables.end()) { @@ -141,6 +141,8 @@ void CompilerMSL::localize_global_variables() // Adds any interface structure variables needed by this shader void CompilerMSL::add_interface_structs() { + auto &execution = get_entry_point(); + stage_in_var_ids.clear(); qual_pos_var_name = ""; @@ -170,6 +172,8 @@ void CompilerMSL::add_interface_structs() // from the binding info provided during compiler construction, matching by location. void CompilerMSL::bind_vertex_attributes(std::set &bindings) { + auto &execution = get_entry_point(); + if (execution.model == ExecutionModelVertex) { for (auto &id : ids) @@ -179,8 +183,8 @@ void CompilerMSL::bind_vertex_attributes(std::set &bindings) auto &var = id.get(); auto &type = get(var.basetype); - if (var.storage == StorageClassInput && (!is_builtin_variable(var)) && !var.remapped_variable && - type.pointer) + if (var.storage == StorageClassInput && interface_variable_exists_in_entry_point(var.self) && + (!is_builtin_variable(var)) && !var.remapped_variable && type.pointer) { auto &dec = meta[var.self].decoration; MSLVertexAttr *p_va = vtx_attrs_by_location[dec.location]; @@ -207,6 +211,7 @@ void CompilerMSL::bind_vertex_attributes(std::set &bindings) // Returns the ID of the newly added variable, or zero if no variable was added. uint32_t CompilerMSL::add_interface_struct(StorageClass storage, uint32_t vtx_binding) { + auto &execution = get_entry_point(); bool incl_builtins = (storage == StorageClassOutput); bool match_binding = (execution.model == ExecutionModelVertex) && (storage == StorageClassInput); @@ -220,8 +225,9 @@ uint32_t CompilerMSL::add_interface_struct(StorageClass storage, uint32_t vtx_bi auto &type = get(var.basetype); auto &dec = meta[var.self].decoration; - if (var.storage == storage && (!is_builtin_variable(var) || incl_builtins) && - (!match_binding || (vtx_binding == dec.binding)) && !var.remapped_variable && type.pointer) + if (var.storage == storage && interface_variable_exists_in_entry_point(var.self) && + (!is_builtin_variable(var) || incl_builtins) && (!match_binding || (vtx_binding == dec.binding)) && + !var.remapped_variable && type.pointer) { vars.push_back(&var); } @@ -264,7 +270,7 @@ uint32_t CompilerMSL::add_interface_struct(StorageClass storage, uint32_t vtx_bi // Add the output interface struct as a local variable to the entry function, // and force the entry function to return the output interface struct from // any blocks that perform a function return. - auto &entry_func = get(execution.entry_point); + auto &entry_func = get(entry_point); entry_func.add_local_variable(ib_var_id); for (auto &blk_id : entry_func.blocks) { @@ -456,7 +462,7 @@ void CompilerMSL::emit_function_declarations() if (id.get_type() == TypeFunction) { auto &func = id.get(); - if (func.self != execution.entry_point) + if (func.self != entry_point) emit_function_prototype(func, true); } @@ -475,7 +481,7 @@ void CompilerMSL::emit_function_prototype(SPIRFunction &func, bool is_decl) local_variable_names = resource_names; string decl; - processing_entry_point = (func.self == execution.entry_point); + processing_entry_point = (func.self == entry_point); auto &type = get(func.return_type); decl += func_type_decl(type); @@ -843,6 +849,8 @@ string CompilerMSL::to_sampler_expression(uint32_t id) // Called automatically at the end of the entry point function void CompilerMSL::emit_fixup() { + auto &execution = get_entry_point(); + if ((execution.model == ExecutionModelVertex) && stage_out_var_id && !qual_pos_var_name.empty()) { if (options.vertex.fixup_clipspace) @@ -868,6 +876,7 @@ string CompilerMSL::member_decl(const SPIRType &type, const SPIRType &membertype // Return a MSL qualifier for the specified function attribute member string CompilerMSL::member_attribute_qualifier(const SPIRType &type, uint32_t index) { + auto &execution = get_entry_point(); BuiltIn builtin; bool is_builtin = is_member_builtin(type, index, &builtin); @@ -1008,7 +1017,7 @@ string CompilerMSL::constant_expression(const SPIRConstant &c) // entry type if the current function is the entry point function string CompilerMSL::func_type_decl(SPIRType &type) { - + auto &execution = get_entry_point(); // The regular function return type. If not processing the entry point function, that's all we need string return_type = type_to_glsl(type); if (!processing_entry_point) @@ -1056,6 +1065,7 @@ string CompilerMSL::clean_func_name(string func_name) // Returns a string containing a comma-delimited list of args for the entry point function string CompilerMSL::entry_point_args(bool append_comma) { + auto &execution = get_entry_point(); string ep_args; // Stage-in structures @@ -1145,7 +1155,7 @@ string CompilerMSL::entry_point_args(bool append_comma) // Returns the Metal index of the resource of the specified type as used by the specified variable. uint32_t CompilerMSL::get_metal_resource_index(SPIRVariable &var, SPIRType::BaseType basetype) { - + auto &execution = get_entry_point(); auto &var_dec = meta[var.self].decoration; uint32_t var_desc_set = (var.storage == StorageClassPushConstant) ? kPushConstDescSet : var_dec.set; uint32_t var_binding = (var.storage == StorageClassPushConstant) ? kPushConstBinding : var_dec.binding; @@ -1189,7 +1199,7 @@ uint32_t CompilerMSL::get_metal_resource_index(SPIRVariable &var, SPIRType::Base // Returns the name of the entry point of this shader string CompilerMSL::get_entry_point_name() { - return clean_func_name(to_name(execution.entry_point)); + return clean_func_name(to_name(entry_point)); } // Returns the name of either the vertex index or instance index builtin @@ -1446,6 +1456,8 @@ string CompilerMSL::builtin_to_glsl(BuiltIn builtin) // Returns an MSL string attribute qualifer for a SPIR-V builtin string CompilerMSL::builtin_qualifier(BuiltIn builtin) { + auto &execution = get_entry_point(); + switch (builtin) { // Vertex function in diff --git a/test_shaders.py b/test_shaders.py index c9d5b95c..3a43bcb3 100755 --- a/test_shaders.py +++ b/test_shaders.py @@ -86,14 +86,14 @@ def cross_compile(shader, vulkan, spirv): # subprocess.check_call(['spirv-val', spirv_path]) spirv_cross_path = './spirv-cross' - subprocess.check_call([spirv_cross_path, '--output', glsl_path, spirv_path]) + subprocess.check_call([spirv_cross_path, '--entry', 'main', '--output', glsl_path, spirv_path]) # A shader might not be possible to make valid GLSL from, skip validation for this case. if (not ('nocompat' in glsl_path)) and (not spirv): validate_shader(glsl_path, False) if vulkan or spirv: - subprocess.check_call([spirv_cross_path, '--vulkan-semantics', '--output', vulkan_glsl_path, spirv_path]) + subprocess.check_call([spirv_cross_path, '--entry', 'main', '--vulkan-semantics', '--output', vulkan_glsl_path, spirv_path]) validate_shader(vulkan_glsl_path, vulkan) return (spirv_path, glsl_path, vulkan_glsl_path if vulkan else None)