From 5bcf02f7c940ce98c86bedf265589722cedd58ed Mon Sep 17 00:00:00 2001 From: Hans-Kristian Arntzen Date: Fri, 5 Oct 2018 11:30:57 +0200 Subject: [PATCH] Hoist out parsing module from spirv_cross::Compiler. This is a large refactor which splits out the SPIR-V parser from Compiler and moves it into its more appropriately named Parser module. The Parser is responsible for building a ParsedIR structure which is then consumed by one or more compilers. Compiler can take a ParsedIR by value or move reference. This should allow for optimal case for both multiple compilations and single compilation scenarios. --- CMakeLists.txt | 4 + main.cpp | 20 +- msvc/SPIRV-Cross.vcxproj | 4 + msvc/SPIRV-Cross.vcxproj.filters | 12 + spirv_common.hpp | 94 +- spirv_cpp.cpp | 40 +- spirv_cpp.hpp | 16 +- spirv_cross.cpp | 1962 ++++++------------------------ spirv_cross.hpp | 84 +- spirv_cross_parsed_ir.cpp | 557 +++++++++ spirv_cross_parsed_ir.hpp | 129 ++ spirv_glsl.cpp | 175 ++- spirv_glsl.hpp | 24 +- spirv_hlsl.cpp | 79 +- spirv_hlsl.hpp | 16 +- spirv_msl.cpp | 157 ++- spirv_msl.hpp | 7 + spirv_parser.cpp | 1025 ++++++++++++++++ spirv_parser.hpp | 94 ++ spirv_reflect.cpp | 8 +- spirv_reflect.hpp | 18 +- 21 files changed, 2606 insertions(+), 1919 deletions(-) create mode 100644 spirv_cross_parsed_ir.cpp create mode 100644 spirv_cross_parsed_ir.hpp create mode 100644 spirv_parser.cpp create mode 100644 spirv_parser.hpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 1f5eaec30..c80ae368b 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -82,6 +82,10 @@ spirv_cross_add_library(spirv-cross-core spirv_cross_core STATIC ${CMAKE_CURRENT_SOURCE_DIR}/spirv.hpp ${CMAKE_CURRENT_SOURCE_DIR}/spirv_cross.hpp ${CMAKE_CURRENT_SOURCE_DIR}/spirv_cross.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/spirv_parser.hpp + ${CMAKE_CURRENT_SOURCE_DIR}/spirv_parser.cpp + ${CMAKE_CURRENT_SOURCE_DIR}/spirv_cross_parsed_ir.hpp + ${CMAKE_CURRENT_SOURCE_DIR}/spirv_cross_parsed_ir.cpp ${CMAKE_CURRENT_SOURCE_DIR}/spirv_cfg.hpp ${CMAKE_CURRENT_SOURCE_DIR}/spirv_cfg.cpp) diff --git a/main.cpp b/main.cpp index 8a86107b4..a8e1cd4bd 100644 --- a/main.cpp +++ b/main.cpp @@ -19,6 +19,7 @@ #include "spirv_glsl.hpp" #include "spirv_hlsl.hpp" #include "spirv_msl.hpp" +#include "spirv_parser.hpp" #include "spirv_reflect.hpp" #include #include @@ -190,7 +191,7 @@ static vector read_spirv_file(const char *path) FILE *file = fopen(path, "rb"); if (!file) { - fprintf(stderr, "Failed to open SPIRV file: %s\n", path); + fprintf(stderr, "Failed to open SPIR-V file: %s\n", path); return {}; } @@ -797,10 +798,17 @@ static int main_inner(int argc, char *argv[]) return EXIT_FAILURE; } + auto spirv_file = read_spirv_file(args.input); + if (spirv_file.empty()) + return EXIT_FAILURE; + Parser spirv_parser(move(spirv_file)); + + spirv_parser.parse(); + // Special case reflection because it has little to do with the path followed by code-outputting compilers if (!args.reflect.empty()) { - CompilerReflection compiler(read_spirv_file(args.input)); + CompilerReflection compiler(move(spirv_parser.get_parsed_ir())); compiler.set_format(args.reflect); auto json = compiler.compile(); if (args.output) @@ -816,13 +824,13 @@ static int main_inner(int argc, char *argv[]) if (args.cpp) { - compiler = unique_ptr(new CompilerCPP(read_spirv_file(args.input))); + compiler.reset(new CompilerCPP(move(spirv_parser.get_parsed_ir()))); if (args.cpp_interface_name) static_cast(compiler.get())->set_interface_name(args.cpp_interface_name); } else if (args.msl) { - compiler = unique_ptr(new CompilerMSL(read_spirv_file(args.input))); + compiler.reset(new CompilerMSL(move(spirv_parser.get_parsed_ir()))); auto *msl_comp = static_cast(compiler.get()); auto msl_opts = msl_comp->get_msl_options(); @@ -834,13 +842,13 @@ static int main_inner(int argc, char *argv[]) msl_comp->set_msl_options(msl_opts); } else if (args.hlsl) - compiler = unique_ptr(new CompilerHLSL(read_spirv_file(args.input))); + compiler.reset(new CompilerHLSL(move(spirv_parser.get_parsed_ir()))); else { combined_image_samplers = !args.vulkan_semantics; if (!args.vulkan_semantics) build_dummy_sampler = true; - compiler = unique_ptr(new CompilerGLSL(read_spirv_file(args.input))); + compiler.reset(new CompilerGLSL(move(spirv_parser.get_parsed_ir()))); } if (!args.variable_type_remaps.empty()) diff --git a/msvc/SPIRV-Cross.vcxproj b/msvc/SPIRV-Cross.vcxproj index d2287dcbe..bb2841ea2 100644 --- a/msvc/SPIRV-Cross.vcxproj +++ b/msvc/SPIRV-Cross.vcxproj @@ -131,6 +131,8 @@ + + @@ -144,6 +146,8 @@ + + diff --git a/msvc/SPIRV-Cross.vcxproj.filters b/msvc/SPIRV-Cross.vcxproj.filters index 3ff2c29a5..0e80b2012 100644 --- a/msvc/SPIRV-Cross.vcxproj.filters +++ b/msvc/SPIRV-Cross.vcxproj.filters @@ -36,6 +36,12 @@ Source Files + + Source Files + + + Source Files + Source Files @@ -71,6 +77,12 @@ Header Files + + Header Files + + + Header Files + Header Files diff --git a/spirv_common.hpp b/spirv_common.hpp index 519438c40..42d464840 100644 --- a/spirv_common.hpp +++ b/spirv_common.hpp @@ -42,8 +42,8 @@ namespace spirv_cross #ifndef _MSC_VER [[noreturn]] #endif - inline void - report_and_abort(const std::string &msg) +inline void +report_and_abort(const std::string &msg) { #ifdef NDEBUG (void)msg; @@ -68,6 +68,17 @@ class CompilerError : public std::runtime_error #define SPIRV_CROSS_THROW(x) throw CompilerError(x) #endif +//#define SPIRV_CROSS_COPY_CONSTRUCTOR_SANITIZE + +// MSVC 2013 does not have noexcept. We need this for Variant to get move constructor to work correctly +// instead of copy constructor. +// MSVC 2013 ignores that move constructors cannot throw in std::vector, so just don't define it. +#if defined(_MSC_VER) && _MSC_VER < 1900 +#define SPIRV_CROSS_NOEXCEPT +#else +#define SPIRV_CROSS_NOEXCEPT noexcept +#endif + #if __cplusplus >= 201402l #define SPIRV_CROSS_DEPRECATED(reason) [[deprecated(reason)]] #elif defined(__GNUC__) @@ -282,21 +293,27 @@ inline std::string convert_to_string(double t) struct Instruction { - Instruction(const std::vector &spirv, uint32_t &index); - - uint16_t op; - uint16_t count; - uint32_t offset; - uint32_t length; + uint16_t op = 0; + uint16_t count = 0; + uint32_t offset = 0; + uint32_t length = 0; }; // Helper for Variant interface. struct IVariant { virtual ~IVariant() = default; + virtual std::unique_ptr clone() = 0; + uint32_t self = 0; }; +#define SPIRV_CROSS_DECLARE_CLONE(T) \ + std::unique_ptr clone() override \ + { \ + return std::unique_ptr(new T(*this)); \ + } + enum Types { TypeNone, @@ -326,6 +343,8 @@ struct SPIRUndef : IVariant { } uint32_t basetype; + + SPIRV_CROSS_DECLARE_CLONE(SPIRUndef) }; // This type is only used by backends which need to access the combined image and sampler IDs separately after @@ -345,6 +364,8 @@ struct SPIRCombinedImageSampler : IVariant uint32_t combined_type; uint32_t image; uint32_t sampler; + + SPIRV_CROSS_DECLARE_CLONE(SPIRCombinedImageSampler) }; struct SPIRConstantOp : IVariant @@ -364,6 +385,8 @@ struct SPIRConstantOp : IVariant spv::Op opcode; std::vector arguments; uint32_t basetype; + + SPIRV_CROSS_DECLARE_CLONE(SPIRConstantOp) }; struct SPIRType : IVariant @@ -438,6 +461,8 @@ struct SPIRType : IVariant // Used in backends to avoid emitting members with conflicting names. std::unordered_set member_name_cache; + + SPIRV_CROSS_DECLARE_CLONE(SPIRType) }; struct SPIRExtension : IVariant @@ -463,6 +488,7 @@ struct SPIRExtension : IVariant } Extension ext; + SPIRV_CROSS_DECLARE_CLONE(SPIRExtension) }; // SPIREntryPoint is not a variant since its IDs are used to decorate OpFunction, @@ -533,6 +559,8 @@ struct SPIRExpression : IVariant // A list of expressions which this expression depends on. std::vector expression_dependencies; + + SPIRV_CROSS_DECLARE_CLONE(SPIRExpression) }; struct SPIRFunctionPrototype : IVariant @@ -549,6 +577,8 @@ struct SPIRFunctionPrototype : IVariant uint32_t return_type; std::vector parameter_types; + + SPIRV_CROSS_DECLARE_CLONE(SPIRFunctionPrototype) }; struct SPIRBlock : IVariant @@ -684,6 +714,8 @@ struct SPIRBlock : IVariant // sub-group-like operations. // Make sure that we only use these expressions in the original block. std::vector invalidate_expressions; + + SPIRV_CROSS_DECLARE_CLONE(SPIRBlock) }; struct SPIRFunction : IVariant @@ -769,6 +801,8 @@ struct SPIRFunction : IVariant bool active = false; bool flush_undeclared = true; bool do_combined_parameters = true; + + SPIRV_CROSS_DECLARE_CLONE(SPIRFunction) }; struct SPIRAccessChain : IVariant @@ -803,6 +837,8 @@ struct SPIRAccessChain : IVariant uint32_t matrix_stride = 0; bool row_major_matrix = false; bool immutable = false; + + SPIRV_CROSS_DECLARE_CLONE(SPIRAccessChain) }; struct SPIRVariable : IVariant @@ -856,6 +892,8 @@ struct SPIRVariable : IVariant bool loop_variable_enable = false; SPIRFunction::Parameter *parameter = nullptr; + + SPIRV_CROSS_DECLARE_CLONE(SPIRVariable) }; struct SPIRConstant : IVariant @@ -1111,6 +1149,8 @@ struct SPIRConstant : IVariant // For composites which are constant arrays, etc. std::vector subconstants; + + SPIRV_CROSS_DECLARE_CLONE(SPIRConstant) }; class Variant @@ -1118,21 +1158,50 @@ class Variant public: // MSVC 2013 workaround, we shouldn't need these constructors. Variant() = default; - Variant(Variant &&other) + + // Marking custom move constructor as noexcept is important. + Variant(Variant &&other) SPIRV_CROSS_NOEXCEPT { *this = std::move(other); } - Variant &operator=(Variant &&other) + + Variant(const Variant &variant) + { + *this = variant; + } + + // Marking custom move constructor as noexcept is important. + Variant &operator=(Variant &&other) SPIRV_CROSS_NOEXCEPT { if (this != &other) { - holder = move(other.holder); + holder = std::move(other.holder); type = other.type; + allow_type_rewrite = other.allow_type_rewrite; other.type = TypeNone; } return *this; } + // This copy/clone should only be called in the Compiler constructor. + // If this is called inside ::compile(), we invalidate any references we took higher in the stack. + // This should never happen. + Variant &operator=(const Variant &other) + { +#ifdef SPIRV_CROSS_COPY_CONSTRUCTOR_SANITIZE + abort(); +#endif + if (this != &other) + { + holder.reset(); + if (other.holder) + holder = other.holder->clone(); + type = other.type; + allow_type_rewrite = other.allow_type_rewrite; + } + return *this; + } + void set(std::unique_ptr val, uint32_t new_type) { holder = std::move(val); @@ -1166,14 +1235,17 @@ class Variant { return type; } + uint32_t get_id() const { return holder ? holder->self : 0; } + bool empty() const { return !holder; } + void reset() { holder.reset(); diff --git a/spirv_cpp.cpp b/spirv_cpp.cpp index 7d781e917..52575f9be 100644 --- a/spirv_cpp.cpp +++ b/spirv_cpp.cpp @@ -27,8 +27,8 @@ void CompilerCPP::emit_buffer_block(const SPIRVariable &var) auto &type = get(var.basetype); auto instance_name = to_name(var.self); - uint32_t descriptor_set = meta[var.self].decoration.set; - uint32_t binding = meta[var.self].decoration.binding; + 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); @@ -49,10 +49,10 @@ void CompilerCPP::emit_interface_block(const SPIRVariable &var) 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 = meta[var.self].decoration.location; + uint32_t location = ir.meta[var.self].decoration.location; string buffer_name; - auto flags = meta[type.self].decoration.decoration_flags; + auto flags = ir.meta[type.self].decoration.decoration_flags; if (flags.get(DecorationBlock)) { emit_block_struct(type); @@ -83,9 +83,9 @@ void CompilerCPP::emit_uniform(const SPIRVariable &var) auto &type = get(var.basetype); auto instance_name = to_name(var.self); - uint32_t descriptor_set = meta[var.self].decoration.set; - uint32_t binding = meta[var.self].decoration.binding; - uint32_t location = meta[var.self].decoration.location; + 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); @@ -114,7 +114,7 @@ void CompilerCPP::emit_push_constant_block(const SPIRVariable &var) add_resource_name(var.self); auto &type = get(var.basetype); - auto &flags = meta[var.self].decoration.decoration_flags; + 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."); @@ -145,14 +145,14 @@ void CompilerCPP::emit_resources() { // Output all basic struct types which are not Block or BufferBlock as these are declared inplace // when such variables are instantiated. - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeType) { auto &type = id.get(); if (type.basetype == SPIRType::Struct && type.array.empty() && !type.pointer && - (!meta[type.self].decoration.decoration_flags.get(DecorationBlock) && - !meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock))) + (!ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock) && + !ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock))) { emit_struct(type); } @@ -163,7 +163,7 @@ void CompilerCPP::emit_resources() begin_scope(); // Output UBOs and SSBOs - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { @@ -172,8 +172,8 @@ void CompilerCPP::emit_resources() if (var.storage != StorageClassFunction && type.pointer && type.storage == StorageClassUniform && !is_hidden_variable(var) && - (meta[type.self].decoration.decoration_flags.get(DecorationBlock) || - meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock))) + (ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock) || + ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock))) { emit_buffer_block(var); } @@ -181,7 +181,7 @@ void CompilerCPP::emit_resources() } // Output push constant blocks - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { @@ -196,7 +196,7 @@ void CompilerCPP::emit_resources() } // Output in/out interfaces. - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { @@ -213,7 +213,7 @@ void CompilerCPP::emit_resources() } // Output Uniform Constants (values, samplers, images, etc). - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { @@ -318,7 +318,7 @@ string CompilerCPP::compile() emit_header(); emit_resources(); - emit_function(get(entry_point), Bitset()); + emit_function(get(ir.default_entry_point), Bitset()); pass_count++; } while (force_recompile); @@ -376,7 +376,7 @@ void CompilerCPP::emit_c_linkage() void CompilerCPP::emit_function_prototype(SPIRFunction &func, const Bitset &) { - if (func.self != entry_point) + if (func.self != ir.default_entry_point) add_function_overload(func); local_variable_names = resource_names; @@ -387,7 +387,7 @@ void CompilerCPP::emit_function_prototype(SPIRFunction &func, const Bitset &) decl += type_to_glsl(type); decl += " "; - if (func.self == entry_point) + if (func.self == ir.default_entry_point) { decl += "main"; processing_entry_point = true; diff --git a/spirv_cpp.hpp b/spirv_cpp.hpp index a61dc8ec3..2560634e0 100644 --- a/spirv_cpp.hpp +++ b/spirv_cpp.hpp @@ -26,13 +26,23 @@ namespace spirv_cross class CompilerCPP : public CompilerGLSL { public: - CompilerCPP(std::vector spirv_) + explicit CompilerCPP(std::vector spirv_) : CompilerGLSL(move(spirv_)) { } - CompilerCPP(const uint32_t *ir, size_t word_count) - : CompilerGLSL(ir, word_count) + CompilerCPP(const uint32_t *ir_, size_t word_count) + : CompilerGLSL(ir_, word_count) + { + } + + explicit CompilerCPP(const ParsedIR &ir_) + : CompilerGLSL(ir_) + { + } + + explicit CompilerCPP(ParsedIR &&ir_) + : CompilerGLSL(std::move(ir_)) { } diff --git a/spirv_cross.cpp b/spirv_cross.cpp index 4082957cf..81002647f 100644 --- a/spirv_cross.cpp +++ b/spirv_cross.cpp @@ -17,6 +17,7 @@ #include "spirv_cross.hpp" #include "GLSL.std.450.h" #include "spirv_cfg.hpp" +#include "spirv_parser.hpp" #include #include #include @@ -25,69 +26,40 @@ using namespace std; using namespace spv; using namespace spirv_cross; -#define log(...) fprintf(stderr, __VA_ARGS__) - -static string ensure_valid_identifier(const string &name, bool member) +Compiler::Compiler(vector ir_) { - // Functions in glslangValidator are mangled with name( stuff. - // Normally, we would never see '(' in any legal identifiers, so just strip them out. - auto str = name.substr(0, name.find('(')); - - for (uint32_t i = 0; i < str.size(); i++) - { - auto &c = str[i]; - - if (member) - { - // _m variables are reserved by the internal implementation, - // otherwise, make sure the name is a valid identifier. - if (i == 0) - c = isalpha(c) ? c : '_'; - else if (i == 2 && str[0] == '_' && str[1] == 'm') - c = isalpha(c) ? c : '_'; - else - c = isalnum(c) ? c : '_'; - } - else - { - // _ variables are reserved by the internal implementation, - // otherwise, make sure the name is a valid identifier. - if (i == 0 || (str[0] == '_' && i == 1)) - c = isalpha(c) ? c : '_'; - else - c = isalnum(c) ? c : '_'; - } - } - return str; + Parser parser(move(ir_)); + parser.parse(); + set_ir(move(parser.get_parsed_ir())); } -Instruction::Instruction(const vector &spirv, uint32_t &index) +Compiler::Compiler(const uint32_t *ir_, size_t word_count) { - op = spirv[index] & 0xffff; - count = (spirv[index] >> 16) & 0xffff; - - if (count == 0) - SPIRV_CROSS_THROW("SPIR-V instructions cannot consume 0 words. Invalid SPIR-V file."); - - offset = index + 1; - length = count - 1; + Parser parser(ir_, word_count); + parser.parse(); + set_ir(move(parser.get_parsed_ir())); +} - index += count; +Compiler::Compiler(const ParsedIR &ir_) +{ + set_ir(ir_); +} - if (index > spirv.size()) - SPIRV_CROSS_THROW("SPIR-V instruction goes out of bounds."); +Compiler::Compiler(ParsedIR &&ir_) +{ + set_ir(move(ir_)); } -Compiler::Compiler(vector ir) - : spirv(move(ir)) +void Compiler::set_ir(ParsedIR &&ir_) { - parse(); + ir = move(ir_); + parse_fixup(); } -Compiler::Compiler(const uint32_t *ir, size_t word_count) - : spirv(ir, ir + word_count) +void Compiler::set_ir(const ParsedIR &ir_) { - parse(); + ir = ir_; + parse_fixup(); } string Compiler::compile() @@ -101,13 +73,13 @@ bool Compiler::variable_storage_is_aliased(const SPIRVariable &v) { auto &type = get(v.basetype); bool ssbo = v.storage == StorageClassStorageBuffer || - meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); + ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); bool image = type.basetype == SPIRType::Image; bool counter = type.basetype == SPIRType::AtomicCounter; bool is_restrict; if (ssbo) - is_restrict = get_buffer_block_flags(v).get(DecorationRestrict); + is_restrict = ir.get_buffer_block_flags(v).get(DecorationRestrict); else is_restrict = has_decoration(v.self, DecorationRestrict); @@ -186,7 +158,7 @@ bool Compiler::block_is_pure(const SPIRBlock &block) string Compiler::to_name(uint32_t id, bool allow_alias) const { - if (allow_alias && ids.at(id).get_type() == TypeType) + if (allow_alias && ir.ids.at(id).get_type() == TypeType) { // If this type is a simple alias, emit the // name of the original type instead. @@ -202,10 +174,10 @@ string Compiler::to_name(uint32_t id, bool allow_alias) const } } - if (meta[id].decoration.alias.empty()) + if (ir.meta[id].decoration.alias.empty()) return join("_", id); else - return meta.at(id).decoration.alias; + return ir.meta[id].decoration.alias; } bool Compiler::function_is_pure(const SPIRFunction &func) @@ -381,7 +353,7 @@ void Compiler::flush_all_active_variables() uint32_t Compiler::expression_type_id(uint32_t id) const { - switch (ids[id].get_type()) + switch (ir.ids[id].get_type()) { case TypeVariable: return get(id).basetype; @@ -431,7 +403,7 @@ bool Compiler::expression_is_lvalue(uint32_t id) const bool Compiler::is_immutable(uint32_t id) const { - if (ids[id].get_type() == TypeVariable) + if (ir.ids[id].get_type() == TypeVariable) { auto &var = get(id); @@ -439,12 +411,12 @@ bool Compiler::is_immutable(uint32_t id) const bool pointer_to_const = var.storage == StorageClassUniformConstant; return pointer_to_const || var.phi_variable || !expression_is_lvalue(id); } - else if (ids[id].get_type() == TypeAccessChain) + else if (ir.ids[id].get_type() == TypeAccessChain) return get(id).immutable; - else if (ids[id].get_type() == TypeExpression) + else if (ir.ids[id].get_type() == TypeExpression) return get(id).immutable; - else if (ids[id].get_type() == TypeConstant || ids[id].get_type() == TypeConstantOp || - ids[id].get_type() == TypeUndef) + else if (ir.ids[id].get_type() == TypeConstant || ir.ids[id].get_type() == TypeConstantOp || + ir.ids[id].get_type() == TypeUndef) return true; else return false; @@ -490,7 +462,7 @@ bool Compiler::is_hidden_variable(const SPIRVariable &var, bool include_builtins bool Compiler::is_builtin_type(const SPIRType &type) const { // We can have builtin structs as well. If one member of a struct is builtin, the struct must also be builtin. - for (auto &m : meta[type.self].members) + for (auto &m : ir.meta[type.self].members) if (m.builtin) return true; @@ -499,7 +471,7 @@ bool Compiler::is_builtin_type(const SPIRType &type) const bool Compiler::is_builtin_variable(const SPIRVariable &var) const { - if (var.compat_builtin || meta[var.self].decoration.builtin) + if (var.compat_builtin || ir.meta[var.self].decoration.builtin) return true; else return is_builtin_type(get(var.basetype)); @@ -507,7 +479,7 @@ bool Compiler::is_builtin_variable(const SPIRVariable &var) const bool Compiler::is_member_builtin(const SPIRType &type, uint32_t index, BuiltIn *builtin) const { - auto &memb = meta[type.self].members; + auto &memb = ir.meta[type.self].members; if (index < memb.size() && memb[index].builtin) { if (builtin) @@ -669,7 +641,7 @@ unordered_set Compiler::get_active_interface_variables() const // Traverse the call graph and find all interface variables which are in use. unordered_set variables; InterfaceVariableAccessHandler handler(*this, variables); - traverse_all_reachable_opcodes(get(entry_point), handler); + traverse_all_reachable_opcodes(get(ir.default_entry_point), handler); // If we needed to create one, we'll need it. if (dummy_sampler_id) @@ -688,7 +660,7 @@ ShaderResources Compiler::get_shader_resources(const unordered_set *ac { ShaderResources res; - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() != TypeVariable) continue; @@ -707,36 +679,36 @@ ShaderResources Compiler::get_shader_resources(const unordered_set *ac // Input if (var.storage == StorageClassInput && interface_variable_exists_in_entry_point(var.self)) { - if (meta[type.self].decoration.decoration_flags.get(DecorationBlock)) + if (ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock)) res.stage_inputs.push_back( { var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) }); else - res.stage_inputs.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias }); + res.stage_inputs.push_back({ var.self, var.basetype, type.self, ir.meta[var.self].decoration.alias }); } // Subpass inputs else if (var.storage == StorageClassUniformConstant && type.image.dim == DimSubpassData) { - res.subpass_inputs.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias }); + res.subpass_inputs.push_back({ var.self, var.basetype, type.self, ir.meta[var.self].decoration.alias }); } // Outputs else if (var.storage == StorageClassOutput && interface_variable_exists_in_entry_point(var.self)) { - if (meta[type.self].decoration.decoration_flags.get(DecorationBlock)) + if (ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock)) res.stage_outputs.push_back( { var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) }); else - res.stage_outputs.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias }); + res.stage_outputs.push_back({ var.self, var.basetype, type.self, ir.meta[var.self].decoration.alias }); } // UBOs else if (type.storage == StorageClassUniform && - (meta[type.self].decoration.decoration_flags.get(DecorationBlock))) + (ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock))) { res.uniform_buffers.push_back( { var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) }); } // Old way to declare SSBOs. else if (type.storage == StorageClassUniform && - (meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock))) + (ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock))) { res.storage_buffers.push_back( { var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self) }); @@ -752,81 +724,41 @@ ShaderResources Compiler::get_shader_resources(const unordered_set *ac { // There can only be one push constant block, but keep the vector in case this restriction is lifted // in the future. - res.push_constant_buffers.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias }); + res.push_constant_buffers.push_back( + { var.self, var.basetype, type.self, ir.meta[var.self].decoration.alias }); } // Images else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Image && type.image.sampled == 2) { - res.storage_images.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias }); + res.storage_images.push_back({ var.self, var.basetype, type.self, ir.meta[var.self].decoration.alias }); } // Separate images else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Image && type.image.sampled == 1) { - res.separate_images.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias }); + res.separate_images.push_back({ var.self, var.basetype, type.self, ir.meta[var.self].decoration.alias }); } // Separate samplers else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Sampler) { - res.separate_samplers.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias }); + res.separate_samplers.push_back({ var.self, var.basetype, type.self, ir.meta[var.self].decoration.alias }); } // Textures else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::SampledImage) { - res.sampled_images.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias }); + res.sampled_images.push_back({ var.self, var.basetype, type.self, ir.meta[var.self].decoration.alias }); } // Atomic counters else if (type.storage == StorageClassAtomicCounter) { - res.atomic_counters.push_back({ var.self, var.basetype, type.self, meta[var.self].decoration.alias }); + res.atomic_counters.push_back({ var.self, var.basetype, type.self, ir.meta[var.self].decoration.alias }); } } return res; } -static inline uint32_t swap_endian(uint32_t v) -{ - return ((v >> 24) & 0x000000ffu) | ((v >> 8) & 0x0000ff00u) | ((v << 8) & 0x00ff0000u) | ((v << 24) & 0xff000000u); -} - -static string extract_string(const vector &spirv, uint32_t offset) -{ - string ret; - for (uint32_t i = offset; i < spirv.size(); i++) - { - uint32_t w = spirv[i]; - - for (uint32_t j = 0; j < 4; j++, w >>= 8) - { - char c = w & 0xff; - if (c == '\0') - return ret; - ret += c; - } - } - - SPIRV_CROSS_THROW("String was not terminated before EOF"); -} - -static bool is_valid_spirv_version(uint32_t version) -{ - switch (version) - { - // Allow v99 since it tends to just work. - case 99: - case 0x10000: // SPIR-V 1.0 - case 0x10100: // SPIR-V 1.1 - case 0x10200: // SPIR-V 1.2 - case 0x10300: // SPIR-V 1.3 - return true; - - default: - return false; - } -} - bool Compiler::type_is_block_like(const SPIRType &type) const { if (type.basetype != SPIRType::Struct) @@ -850,7 +782,7 @@ void Compiler::fixup_type_alias() // Due to how some backends work, the "master" type of type_alias must be a block-like type if it exists. // FIXME: Multiple alias types which are both block-like will be awkward, for now, it's best to just drop the type // alias if the slave type is a block type. - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() != TypeType) continue; @@ -860,7 +792,7 @@ void Compiler::fixup_type_alias() if (type.type_alias && type_is_block_like(type)) { // Become the master. - for (auto &other_id : ids) + for (auto &other_id : ir.ids) { if (other_id.get_type() != TypeType) continue; @@ -877,7 +809,7 @@ void Compiler::fixup_type_alias() } } - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() != TypeType) continue; @@ -891,48 +823,19 @@ void Compiler::fixup_type_alias() } } -void Compiler::parse() +void Compiler::parse_fixup() { - auto len = spirv.size(); - if (len < 5) - SPIRV_CROSS_THROW("SPIRV file too small."); - - auto s = spirv.data(); - - // Endian-swap if we need to. - if (s[0] == swap_endian(MagicNumber)) - transform(begin(spirv), end(spirv), begin(spirv), [](uint32_t c) { return swap_endian(c); }); - - if (s[0] != MagicNumber || !is_valid_spirv_version(s[1])) - SPIRV_CROSS_THROW("Invalid SPIRV format."); - - uint32_t bound = s[3]; - ids.resize(bound); - meta.resize(bound); - - uint32_t offset = 5; - while (offset < len) - inst.emplace_back(spirv, offset); - - for (auto &i : inst) - parse(i); - - if (current_function) - SPIRV_CROSS_THROW("Function was not terminated."); - if (current_block) - SPIRV_CROSS_THROW("Block was not terminated."); - // Figure out specialization constants for work group sizes. - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeConstant) { auto &c = id.get(); - if (meta[c.self].decoration.builtin && meta[c.self].decoration.builtin_type == BuiltInWorkgroupSize) + if (ir.meta[c.self].decoration.builtin && ir.meta[c.self].decoration.builtin_type == BuiltInWorkgroupSize) { // In current SPIR-V, there can be just one constant like this. // All entry points will receive the constant value. - for (auto &entry : entry_points) + for (auto &entry : ir.entry_points) { entry.second.workgroup_size.constant = c.self; entry.second.workgroup_size.x = c.scalar(0, 0); @@ -941,6 +844,15 @@ void Compiler::parse() } } } + else if (id.get_type() == TypeVariable) + { + auto &var = id.get(); + if (var.storage == StorageClassPrivate || var.storage == StorageClassWorkgroup || + var.storage == StorageClassOutput) + global_variables.push_back(var.self); + if (variable_storage_is_aliased(var)) + aliased_variables.push_back(var.self); + } } fixup_type_alias(); @@ -950,7 +862,7 @@ void Compiler::flatten_interface_block(uint32_t id) { auto &var = get(id); auto &type = get(var.basetype); - auto &flags = meta.at(type.self).decoration.decoration_flags; + auto &flags = ir.meta.at(type.self).decoration.decoration_flags; if (!type.array.empty()) SPIRV_CROSS_THROW("Type is array of UBOs."); @@ -972,1378 +884,252 @@ void Compiler::flatten_interface_block(uint32_t id) if (mtype.basetype == SPIRType::Struct) SPIRV_CROSS_THROW("Member type cannot be struct."); - // Inherit variable name from interface block name. - meta.at(var.self).decoration.alias = meta.at(type.self).decoration.alias; - - auto storage = var.storage; - if (storage == StorageClassUniform) - storage = StorageClassUniformConstant; - - // Change type definition in-place into an array instead. - // Access chains will still work as-is. - uint32_t array_size = uint32_t(type.member_types.size()); - type = mtype; - type.array.push_back(array_size); - type.pointer = true; - type.storage = storage; - var.storage = storage; -} - -void Compiler::update_name_cache(unordered_set &cache, string &name) -{ - if (name.empty()) - return; - - if (cache.find(name) == end(cache)) - { - cache.insert(name); - return; - } - - uint32_t counter = 0; - auto tmpname = name; - - bool use_linked_underscore = true; - - if (tmpname == "_") - { - // We cannot just append numbers, as we will end up creating internally reserved names. - // Make it like _0_ instead. - tmpname += "0"; - } - else if (tmpname.back() == '_') - { - // The last_character is an underscore, so we don't need to link in underscore. - // This would violate double underscore rules. - use_linked_underscore = false; - } - - // If there is a collision (very rare), - // keep tacking on extra identifier until it's unique. - do - { - counter++; - name = tmpname + (use_linked_underscore ? "_" : "") + convert_to_string(counter); - } while (cache.find(name) != end(cache)); - cache.insert(name); -} - -void Compiler::set_name(uint32_t id, const std::string &name) -{ - auto &str = meta.at(id).decoration.alias; - str.clear(); - - if (name.empty()) - return; - - // glslang uses identifiers to pass along meaningful information - // about HLSL reflection. - // FIXME: This should be deprecated eventually. - auto &m = meta.at(id); - if (source.hlsl && name.size() >= 6 && name.find("@count") == name.size() - 6) - { - m.hlsl_magic_counter_buffer_candidate = true; - m.hlsl_magic_counter_buffer_name = name.substr(0, name.find("@count")); - } - else - { - m.hlsl_magic_counter_buffer_candidate = false; - m.hlsl_magic_counter_buffer_name.clear(); - } - - // Reserved for temporaries. - if (name[0] == '_' && name.size() >= 2 && isdigit(name[1])) - return; - - str = ensure_valid_identifier(name, false); -} - -const SPIRType &Compiler::get_type(uint32_t id) const -{ - return get(id); -} - -const SPIRType &Compiler::get_type_from_variable(uint32_t id) const -{ - return get(get(id).basetype); -} - -uint32_t Compiler::get_non_pointer_type_id(uint32_t type_id) const -{ - auto *p_type = &get(type_id); - while (p_type->pointer) - { - assert(p_type->parent_type); - type_id = p_type->parent_type; - p_type = &get(type_id); - } - return type_id; -} - -const SPIRType &Compiler::get_non_pointer_type(const SPIRType &type) const -{ - auto *p_type = &type; - while (p_type->pointer) - { - assert(p_type->parent_type); - p_type = &get(p_type->parent_type); - } - return *p_type; -} - -const SPIRType &Compiler::get_non_pointer_type(uint32_t type_id) const -{ - return get_non_pointer_type(get(type_id)); -} - -bool Compiler::is_sampled_image_type(const SPIRType &type) -{ - return (type.basetype == SPIRType::Image || type.basetype == SPIRType::SampledImage) && type.image.sampled == 1 && - type.image.dim != DimBuffer; -} - -void Compiler::set_member_decoration_string(uint32_t id, uint32_t index, spv::Decoration decoration, - const std::string &argument) -{ - meta.at(id).members.resize(max(meta[id].members.size(), size_t(index) + 1)); - auto &dec = meta.at(id).members[index]; - dec.decoration_flags.set(decoration); - - switch (decoration) - { - case DecorationHlslSemanticGOOGLE: - dec.hlsl_semantic = argument; - break; - - default: - break; - } -} - -void Compiler::set_member_decoration(uint32_t id, uint32_t index, Decoration decoration, uint32_t argument) -{ - meta.at(id).members.resize(max(meta[id].members.size(), size_t(index) + 1)); - auto &dec = meta.at(id).members[index]; - dec.decoration_flags.set(decoration); - - switch (decoration) - { - case DecorationBuiltIn: - dec.builtin = true; - dec.builtin_type = static_cast(argument); - break; - - case DecorationLocation: - dec.location = argument; - break; - - case DecorationComponent: - dec.component = argument; - break; - - case DecorationBinding: - dec.binding = argument; - break; - - case DecorationOffset: - dec.offset = argument; - break; - - case DecorationSpecId: - dec.spec_id = argument; - break; - - case DecorationMatrixStride: - dec.matrix_stride = argument; - break; - - case DecorationIndex: - dec.index = argument; - break; - - default: - break; - } -} - -void Compiler::set_member_name(uint32_t id, uint32_t index, const std::string &name) -{ - meta.at(id).members.resize(max(meta[id].members.size(), size_t(index) + 1)); - - auto &str = meta.at(id).members[index].alias; - str.clear(); - if (name.empty()) - return; - - // Reserved for unnamed members. - if (name[0] == '_' && name.size() >= 3 && name[1] == 'm' && isdigit(name[2])) - return; - - str = ensure_valid_identifier(name, true); -} - -const std::string &Compiler::get_member_name(uint32_t id, uint32_t index) const -{ - auto &m = meta.at(id); - if (index >= m.members.size()) - { - static string empty; - return empty; - } - - return m.members[index].alias; -} - -void Compiler::set_member_qualified_name(uint32_t type_id, uint32_t index, const std::string &name) -{ - meta.at(type_id).members.resize(max(meta[type_id].members.size(), size_t(index) + 1)); - meta.at(type_id).members[index].qualified_alias = name; -} - -const std::string &Compiler::get_member_qualified_name(uint32_t type_id, uint32_t index) const -{ - const static string empty; - - auto &m = meta.at(type_id); - if (index < m.members.size()) - return m.members[index].qualified_alias; - else - return empty; -} - -uint32_t Compiler::get_member_decoration(uint32_t id, uint32_t index, Decoration decoration) const -{ - auto &m = meta.at(id); - if (index >= m.members.size()) - return 0; - - auto &dec = m.members[index]; - if (!dec.decoration_flags.get(decoration)) - return 0; - - switch (decoration) - { - case DecorationBuiltIn: - return dec.builtin_type; - case DecorationLocation: - return dec.location; - case DecorationComponent: - return dec.component; - case DecorationBinding: - return dec.binding; - case DecorationOffset: - return dec.offset; - case DecorationSpecId: - return dec.spec_id; - case DecorationIndex: - return dec.index; - default: - return 1; - } -} - -uint64_t Compiler::get_member_decoration_mask(uint32_t id, uint32_t index) const -{ - return get_member_decoration_bitset(id, index).get_lower(); -} - -const Bitset &Compiler::get_member_decoration_bitset(uint32_t id, uint32_t index) const -{ - auto &m = meta.at(id); - if (index >= m.members.size()) - { - static const Bitset cleared = {}; - return cleared; - } - - return m.members[index].decoration_flags; -} - -bool Compiler::has_member_decoration(uint32_t id, uint32_t index, Decoration decoration) const -{ - return get_member_decoration_bitset(id, index).get(decoration); -} - -void Compiler::unset_member_decoration(uint32_t id, uint32_t index, Decoration decoration) -{ - auto &m = meta.at(id); - if (index >= m.members.size()) - return; - - auto &dec = m.members[index]; - - dec.decoration_flags.clear(decoration); - switch (decoration) - { - case DecorationBuiltIn: - dec.builtin = false; - break; - - case DecorationLocation: - dec.location = 0; - break; - - case DecorationComponent: - dec.component = 0; - break; - - case DecorationOffset: - dec.offset = 0; - break; - - case DecorationSpecId: - dec.spec_id = 0; - break; - - case DecorationHlslSemanticGOOGLE: - dec.hlsl_semantic.clear(); - break; - - default: - break; - } -} - -void Compiler::set_decoration_string(uint32_t id, spv::Decoration decoration, const std::string &argument) -{ - auto &dec = meta.at(id).decoration; - dec.decoration_flags.set(decoration); - - switch (decoration) - { - case DecorationHlslSemanticGOOGLE: - dec.hlsl_semantic = argument; - break; - - default: - break; - } -} - -void Compiler::set_decoration(uint32_t id, Decoration decoration, uint32_t argument) -{ - auto &dec = meta.at(id).decoration; - dec.decoration_flags.set(decoration); - - switch (decoration) - { - case DecorationBuiltIn: - dec.builtin = true; - dec.builtin_type = static_cast(argument); - break; - - case DecorationLocation: - dec.location = argument; - break; - - case DecorationComponent: - dec.component = argument; - break; - - case DecorationOffset: - dec.offset = argument; - break; - - case DecorationArrayStride: - dec.array_stride = argument; - break; - - case DecorationMatrixStride: - dec.matrix_stride = argument; - break; - - case DecorationBinding: - dec.binding = argument; - break; - - case DecorationDescriptorSet: - dec.set = argument; - break; - - case DecorationInputAttachmentIndex: - dec.input_attachment = argument; - break; - - case DecorationSpecId: - dec.spec_id = argument; - break; - - case DecorationIndex: - dec.index = argument; - break; - - case DecorationHlslCounterBufferGOOGLE: - meta.at(id).hlsl_magic_counter_buffer = argument; - meta.at(argument).hlsl_is_magic_counter_buffer = true; - break; - - default: - break; - } -} - -StorageClass Compiler::get_storage_class(uint32_t id) const -{ - return get(id).storage; -} - -const std::string &Compiler::get_name(uint32_t id) const -{ - return meta.at(id).decoration.alias; -} - -const std::string Compiler::get_fallback_name(uint32_t id) const -{ - return join("_", id); -} - -const std::string Compiler::get_block_fallback_name(uint32_t id) const -{ - auto &var = get(id); - if (get_name(id).empty()) - return join("_", get(var.basetype).self, "_", id); - else - return get_name(id); -} - -uint64_t Compiler::get_decoration_mask(uint32_t id) const -{ - return get_decoration_bitset(id).get_lower(); -} - -const Bitset &Compiler::get_decoration_bitset(uint32_t id) const -{ - auto &dec = meta.at(id).decoration; - return dec.decoration_flags; -} - -bool Compiler::has_decoration(uint32_t id, Decoration decoration) const -{ - return get_decoration_bitset(id).get(decoration); -} - -const string &Compiler::get_decoration_string(uint32_t id, spv::Decoration decoration) const -{ - auto &dec = meta.at(id).decoration; - static const string empty; - - if (!dec.decoration_flags.get(decoration)) - return empty; - - switch (decoration) - { - case DecorationHlslSemanticGOOGLE: - return dec.hlsl_semantic; - - default: - return empty; - } -} - -uint32_t Compiler::get_decoration(uint32_t id, Decoration decoration) const -{ - auto &dec = meta.at(id).decoration; - if (!dec.decoration_flags.get(decoration)) - return 0; - - switch (decoration) - { - case DecorationBuiltIn: - return dec.builtin_type; - case DecorationLocation: - return dec.location; - case DecorationComponent: - return dec.component; - case DecorationOffset: - return dec.offset; - case DecorationBinding: - return dec.binding; - case DecorationDescriptorSet: - return dec.set; - case DecorationInputAttachmentIndex: - return dec.input_attachment; - case DecorationSpecId: - return dec.spec_id; - case DecorationArrayStride: - return dec.array_stride; - case DecorationMatrixStride: - return dec.matrix_stride; - case DecorationIndex: - return dec.index; - default: - return 1; - } -} - -void Compiler::unset_decoration(uint32_t id, Decoration decoration) -{ - auto &dec = meta.at(id).decoration; - dec.decoration_flags.clear(decoration); - switch (decoration) - { - case DecorationBuiltIn: - dec.builtin = false; - break; - - case DecorationLocation: - dec.location = 0; - break; - - case DecorationComponent: - dec.component = 0; - break; - - case DecorationOffset: - dec.offset = 0; - break; - - case DecorationBinding: - dec.binding = 0; - break; - - case DecorationDescriptorSet: - dec.set = 0; - break; - - case DecorationInputAttachmentIndex: - dec.input_attachment = 0; - break; - - case DecorationSpecId: - dec.spec_id = 0; - break; - - case DecorationHlslSemanticGOOGLE: - dec.hlsl_semantic.clear(); - break; - - case DecorationHlslCounterBufferGOOGLE: - { - auto &counter = meta.at(id).hlsl_magic_counter_buffer; - if (counter) - { - meta.at(counter).hlsl_is_magic_counter_buffer = false; - counter = 0; - } - break; - } - - default: - break; - } -} - -bool Compiler::get_binary_offset_for_decoration(uint32_t id, spv::Decoration decoration, uint32_t &word_offset) const -{ - auto &word_offsets = meta.at(id).decoration_word_offset; - auto itr = word_offsets.find(decoration); - if (itr == end(word_offsets)) - return false; - - word_offset = itr->second; - return true; -} - -void Compiler::parse(const Instruction &instruction) -{ - auto ops = stream(instruction); - auto op = static_cast(instruction.op); - uint32_t length = instruction.length; - - switch (op) - { - case OpMemoryModel: - case OpSourceExtension: - case OpNop: - case OpLine: - case OpNoLine: - case OpString: - break; - - case OpSource: - { - auto lang = static_cast(ops[0]); - switch (lang) - { - case SourceLanguageESSL: - source.es = true; - source.version = ops[1]; - source.known = true; - source.hlsl = false; - break; - - case SourceLanguageGLSL: - source.es = false; - source.version = ops[1]; - source.known = true; - source.hlsl = false; - break; - - case SourceLanguageHLSL: - // For purposes of cross-compiling, this is GLSL 450. - source.es = false; - source.version = 450; - source.known = true; - source.hlsl = true; - break; - - default: - source.known = false; - break; - } - break; - } - - case OpUndef: - { - uint32_t result_type = ops[0]; - uint32_t id = ops[1]; - set(id, result_type); - break; - } - - case OpCapability: - { - uint32_t cap = ops[0]; - if (cap == CapabilityKernel) - SPIRV_CROSS_THROW("Kernel capability not supported."); - - declared_capabilities.push_back(static_cast(ops[0])); - break; - } - - case OpExtension: - { - auto ext = extract_string(spirv, instruction.offset); - declared_extensions.push_back(move(ext)); - break; - } - - case OpExtInstImport: - { - uint32_t id = ops[0]; - auto ext = extract_string(spirv, instruction.offset + 1); - if (ext == "GLSL.std.450") - set(id, SPIRExtension::GLSL); - else if (ext == "SPV_AMD_shader_ballot") - set(id, SPIRExtension::SPV_AMD_shader_ballot); - else if (ext == "SPV_AMD_shader_explicit_vertex_parameter") - set(id, SPIRExtension::SPV_AMD_shader_explicit_vertex_parameter); - else if (ext == "SPV_AMD_shader_trinary_minmax") - set(id, SPIRExtension::SPV_AMD_shader_trinary_minmax); - else if (ext == "SPV_AMD_gcn_shader") - set(id, SPIRExtension::SPV_AMD_gcn_shader); - else - set(id, SPIRExtension::Unsupported); - - // Other SPIR-V extensions currently not supported. - - break; - } - - case OpEntryPoint: - { - auto itr = - entry_points.insert(make_pair(ops[1], SPIREntryPoint(ops[1], static_cast(ops[0]), - extract_string(spirv, instruction.offset + 2)))); - auto &e = itr.first->second; - - // Strings need nul-terminator and consume the whole word. - uint32_t strlen_words = uint32_t((e.name.size() + 1 + 3) >> 2); - e.interface_variables.insert(end(e.interface_variables), ops + strlen_words + 2, ops + instruction.length); - - // Set the name of the entry point in case OpName is not provided later - set_name(ops[1], e.name); - - // If we don't have an entry, make the first one our "default". - if (!entry_point) - entry_point = ops[1]; - break; - } - - case OpExecutionMode: - { - auto &execution = entry_points[ops[0]]; - auto mode = static_cast(ops[1]); - execution.flags.set(mode); - - switch (mode) - { - case ExecutionModeInvocations: - execution.invocations = ops[2]; - break; - - case ExecutionModeLocalSize: - execution.workgroup_size.x = ops[2]; - execution.workgroup_size.y = ops[3]; - execution.workgroup_size.z = ops[4]; - break; - - case ExecutionModeOutputVertices: - execution.output_vertices = ops[2]; - break; - - default: - break; - } - break; - } - - case OpName: - { - uint32_t id = ops[0]; - set_name(id, extract_string(spirv, instruction.offset + 1)); - break; - } - - case OpMemberName: - { - uint32_t id = ops[0]; - uint32_t member = ops[1]; - set_member_name(id, member, extract_string(spirv, instruction.offset + 2)); - break; - } - - case OpDecorate: - case OpDecorateId: - { - uint32_t id = ops[0]; - - auto decoration = static_cast(ops[1]); - if (length >= 3) - { - meta[id].decoration_word_offset[decoration] = uint32_t(&ops[2] - spirv.data()); - set_decoration(id, decoration, ops[2]); - } - else - set_decoration(id, decoration); - - break; - } - - case OpDecorateStringGOOGLE: - { - uint32_t id = ops[0]; - auto decoration = static_cast(ops[1]); - set_decoration_string(id, decoration, extract_string(spirv, instruction.offset + 2)); - break; - } - - case OpMemberDecorate: - { - uint32_t id = ops[0]; - uint32_t member = ops[1]; - auto decoration = static_cast(ops[2]); - if (length >= 4) - set_member_decoration(id, member, decoration, ops[3]); - else - set_member_decoration(id, member, decoration); - break; - } - - case OpMemberDecorateStringGOOGLE: - { - uint32_t id = ops[0]; - uint32_t member = ops[1]; - auto decoration = static_cast(ops[2]); - set_member_decoration_string(id, member, decoration, extract_string(spirv, instruction.offset + 3)); - break; - } - - // Build up basic types. - case OpTypeVoid: - { - uint32_t id = ops[0]; - auto &type = set(id); - type.basetype = SPIRType::Void; - break; - } - - case OpTypeBool: - { - uint32_t id = ops[0]; - auto &type = set(id); - type.basetype = SPIRType::Boolean; - type.width = 1; - break; - } - - case OpTypeFloat: - { - uint32_t id = ops[0]; - uint32_t width = ops[1]; - auto &type = set(id); - if (width == 64) - type.basetype = SPIRType::Double; - else if (width == 32) - type.basetype = SPIRType::Float; - else if (width == 16) - type.basetype = SPIRType::Half; - else - SPIRV_CROSS_THROW("Unrecognized bit-width of floating point type."); - type.width = width; - break; - } - - case OpTypeInt: - { - uint32_t id = ops[0]; - uint32_t width = ops[1]; - auto &type = set(id); - type.basetype = - ops[2] ? (width > 32 ? SPIRType::Int64 : SPIRType::Int) : (width > 32 ? SPIRType::UInt64 : SPIRType::UInt); - type.width = width; - break; - } - - // Build composite types by "inheriting". - // NOTE: The self member is also copied! For pointers and array modifiers this is a good thing - // since we can refer to decorations on pointee classes which is needed for UBO/SSBO, I/O blocks in geometry/tess etc. - case OpTypeVector: - { - uint32_t id = ops[0]; - uint32_t vecsize = ops[2]; - - auto &base = get(ops[1]); - auto &vecbase = set(id); - - vecbase = base; - vecbase.vecsize = vecsize; - vecbase.self = id; - vecbase.parent_type = ops[1]; - break; - } - - case OpTypeMatrix: - { - uint32_t id = ops[0]; - uint32_t colcount = ops[2]; - - auto &base = get(ops[1]); - auto &matrixbase = set(id); - - matrixbase = base; - matrixbase.columns = colcount; - matrixbase.self = id; - matrixbase.parent_type = ops[1]; - break; - } - - case OpTypeArray: - { - uint32_t id = ops[0]; - auto &arraybase = set(id); - - uint32_t tid = ops[1]; - auto &base = get(tid); - - arraybase = base; - arraybase.parent_type = tid; - - uint32_t cid = ops[2]; - mark_used_as_array_length(cid); - auto *c = maybe_get(cid); - bool literal = c && !c->specialization; - - arraybase.array_size_literal.push_back(literal); - arraybase.array.push_back(literal ? c->scalar() : cid); - // Do NOT set arraybase.self! - break; - } - - case OpTypeRuntimeArray: - { - uint32_t id = ops[0]; - - auto &base = get(ops[1]); - auto &arraybase = set(id); - - arraybase = base; - arraybase.array.push_back(0); - arraybase.array_size_literal.push_back(true); - arraybase.parent_type = ops[1]; - // Do NOT set arraybase.self! - break; - } - - case OpTypeImage: - { - uint32_t id = ops[0]; - auto &type = set(id); - type.basetype = SPIRType::Image; - type.image.type = ops[1]; - type.image.dim = static_cast(ops[2]); - type.image.depth = ops[3] == 1; - type.image.arrayed = ops[4] != 0; - type.image.ms = ops[5] != 0; - type.image.sampled = ops[6]; - type.image.format = static_cast(ops[7]); - type.image.access = (length >= 9) ? static_cast(ops[8]) : AccessQualifierMax; - - if (type.image.sampled == 0) - SPIRV_CROSS_THROW("OpTypeImage Sampled parameter must not be zero."); - - break; - } - - case OpTypeSampledImage: - { - uint32_t id = ops[0]; - uint32_t imagetype = ops[1]; - auto &type = set(id); - type = get(imagetype); - type.basetype = SPIRType::SampledImage; - type.self = id; - break; - } - - case OpTypeSampler: - { - uint32_t id = ops[0]; - auto &type = set(id); - type.basetype = SPIRType::Sampler; - break; - } - - case OpTypePointer: - { - uint32_t id = ops[0]; - - auto &base = get(ops[2]); - auto &ptrbase = set(id); - - ptrbase = base; - if (ptrbase.pointer) - SPIRV_CROSS_THROW("Cannot make pointer-to-pointer type."); - ptrbase.pointer = true; - ptrbase.storage = static_cast(ops[1]); - - if (ptrbase.storage == StorageClassAtomicCounter) - ptrbase.basetype = SPIRType::AtomicCounter; - - ptrbase.parent_type = ops[2]; - - // Do NOT set ptrbase.self! - break; - } - - case OpTypeStruct: - { - uint32_t id = ops[0]; - auto &type = set(id); - type.basetype = SPIRType::Struct; - for (uint32_t i = 1; i < length; i++) - type.member_types.push_back(ops[i]); - - // Check if we have seen this struct type before, with just different - // decorations. - // - // Add workaround for issue #17 as well by looking at OpName for the struct - // types, which we shouldn't normally do. - // We should not normally have to consider type aliases like this to begin with - // however ... glslang issues #304, #307 cover this. - - // For stripped names, never consider struct type aliasing. - // We risk declaring the same struct multiple times, but type-punning is not allowed - // so this is safe. - bool consider_aliasing = !get_name(type.self).empty(); - if (consider_aliasing) - { - for (auto &other : global_struct_cache) - { - if (get_name(type.self) == get_name(other) && - types_are_logically_equivalent(type, get(other))) - { - type.type_alias = other; - break; - } - } - - if (type.type_alias == 0) - global_struct_cache.push_back(id); - } - break; - } - - case OpTypeFunction: - { - uint32_t id = ops[0]; - uint32_t ret = ops[1]; - - auto &func = set(id, ret); - for (uint32_t i = 2; i < length; i++) - func.parameter_types.push_back(ops[i]); - break; - } - - // Variable declaration - // All variables are essentially pointers with a storage qualifier. - case OpVariable: - { - uint32_t type = ops[0]; - uint32_t id = ops[1]; - auto storage = static_cast(ops[2]); - uint32_t initializer = length == 4 ? ops[3] : 0; - - if (storage == StorageClassFunction) - { - if (!current_function) - SPIRV_CROSS_THROW("No function currently in scope"); - current_function->add_local_variable(id); - } - else if (storage == StorageClassPrivate || storage == StorageClassWorkgroup || storage == StorageClassOutput) - { - global_variables.push_back(id); - } - - auto &var = set(id, type, storage, initializer); - - // hlsl based shaders don't have those decorations. force them and then reset when reading/writing images - auto &ttype = get(type); - if (ttype.basetype == SPIRType::BaseType::Image) - { - set_decoration(id, DecorationNonWritable); - set_decoration(id, DecorationNonReadable); - } - - if (variable_storage_is_aliased(var)) - aliased_variables.push_back(var.self); - - break; - } - - // OpPhi - // OpPhi is a fairly magical opcode. - // It selects temporary variables based on which parent block we *came from*. - // In high-level languages we can "de-SSA" by creating a function local, and flush out temporaries to this function-local - // variable to emulate SSA Phi. - case OpPhi: - { - if (!current_function) - SPIRV_CROSS_THROW("No function currently in scope"); - if (!current_block) - SPIRV_CROSS_THROW("No block currently in scope"); - - uint32_t result_type = ops[0]; - uint32_t id = ops[1]; - - // Instead of a temporary, create a new function-wide temporary with this ID instead. - auto &var = set(id, result_type, spv::StorageClassFunction); - var.phi_variable = true; - - current_function->add_local_variable(id); - - for (uint32_t i = 2; i + 2 <= length; i += 2) - current_block->phi_variables.push_back({ ops[i], ops[i + 1], id }); - break; - } - - // Constants - case OpSpecConstant: - case OpConstant: - { - uint32_t id = ops[1]; - auto &type = get(ops[0]); - - if (type.width > 32) - set(id, ops[0], ops[2] | (uint64_t(ops[3]) << 32), op == OpSpecConstant); - else - set(id, ops[0], ops[2], op == OpSpecConstant); - break; - } - - case OpSpecConstantFalse: - case OpConstantFalse: - { - uint32_t id = ops[1]; - set(id, ops[0], uint32_t(0), op == OpSpecConstantFalse); - break; - } - - case OpSpecConstantTrue: - case OpConstantTrue: - { - uint32_t id = ops[1]; - set(id, ops[0], uint32_t(1), op == OpSpecConstantTrue); - break; - } - - case OpConstantNull: - { - uint32_t id = ops[1]; - uint32_t type = ops[0]; - make_constant_null(id, type); - break; - } - - case OpSpecConstantComposite: - case OpConstantComposite: - { - uint32_t id = ops[1]; - uint32_t type = ops[0]; - - auto &ctype = get(type); - - // We can have constants which are structs and arrays. - // In this case, our SPIRConstant will be a list of other SPIRConstant ids which we - // can refer to. - if (ctype.basetype == SPIRType::Struct || !ctype.array.empty()) - { - set(id, type, ops + 2, length - 2, op == OpSpecConstantComposite); - } - else - { - uint32_t elements = length - 2; - if (elements > 4) - SPIRV_CROSS_THROW("OpConstantComposite only supports 1, 2, 3 and 4 elements."); - - SPIRConstant remapped_constant_ops[4]; - const SPIRConstant *c[4]; - for (uint32_t i = 0; i < elements; i++) - { - // Specialization constants operations can also be part of this. - // We do not know their value, so any attempt to query SPIRConstant later - // will fail. We can only propagate the ID of the expression and use to_expression on it. - auto *constant_op = maybe_get(ops[2 + i]); - if (constant_op) - { - if (op == OpConstantComposite) - SPIRV_CROSS_THROW("Specialization constant operation used in OpConstantComposite."); - - remapped_constant_ops[i].make_null(get(constant_op->basetype)); - remapped_constant_ops[i].self = constant_op->self; - remapped_constant_ops[i].constant_type = constant_op->basetype; - remapped_constant_ops[i].specialization = true; - c[i] = &remapped_constant_ops[i]; - } - else - c[i] = &get(ops[2 + i]); - } - set(id, type, c, elements, op == OpSpecConstantComposite); - } - break; - } - - // Functions - case OpFunction: - { - uint32_t res = ops[0]; - uint32_t id = ops[1]; - // Control - uint32_t type = ops[3]; + // Inherit variable name from interface block name. + ir.meta.at(var.self).decoration.alias = ir.meta.at(type.self).decoration.alias; - if (current_function) - SPIRV_CROSS_THROW("Must end a function before starting a new one!"); + auto storage = var.storage; + if (storage == StorageClassUniform) + storage = StorageClassUniformConstant; - current_function = &set(id, res, type); - break; - } + // Change type definition in-place into an array instead. + // Access chains will still work as-is. + uint32_t array_size = uint32_t(type.member_types.size()); + type = mtype; + type.array.push_back(array_size); + type.pointer = true; + type.storage = storage; + var.storage = storage; +} - case OpFunctionParameter: +void Compiler::update_name_cache(unordered_set &cache, string &name) +{ + if (name.empty()) + return; + + if (cache.find(name) == end(cache)) { - uint32_t type = ops[0]; - uint32_t id = ops[1]; + cache.insert(name); + return; + } - if (!current_function) - SPIRV_CROSS_THROW("Must be in a function!"); + uint32_t counter = 0; + auto tmpname = name; - current_function->add_parameter(type, id); - set(id, type, StorageClassFunction); - break; - } + bool use_linked_underscore = true; - case OpFunctionEnd: + if (tmpname == "_") { - if (current_block) - { - // Very specific error message, but seems to come up quite often. - SPIRV_CROSS_THROW( - "Cannot end a function before ending the current block.\n" - "Likely cause: If this SPIR-V was created from glslang HLSL, make sure the entry point is valid."); - } - current_function = nullptr; - break; + // We cannot just append numbers, as we will end up creating internally reserved names. + // Make it like _0_ instead. + tmpname += "0"; + } + else if (tmpname.back() == '_') + { + // The last_character is an underscore, so we don't need to link in underscore. + // This would violate double underscore rules. + use_linked_underscore = false; } - // Blocks - case OpLabel: + // If there is a collision (very rare), + // keep tacking on extra identifier until it's unique. + do { - // OpLabel always starts a block. - if (!current_function) - SPIRV_CROSS_THROW("Blocks cannot exist outside functions!"); + counter++; + name = tmpname + (use_linked_underscore ? "_" : "") + convert_to_string(counter); + } while (cache.find(name) != end(cache)); + cache.insert(name); +} - uint32_t id = ops[0]; +void Compiler::set_name(uint32_t id, const std::string &name) +{ + ir.set_name(id, name); +} - current_function->blocks.push_back(id); - if (!current_function->entry_block) - current_function->entry_block = id; +const SPIRType &Compiler::get_type(uint32_t id) const +{ + return get(id); +} - if (current_block) - SPIRV_CROSS_THROW("Cannot start a block before ending the current block."); +const SPIRType &Compiler::get_type_from_variable(uint32_t id) const +{ + return get(get(id).basetype); +} - current_block = &set(id); - break; +uint32_t Compiler::get_non_pointer_type_id(uint32_t type_id) const +{ + auto *p_type = &get(type_id); + while (p_type->pointer) + { + assert(p_type->parent_type); + type_id = p_type->parent_type; + p_type = &get(type_id); } + return type_id; +} - // Branch instructions end blocks. - case OpBranch: +const SPIRType &Compiler::get_non_pointer_type(const SPIRType &type) const +{ + auto *p_type = &type; + while (p_type->pointer) { - if (!current_block) - SPIRV_CROSS_THROW("Trying to end a non-existing block."); - - uint32_t target = ops[0]; - current_block->terminator = SPIRBlock::Direct; - current_block->next_block = target; - current_block = nullptr; - break; + assert(p_type->parent_type); + p_type = &get(p_type->parent_type); } + return *p_type; +} - case OpBranchConditional: - { - if (!current_block) - SPIRV_CROSS_THROW("Trying to end a non-existing block."); +const SPIRType &Compiler::get_non_pointer_type(uint32_t type_id) const +{ + return get_non_pointer_type(get(type_id)); +} - current_block->condition = ops[0]; - current_block->true_block = ops[1]; - current_block->false_block = ops[2]; +bool Compiler::is_sampled_image_type(const SPIRType &type) +{ + return (type.basetype == SPIRType::Image || type.basetype == SPIRType::SampledImage) && type.image.sampled == 1 && + type.image.dim != DimBuffer; +} - current_block->terminator = SPIRBlock::Select; - current_block = nullptr; - break; - } +void Compiler::set_member_decoration_string(uint32_t id, uint32_t index, spv::Decoration decoration, + const std::string &argument) +{ + ir.set_member_decoration_string(id, index, decoration, argument); +} - case OpSwitch: - { - if (!current_block) - SPIRV_CROSS_THROW("Trying to end a non-existing block."); +void Compiler::set_member_decoration(uint32_t id, uint32_t index, Decoration decoration, uint32_t argument) +{ + ir.set_member_decoration(id, index, decoration, argument); +} - if (current_block->merge == SPIRBlock::MergeNone) - SPIRV_CROSS_THROW("Switch statement is not structured"); +void Compiler::set_member_name(uint32_t id, uint32_t index, const std::string &name) +{ + ir.set_member_name(id, index, name); +} - current_block->terminator = SPIRBlock::MultiSelect; +const std::string &Compiler::get_member_name(uint32_t id, uint32_t index) const +{ + return ir.get_member_name(id, index); +} - current_block->condition = ops[0]; - current_block->default_block = ops[1]; +void Compiler::set_member_qualified_name(uint32_t type_id, uint32_t index, const std::string &name) +{ + ir.meta.at(type_id).members.resize(max(ir.meta[type_id].members.size(), size_t(index) + 1)); + ir.meta.at(type_id).members[index].qualified_alias = name; +} - for (uint32_t i = 2; i + 2 <= length; i += 2) - current_block->cases.push_back({ ops[i], ops[i + 1] }); +const std::string &Compiler::get_member_qualified_name(uint32_t type_id, uint32_t index) const +{ + const static string empty; - // If we jump to next block, make it break instead since we're inside a switch case block at that point. - multiselect_merge_targets.insert(current_block->next_block); + auto &m = ir.meta.at(type_id); + if (index < m.members.size()) + return m.members[index].qualified_alias; + else + return empty; +} - current_block = nullptr; - break; - } +uint32_t Compiler::get_member_decoration(uint32_t id, uint32_t index, Decoration decoration) const +{ + return ir.get_member_decoration(id, index, decoration); +} - case OpKill: - { - if (!current_block) - SPIRV_CROSS_THROW("Trying to end a non-existing block."); - current_block->terminator = SPIRBlock::Kill; - current_block = nullptr; - break; - } +uint64_t Compiler::get_member_decoration_mask(uint32_t id, uint32_t index) const +{ + return get_member_decoration_bitset(id, index).get_lower(); +} - case OpReturn: - { - if (!current_block) - SPIRV_CROSS_THROW("Trying to end a non-existing block."); - current_block->terminator = SPIRBlock::Return; - current_block = nullptr; - break; - } +const Bitset &Compiler::get_member_decoration_bitset(uint32_t id, uint32_t index) const +{ + return ir.get_member_decoration_bitset(id, index); +} - case OpReturnValue: - { - if (!current_block) - SPIRV_CROSS_THROW("Trying to end a non-existing block."); - current_block->terminator = SPIRBlock::Return; - current_block->return_value = ops[0]; - current_block = nullptr; - break; - } +bool Compiler::has_member_decoration(uint32_t id, uint32_t index, Decoration decoration) const +{ + return ir.has_member_decoration(id, index, decoration); +} - case OpUnreachable: - { - if (!current_block) - SPIRV_CROSS_THROW("Trying to end a non-existing block."); - current_block->terminator = SPIRBlock::Unreachable; - current_block = nullptr; - break; - } +void Compiler::unset_member_decoration(uint32_t id, uint32_t index, Decoration decoration) +{ + ir.unset_member_decoration(id, index, decoration); +} - case OpSelectionMerge: - { - if (!current_block) - SPIRV_CROSS_THROW("Trying to modify a non-existing block."); +void Compiler::set_decoration_string(uint32_t id, spv::Decoration decoration, const std::string &argument) +{ + ir.set_decoration_string(id, decoration, argument); +} - current_block->next_block = ops[0]; - current_block->merge = SPIRBlock::MergeSelection; - selection_merge_targets.insert(current_block->next_block); +void Compiler::set_decoration(uint32_t id, Decoration decoration, uint32_t argument) +{ + ir.set_decoration(id, decoration, argument); +} - if (length >= 2) - { - if (ops[1] & SelectionControlFlattenMask) - current_block->hint = SPIRBlock::HintFlatten; - else if (ops[1] & SelectionControlDontFlattenMask) - current_block->hint = SPIRBlock::HintDontFlatten; - } - break; - } +StorageClass Compiler::get_storage_class(uint32_t id) const +{ + return get(id).storage; +} - case OpLoopMerge: - { - if (!current_block) - SPIRV_CROSS_THROW("Trying to modify a non-existing block."); +const std::string &Compiler::get_name(uint32_t id) const +{ + return ir.meta.at(id).decoration.alias; +} - current_block->merge_block = ops[0]; - current_block->continue_block = ops[1]; - current_block->merge = SPIRBlock::MergeLoop; +const std::string Compiler::get_fallback_name(uint32_t id) const +{ + return join("_", id); +} - loop_blocks.insert(current_block->self); - loop_merge_targets.insert(current_block->merge_block); +const std::string Compiler::get_block_fallback_name(uint32_t id) const +{ + auto &var = get(id); + if (get_name(id).empty()) + return join("_", get(var.basetype).self, "_", id); + else + return get_name(id); +} - continue_block_to_loop_header[current_block->continue_block] = current_block->self; +uint64_t Compiler::get_decoration_mask(uint32_t id) const +{ + return get_decoration_bitset(id).get_lower(); +} - // Don't add loop headers to continue blocks, - // which would make it impossible branch into the loop header since - // they are treated as continues. - if (current_block->continue_block != current_block->self) - continue_blocks.insert(current_block->continue_block); +const Bitset &Compiler::get_decoration_bitset(uint32_t id) const +{ + return ir.get_decoration_bitset(id); +} - if (length >= 3) - { - if (ops[2] & LoopControlUnrollMask) - current_block->hint = SPIRBlock::HintUnroll; - else if (ops[2] & LoopControlDontUnrollMask) - current_block->hint = SPIRBlock::HintDontUnroll; - } - break; - } +bool Compiler::has_decoration(uint32_t id, Decoration decoration) const +{ + return ir.has_decoration(id, decoration); +} - case OpSpecConstantOp: - { - if (length < 3) - SPIRV_CROSS_THROW("OpSpecConstantOp not enough arguments."); +const string &Compiler::get_decoration_string(uint32_t id, Decoration decoration) const +{ + return ir.get_decoration_string(id, decoration); +} - uint32_t result_type = ops[0]; - uint32_t id = ops[1]; - auto spec_op = static_cast(ops[2]); +const string &Compiler::get_member_decoration_string(uint32_t id, uint32_t index, Decoration decoration) const +{ + return ir.get_member_decoration_string(id, index, decoration); +} - set(id, result_type, spec_op, ops + 3, length - 3); - break; - } +uint32_t Compiler::get_decoration(uint32_t id, Decoration decoration) const +{ + return ir.get_decoration(id, decoration); +} - // Actual opcodes. - default: - { - if (!current_block) - SPIRV_CROSS_THROW("Currently no block to insert opcode."); +void Compiler::unset_decoration(uint32_t id, Decoration decoration) +{ + ir.unset_decoration(id, decoration); +} - current_block->ops.push_back(instruction); - break; - } - } +bool Compiler::get_binary_offset_for_decoration(uint32_t id, spv::Decoration decoration, uint32_t &word_offset) const +{ + auto &word_offsets = ir.meta.at(id).decoration_word_offset; + auto itr = word_offsets.find(decoration); + if (itr == end(word_offsets)) + return false; + + word_offset = itr->second; + return true; } bool Compiler::block_is_loop_candidate(const SPIRBlock &block, SPIRBlock::Method method) const @@ -2567,7 +1353,7 @@ bool Compiler::traverse_all_reachable_opcodes(const SPIRFunction &func, OpcodeHa uint32_t Compiler::type_struct_member_offset(const SPIRType &type, uint32_t index) const { // Decoration must be set in valid SPIR-V, otherwise throw. - auto &dec = meta[type.self].members.at(index); + auto &dec = ir.meta[type.self].members.at(index); if (dec.decoration_flags.get(DecorationOffset)) return dec.offset; else @@ -2578,7 +1364,7 @@ uint32_t Compiler::type_struct_member_array_stride(const SPIRType &type, uint32_ { // Decoration must be set in valid SPIR-V, otherwise throw. // ArrayStride is part of the array type not OpMemberDecorate. - auto &dec = meta[type.member_types[index]].decoration; + auto &dec = ir.meta[type.member_types[index]].decoration; if (dec.decoration_flags.get(DecorationArrayStride)) return dec.array_stride; else @@ -2589,7 +1375,7 @@ uint32_t Compiler::type_struct_member_matrix_stride(const SPIRType &type, uint32 { // Decoration must be set in valid SPIR-V, otherwise throw. // MatrixStride is part of OpMemberDecorate. - auto &dec = meta[type.self].members[index]; + auto &dec = ir.meta[type.self].members[index]; if (dec.decoration_flags.get(DecorationMatrixStride)) return dec.matrix_stride; else @@ -2728,21 +1514,10 @@ std::vector Compiler::get_active_buffer_ranges(uint32_t id) const { std::vector ranges; BufferAccessHandler handler(*this, ranges, id); - traverse_all_reachable_opcodes(get(entry_point), handler); + traverse_all_reachable_opcodes(get(ir.default_entry_point), handler); return ranges; } -// Increase the number of IDs by the specified incremental amount. -// Returns the value of the first ID available for use in the expanded bound. -uint32_t Compiler::increase_bound_by(uint32_t incr_amount) -{ - auto curr_bound = ids.size(); - auto new_bound = curr_bound + incr_amount; - ids.resize(new_bound); - meta.resize(new_bound); - return uint32_t(curr_bound); -} - bool Compiler::types_are_logically_equivalent(const SPIRType &a, const SPIRType &b) const { if (a.basetype != b.basetype) @@ -2948,7 +1723,7 @@ void Compiler::inherit_expression_dependencies(uint32_t dst, uint32_t source_exp vector Compiler::get_entry_points() const { vector entries; - for (auto &entry : entry_points) + for (auto &entry : ir.entry_points) entries.push_back(entry.second.orig_name); return entries; } @@ -2956,7 +1731,7 @@ vector Compiler::get_entry_points() const vector Compiler::get_entry_points_and_stages() const { vector entries; - for (auto &entry : entry_points) + for (auto &entry : ir.entry_points) entries.push_back({ entry.second.orig_name, entry.second.model }); return entries; } @@ -2978,13 +1753,13 @@ void Compiler::rename_entry_point(const std::string &old_name, const std::string void Compiler::set_entry_point(const std::string &name) { auto &entry = get_first_entry_point(name); - entry_point = entry.self; + ir.default_entry_point = entry.self; } void Compiler::set_entry_point(const std::string &name, spv::ExecutionModel model) { auto &entry = get_entry_point(name, model); - entry_point = entry.self; + ir.default_entry_point = entry.self; } SPIREntryPoint &Compiler::get_entry_point(const std::string &name) @@ -2999,12 +1774,11 @@ const SPIREntryPoint &Compiler::get_entry_point(const std::string &name) const SPIREntryPoint &Compiler::get_first_entry_point(const std::string &name) { - auto itr = - find_if(begin(entry_points), end(entry_points), [&](const std::pair &entry) -> bool { - return entry.second.orig_name == name; - }); + auto itr = find_if( + begin(ir.entry_points), end(ir.entry_points), + [&](const std::pair &entry) -> bool { return entry.second.orig_name == name; }); - if (itr == end(entry_points)) + if (itr == end(ir.entry_points)) SPIRV_CROSS_THROW("Entry point does not exist."); return itr->second; @@ -3012,12 +1786,11 @@ SPIREntryPoint &Compiler::get_first_entry_point(const std::string &name) const SPIREntryPoint &Compiler::get_first_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.orig_name == name; - }); + auto itr = find_if( + begin(ir.entry_points), end(ir.entry_points), + [&](const std::pair &entry) -> bool { return entry.second.orig_name == name; }); - if (itr == end(entry_points)) + if (itr == end(ir.entry_points)) SPIRV_CROSS_THROW("Entry point does not exist."); return itr->second; @@ -3025,12 +1798,12 @@ const SPIREntryPoint &Compiler::get_first_entry_point(const std::string &name) c SPIREntryPoint &Compiler::get_entry_point(const std::string &name, ExecutionModel model) { - auto itr = - find_if(begin(entry_points), end(entry_points), [&](const std::pair &entry) -> bool { - return entry.second.orig_name == name && entry.second.model == model; - }); + auto itr = find_if(begin(ir.entry_points), end(ir.entry_points), + [&](const std::pair &entry) -> bool { + return entry.second.orig_name == name && entry.second.model == model; + }); - if (itr == end(entry_points)) + if (itr == end(ir.entry_points)) SPIRV_CROSS_THROW("Entry point does not exist."); return itr->second; @@ -3038,12 +1811,12 @@ SPIREntryPoint &Compiler::get_entry_point(const std::string &name, ExecutionMode const SPIREntryPoint &Compiler::get_entry_point(const std::string &name, ExecutionModel model) const { - auto itr = - find_if(begin(entry_points), end(entry_points), [&](const std::pair &entry) -> bool { - return entry.second.orig_name == name && entry.second.model == model; - }); + auto itr = find_if(begin(ir.entry_points), end(ir.entry_points), + [&](const std::pair &entry) -> bool { + return entry.second.orig_name == name && entry.second.model == model; + }); - if (itr == end(entry_points)) + if (itr == end(ir.entry_points)) SPIRV_CROSS_THROW("Entry point does not exist."); return itr->second; @@ -3061,12 +1834,12 @@ const string &Compiler::get_cleansed_entry_point_name(const std::string &name, E const SPIREntryPoint &Compiler::get_entry_point() const { - return entry_points.find(entry_point)->second; + return ir.entry_points.find(ir.default_entry_point)->second; } SPIREntryPoint &Compiler::get_entry_point() { - return entry_points.find(entry_point)->second; + return ir.entry_points.find(ir.default_entry_point)->second; } bool Compiler::interface_variable_exists_in_entry_point(uint32_t id) const @@ -3080,7 +1853,7 @@ bool Compiler::interface_variable_exists_in_entry_point(uint32_t id) const // 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) + if (ir.entry_points.size() <= 1) return true; auto &execution = get_entry_point(); @@ -3219,7 +1992,7 @@ void Compiler::CombinedImageSamplerHandler::register_combined_image_sampler(SPIR if (itr == end(caller.combined_parameters)) { - uint32_t id = compiler.increase_bound_by(3); + uint32_t id = compiler.ir.increase_bound_by(3); auto type_id = id + 0; auto ptr_type_id = id + 1; auto combined_id = id + 2; @@ -3242,8 +2015,8 @@ void Compiler::CombinedImageSamplerHandler::register_combined_image_sampler(SPIR compiler.set(combined_id, ptr_type_id, StorageClassFunction, 0); // Inherit RelaxedPrecision (and potentially other useful flags if deemed relevant). - auto &new_flags = compiler.meta[combined_id].decoration.decoration_flags; - auto &old_flags = compiler.meta[sampler_id].decoration.decoration_flags; + auto &new_flags = compiler.ir.meta[combined_id].decoration.decoration_flags; + auto &old_flags = compiler.ir.meta[sampler_id].decoration.decoration_flags; new_flags.reset(); if (old_flags.get(DecorationRelaxedPrecision)) new_flags.set(DecorationRelaxedPrecision); @@ -3327,7 +2100,7 @@ bool Compiler::DummySamplerForCombinedImageHandler::handle(Op opcode, const uint compiler.register_read(id, ptr, true); // Other backends might use SPIRAccessChain for this later. - compiler.ids[id].set_allow_type_rewrite(); + compiler.ir.ids[id].set_allow_type_rewrite(); break; } @@ -3474,7 +2247,7 @@ bool Compiler::CombinedImageSamplerHandler::handle(Op opcode, const uint32_t *ar if (is_fetch) { // Have to invent the sampled image type. - sampled_type = compiler.increase_bound_by(1); + sampled_type = compiler.ir.increase_bound_by(1); auto &type = compiler.set(sampled_type); type = compiler.expression_type(args[2]); type.self = sampled_type; @@ -3486,7 +2259,7 @@ bool Compiler::CombinedImageSamplerHandler::handle(Op opcode, const uint32_t *ar sampled_type = args[0]; } - auto id = compiler.increase_bound_by(2); + auto id = compiler.ir.increase_bound_by(2); auto type_id = id + 0; auto combined_id = id + 1; @@ -3502,9 +2275,9 @@ bool Compiler::CombinedImageSamplerHandler::handle(Op opcode, const uint32_t *ar compiler.set(combined_id, type_id, StorageClassUniformConstant, 0); // Inherit RelaxedPrecision (and potentially other useful flags if deemed relevant). - auto &new_flags = compiler.meta[combined_id].decoration.decoration_flags; + auto &new_flags = compiler.ir.meta[combined_id].decoration.decoration_flags; // Fetch inherits precision from the image, not sampler (there is no sampler). - auto &old_flags = compiler.meta[is_fetch ? image_id : sampler_id].decoration.decoration_flags; + auto &old_flags = compiler.ir.meta[is_fetch ? image_id : sampler_id].decoration.decoration_flags; new_flags.reset(); if (old_flags.get(DecorationRelaxedPrecision)) new_flags.set(DecorationRelaxedPrecision); @@ -3527,10 +2300,10 @@ bool Compiler::CombinedImageSamplerHandler::handle(Op opcode, const uint32_t *ar uint32_t Compiler::build_dummy_sampler_for_combined_images() { DummySamplerForCombinedImageHandler handler(*this); - traverse_all_reachable_opcodes(get(entry_point), handler); + traverse_all_reachable_opcodes(get(ir.default_entry_point), handler); if (handler.need_dummy_sampler) { - uint32_t offset = increase_bound_by(3); + uint32_t offset = ir.increase_bound_by(3); auto type_id = offset + 0; auto ptr_type_id = offset + 1; auto var_id = offset + 2; @@ -3556,7 +2329,7 @@ uint32_t Compiler::build_dummy_sampler_for_combined_images() void Compiler::build_combined_image_samplers() { - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeFunction) { @@ -3569,13 +2342,13 @@ void Compiler::build_combined_image_samplers() combined_image_samplers.clear(); CombinedImageSamplerHandler handler(*this); - traverse_all_reachable_opcodes(get(entry_point), handler); + traverse_all_reachable_opcodes(get(ir.default_entry_point), handler); } vector Compiler::get_specialization_constants() const { vector spec_consts; - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeConstant) { @@ -3597,31 +2370,6 @@ const SPIRConstant &Compiler::get_constant(uint32_t id) const return get(id); } -// Recursively marks any constants referenced by the specified constant instruction as being used -// as an array length. The id must be a constant instruction (SPIRConstant or SPIRConstantOp). -void Compiler::mark_used_as_array_length(uint32_t id) -{ - switch (ids[id].get_type()) - { - case TypeConstant: - get(id).is_used_as_array_length = true; - break; - - case TypeConstantOp: - { - auto &cop = get(id); - for (uint32_t arg_id : cop.arguments) - mark_used_as_array_length(arg_id); - } - - case TypeUndef: - return; - - default: - SPIRV_CROSS_THROW("Array lengths must be a constant instruction (OpConstant.. or OpSpecConstant...)."); - } -} - static bool exists_unaccessed_path_to_return(const CFG &cfg, uint32_t block, const unordered_set &blocks) { // This block accesses the variable. @@ -3781,7 +2529,7 @@ bool Compiler::AnalyzeVariableScopeAccessHandler::id_is_potential_temporary(uint return false; // Temporaries are not created before we start emitting code. - return compiler.ids[id].empty() || (compiler.ids[id].get_type() == TypeExpression); + return compiler.ir.ids[id].empty() || (compiler.ir.ids[id].get_type() == TypeExpression); } bool Compiler::AnalyzeVariableScopeAccessHandler::handle(spv::Op op, const uint32_t *args, uint32_t length) @@ -3836,7 +2584,7 @@ bool Compiler::AnalyzeVariableScopeAccessHandler::handle(spv::Op op, const uint3 e.loaded_from = backing_variable ? backing_variable->self : 0; // Other backends might use SPIRAccessChain for this later. - compiler.ids[args[1]].set_allow_type_rewrite(); + compiler.ir.ids[args[1]].set_allow_type_rewrite(); break; } @@ -4048,7 +2796,7 @@ void Compiler::find_function_local_luts(SPIRFunction &entry, const AnalyzeVariab uint32_t static_constant_expression = 0; if (var.initializer) { - if (ids[var.initializer].get_type() != TypeConstant) + if (ir.ids[var.initializer].get_type() != TypeConstant) continue; static_constant_expression = var.initializer; @@ -4095,7 +2843,7 @@ void Compiler::find_function_local_luts(SPIRFunction &entry, const AnalyzeVariab continue; // Is it a constant expression? - if (ids[static_expression_handler.static_expression].get_type() != TypeConstant) + if (ir.ids[static_expression_handler.static_expression].get_type() != TypeConstant) continue; // We found a LUT! @@ -4147,7 +2895,7 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry, AnalyzeVariableScopeA // The continue block is dominated by the inner part of the loop, which does not make sense in high-level // language output because it will be declared before the body, // so we will have to lift the dominator up to the relevant loop header instead. - builder.add_block(continue_block_to_loop_header[block]); + builder.add_block(ir.continue_block_to_loop_header[block]); // Arrays or structs cannot be loop variables. if (type.vecsize == 1 && type.columns == 1 && type.basetype != SPIRType::Struct && type.array.empty()) @@ -4203,7 +2951,7 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry, AnalyzeVariableScopeA // If a temporary is used in more than one block, we might have to lift continue block // access up to loop header like we did for variables. if (blocks.size() != 1 && is_continue(block)) - builder.add_block(continue_block_to_loop_header[block]); + builder.add_block(ir.continue_block_to_loop_header[block]); else if (blocks.size() != 1 && is_single_block_loop(block)) { // Awkward case, because the loop header is also the continue block. @@ -4261,14 +3009,17 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry, AnalyzeVariableScopeA uint32_t header = 0; - // Find the loop header for this block. - for (auto b : loop_blocks) + // Find the loop header for this block if we are a continue block. { - auto &potential_header = get(b); - if (potential_header.continue_block == block) + auto itr = ir.continue_block_to_loop_header.find(block); + if (itr != end(ir.continue_block_to_loop_header)) { - header = b; - break; + header = itr->second; + } + else if (get(block).continue_block == block) + { + // Also check for self-referential continue block. + header = block; } } @@ -4335,28 +3086,7 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry, AnalyzeVariableScopeA Bitset Compiler::get_buffer_block_flags(uint32_t id) const { - return get_buffer_block_flags(get(id)); -} - -Bitset Compiler::get_buffer_block_flags(const SPIRVariable &var) const -{ - auto &type = get(var.basetype); - assert(type.basetype == SPIRType::Struct); - - // Some flags like non-writable, non-readable are actually found - // as member decorations. If all members have a decoration set, propagate - // the decoration up as a regular variable decoration. - Bitset base_flags = meta[var.self].decoration.decoration_flags; - - if (type.member_types.empty()) - return base_flags; - - Bitset all_members_flags = get_member_decoration_bitset(type.self, 0); - for (uint32_t i = 1; i < uint32_t(type.member_types.size()); i++) - all_members_flags.merge_and(get_member_decoration_bitset(type.self, i)); - - base_flags.merge_or(all_members_flags); - return base_flags; + return ir.get_buffer_block_flags(get(id)); } bool Compiler::get_common_basic_type(const SPIRType &type, SPIRType::BaseType &base_type) @@ -4420,7 +3150,7 @@ bool Compiler::ActiveBuiltinHandler::handle(spv::Op opcode, const uint32_t *args // Only handles variables here. // Builtins which are part of a block are handled in AccessChain. auto *var = compiler.maybe_get(id); - auto &decorations = compiler.meta[id].decoration; + auto &decorations = compiler.ir.meta[id].decoration; if (var && decorations.builtin) { auto &type = compiler.get(var->basetype); @@ -4503,9 +3233,9 @@ bool Compiler::ActiveBuiltinHandler::handle(spv::Op opcode, const uint32_t *args { uint32_t index = compiler.get(args[i]).scalar(); - if (index < uint32_t(compiler.meta[type->self].members.size())) + if (index < uint32_t(compiler.ir.meta[type->self].members.size())) { - auto &decorations = compiler.meta[type->self].members[index]; + auto &decorations = compiler.ir.meta[type->self].members[index]; if (decorations.builtin) { flags.set(decorations.builtin_type); @@ -4539,7 +3269,7 @@ void Compiler::update_active_builtins() cull_distance_count = 0; clip_distance_count = 0; ActiveBuiltinHandler handler(*this); - traverse_all_reachable_opcodes(get(entry_point), handler); + traverse_all_reachable_opcodes(get(ir.default_entry_point), handler); } // Returns whether this shader uses a builtin of the storage class @@ -4564,10 +3294,10 @@ bool Compiler::has_active_builtin(BuiltIn builtin, StorageClass storage) void Compiler::analyze_image_and_sampler_usage() { CombinedImageSamplerDrefHandler dref_handler(*this); - traverse_all_reachable_opcodes(get(entry_point), dref_handler); + traverse_all_reachable_opcodes(get(ir.default_entry_point), dref_handler); CombinedImageSamplerUsageHandler handler(*this, dref_handler.dref_combined_samplers); - traverse_all_reachable_opcodes(get(entry_point), handler); + traverse_all_reachable_opcodes(get(ir.default_entry_point), handler); comparison_ids = move(handler.comparison_ids); need_subpass_input = handler.need_subpass_input; @@ -4605,8 +3335,8 @@ bool Compiler::CombinedImageSamplerDrefHandler::handle(spv::Op opcode, const uin void Compiler::build_function_control_flow_graphs_and_analyze() { CFGBuilder handler(*this); - handler.function_cfgs[entry_point].reset(new CFG(*this, get(entry_point))); - traverse_all_reachable_opcodes(get(entry_point), handler); + handler.function_cfgs[ir.default_entry_point].reset(new CFG(*this, get(ir.default_entry_point))); + traverse_all_reachable_opcodes(get(ir.default_entry_point), handler); function_cfgs = move(handler.function_cfgs); for (auto &f : function_cfgs) @@ -4752,13 +3482,13 @@ bool Compiler::CombinedImageSamplerUsageHandler::handle(Op opcode, const uint32_ bool Compiler::buffer_is_hlsl_counter_buffer(uint32_t id) const { // First, check for the proper decoration. - if (meta.at(id).hlsl_is_magic_counter_buffer) + if (ir.meta.at(id).hlsl_is_magic_counter_buffer) return true; // Check for legacy fallback method. // FIXME: This should be deprecated eventually. - if (meta.at(id).hlsl_magic_counter_buffer_candidate) + if (ir.meta.at(id).hlsl_magic_counter_buffer_candidate) { auto *var = maybe_get(id); // Ensure that this is actually a buffer object. @@ -4772,9 +3502,9 @@ bool Compiler::buffer_is_hlsl_counter_buffer(uint32_t id) const bool Compiler::buffer_get_hlsl_counter_buffer(uint32_t id, uint32_t &counter_id) const { // First, check for the proper decoration. - if (meta[id].hlsl_magic_counter_buffer != 0) + if (ir.meta[id].hlsl_magic_counter_buffer != 0) { - counter_id = meta[id].hlsl_magic_counter_buffer; + counter_id = ir.meta[id].hlsl_magic_counter_buffer; return true; } @@ -4785,7 +3515,7 @@ bool Compiler::buffer_get_hlsl_counter_buffer(uint32_t id, uint32_t &counter_id) uint32_t id_bound = get_current_id_bound(); for (uint32_t i = 0; i < id_bound; i++) { - if (meta[i].hlsl_magic_counter_buffer_candidate && meta[i].hlsl_magic_counter_buffer_name == name) + if (ir.meta[i].hlsl_magic_counter_buffer_candidate && ir.meta[i].hlsl_magic_counter_buffer_name == name) { auto *var = maybe_get(i); // Ensure that this is actually a buffer object. @@ -4807,7 +3537,7 @@ void Compiler::make_constant_null(uint32_t id, uint32_t type) if (!constant_type.array.empty()) { assert(constant_type.parent_type); - uint32_t parent_id = increase_bound_by(1); + uint32_t parent_id = ir.increase_bound_by(1); make_constant_null(parent_id, constant_type.parent_type); if (!constant_type.array_size_literal.back()) @@ -4820,7 +3550,7 @@ void Compiler::make_constant_null(uint32_t id, uint32_t type) } else if (!constant_type.member_types.empty()) { - uint32_t member_ids = increase_bound_by(uint32_t(constant_type.member_types.size())); + uint32_t member_ids = ir.increase_bound_by(uint32_t(constant_type.member_types.size())); vector elements(constant_type.member_types.size()); for (uint32_t i = 0; i < constant_type.member_types.size(); i++) { @@ -4838,12 +3568,12 @@ void Compiler::make_constant_null(uint32_t id, uint32_t type) const std::vector &Compiler::get_declared_capabilities() const { - return declared_capabilities; + return ir.declared_capabilities; } const std::vector &Compiler::get_declared_extensions() const { - return declared_extensions; + return ir.declared_extensions; } std::string Compiler::get_remapped_declared_block_name(uint32_t id) const @@ -4855,7 +3585,7 @@ std::string Compiler::get_remapped_declared_block_name(uint32_t id) const { auto &var = get(id); auto &type = get(var.basetype); - auto &block_name = meta[type.self].decoration.alias; + auto &block_name = ir.meta[type.self].decoration.alias; return block_name.empty() ? get_block_fallback_name(id) : block_name; } } @@ -4903,7 +3633,7 @@ bool Compiler::instruction_to_result_type(uint32_t &result_type, uint32_t &resul Bitset Compiler::combined_decoration_for_member(const SPIRType &type, uint32_t index) const { Bitset flags; - auto &memb = meta[type.self].members; + auto &memb = ir.meta[type.self].members; if (index >= memb.size()) return flags; auto &dec = memb[index]; diff --git a/spirv_cross.hpp b/spirv_cross.hpp index 1ca207be5..82ff0418e 100644 --- a/spirv_cross.hpp +++ b/spirv_cross.hpp @@ -19,7 +19,7 @@ #include "spirv.hpp" #include "spirv_cfg.hpp" -#include "spirv_common.hpp" +#include "spirv_cross_parsed_ir.hpp" namespace spirv_cross { @@ -121,9 +121,16 @@ class Compiler friend class DominatorBuilder; // The constructor takes a buffer of SPIR-V words and parses it. - Compiler(std::vector ir); + // It will create its own parser, parse the SPIR-V and move the parsed IR + // as if you had called the constructors taking ParsedIR directly. + explicit Compiler(std::vector ir); Compiler(const uint32_t *ir, size_t word_count); + // This is more modular. We can also consume a ParsedIR structure directly, either as a move, or copy. + // With copy, we can reuse the same parsed IR for multiple Compiler instances. + explicit Compiler(const ParsedIR &ir); + explicit Compiler(ParsedIR &&ir); + virtual ~Compiler() = default; // After parsing, API users can modify the SPIR-V via reflection and call this @@ -160,7 +167,7 @@ class Compiler uint32_t get_decoration(uint32_t id, spv::Decoration decoration) const; const std::string &get_decoration_string(uint32_t id, spv::Decoration decoration) const; - // Removes the decoration for a an ID. + // Removes the decoration for an ID. void unset_decoration(uint32_t id, spv::Decoration decoration); // Gets the SPIR-V type associated with ID. @@ -444,7 +451,7 @@ class Compiler uint32_t get_current_id_bound() const { - return uint32_t(ids.size()); + return uint32_t(ir.ids.size()); } // API for querying buffer objects. @@ -522,20 +529,19 @@ class Compiler if (!instr.length) return nullptr; - if (instr.offset + instr.length > spirv.size()) + if (instr.offset + instr.length > ir.spirv.size()) SPIRV_CROSS_THROW("Compiler::stream() out of range."); - return &spirv[instr.offset]; + return &ir.spirv[instr.offset]; } - std::vector spirv; - std::vector inst; - std::vector ids; - std::vector meta; + ParsedIR ir; + // Marks variables which have global scope and variables which can alias with other variables + // (SSBO, image load store, etc) + std::vector global_variables; + std::vector aliased_variables; SPIRFunction *current_function = nullptr; SPIRBlock *current_block = nullptr; - std::vector global_variables; - std::vector aliased_variables; std::unordered_set active_interface_variables; bool check_active_interface_variables = false; @@ -544,7 +550,7 @@ class Compiler template T &set(uint32_t id, P &&... args) { - auto &var = variant_set(ids.at(id), std::forward

(args)...); + auto &var = variant_set(ir.ids.at(id), std::forward

(args)...); var.self = id; return var; } @@ -552,13 +558,13 @@ class Compiler template T &get(uint32_t id) { - return variant_get(ids.at(id)); + return variant_get(ir.ids.at(id)); } template T *maybe_get(uint32_t id) { - if (ids.at(id).get_type() == T::type) + if (ir.ids.at(id).get_type() == T::type) return &get(id); else return nullptr; @@ -567,42 +573,21 @@ class Compiler template const T &get(uint32_t id) const { - return variant_get(ids.at(id)); + return variant_get(ir.ids.at(id)); } template const T *maybe_get(uint32_t id) const { - if (ids.at(id).get_type() == T::type) + if (ir.ids.at(id).get_type() == T::type) return &get(id); else return nullptr; } - 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 - { - uint32_t version = 0; - bool es = false; - bool known = false; - bool hlsl = false; - - Source() = default; - } source; - - std::unordered_set loop_blocks; - std::unordered_set continue_blocks; - std::unordered_set loop_merge_targets; - std::unordered_set selection_merge_targets; - std::unordered_set multiselect_merge_targets; - std::unordered_map continue_block_to_loop_header; - virtual std::string to_name(uint32_t id, bool allow_alias = true) const; bool is_builtin_variable(const SPIRVariable &var) const; bool is_builtin_type(const SPIRType &type) const; @@ -618,14 +603,13 @@ class Compiler bool expression_is_lvalue(uint32_t id) const; bool variable_storage_is_aliased(const SPIRVariable &var); SPIRVariable *maybe_get_backing_variable(uint32_t chain); - void mark_used_as_array_length(uint32_t id); void register_read(uint32_t expr, uint32_t chain, bool forwarded); void register_write(uint32_t chain); inline bool is_continue(uint32_t next) const { - return continue_blocks.find(next) != end(continue_blocks); + return (ir.block_meta[next] & ParsedIR::BLOCK_META_CONTINUE_BIT) != 0; } inline bool is_single_block_loop(uint32_t next) const @@ -636,19 +620,19 @@ class Compiler inline bool is_break(uint32_t next) const { - return loop_merge_targets.find(next) != end(loop_merge_targets) || - multiselect_merge_targets.find(next) != end(multiselect_merge_targets); + return (ir.block_meta[next] & + (ParsedIR::BLOCK_META_LOOP_MERGE_BIT | ParsedIR::BLOCK_META_MULTISELECT_MERGE_BIT)) != 0; } inline bool is_loop_break(uint32_t next) const { - return loop_merge_targets.find(next) != end(loop_merge_targets); + return (ir.block_meta[next] & ParsedIR::BLOCK_META_LOOP_MERGE_BIT) != 0; } inline bool is_conditional(uint32_t next) const { - return selection_merge_targets.find(next) != end(selection_merge_targets) && - multiselect_merge_targets.find(next) == end(multiselect_merge_targets); + return (ir.block_meta[next] & + (ParsedIR::BLOCK_META_SELECTION_MERGE_BIT | ParsedIR::BLOCK_META_MULTISELECT_MERGE_BIT)) != 0; } // Dependency tracking for temporaries read from variables. @@ -675,8 +659,6 @@ class Compiler bool block_is_loop_candidate(const SPIRBlock &block, SPIRBlock::Method method) const; - uint32_t increase_bound_by(uint32_t incr_amount); - bool types_are_logically_equivalent(const SPIRType &a, const SPIRType &b) const; void inherit_expression_dependencies(uint32_t dst, uint32_t source); @@ -692,8 +674,9 @@ class Compiler variable_remap_callback(type, var_name, type_name); } - void parse(); - void parse(const Instruction &i); + void set_ir(const ParsedIR &parsed); + void set_ir(ParsedIR &&parsed); + void parse_fixup(); // Used internally to implement various traversals for queries. struct OpcodeHandler @@ -813,7 +796,6 @@ class Compiler VariableTypeRemapCallback variable_remap_callback; - Bitset get_buffer_block_flags(const SPIRVariable &var) const; bool get_common_basic_type(const SPIRType &type, SPIRType::BaseType &base_type); std::unordered_set forced_temporaries; @@ -934,8 +916,6 @@ class Compiler void make_constant_null(uint32_t id, uint32_t type); - std::vector declared_capabilities; - std::vector declared_extensions; std::unordered_map declared_block_names; bool instruction_to_result_type(uint32_t &result_type, uint32_t &result_id, spv::Op op, const uint32_t *args, diff --git a/spirv_cross_parsed_ir.cpp b/spirv_cross_parsed_ir.cpp new file mode 100644 index 000000000..df24a1e32 --- /dev/null +++ b/spirv_cross_parsed_ir.cpp @@ -0,0 +1,557 @@ +/* + * Copyright 2018 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_cross_parsed_ir.hpp" +#include + +using namespace std; +using namespace spv; + +namespace spirv_cross +{ +void ParsedIR::set_id_bounds(uint32_t bounds) +{ + ids.resize(bounds); + meta.resize(bounds); + block_meta.resize(bounds); +} + +static string ensure_valid_identifier(const string &name, bool member) +{ + // Functions in glslangValidator are mangled with name( stuff. + // Normally, we would never see '(' in any legal identifiers, so just strip them out. + auto str = name.substr(0, name.find('(')); + + for (uint32_t i = 0; i < str.size(); i++) + { + auto &c = str[i]; + + if (member) + { + // _m variables are reserved by the internal implementation, + // otherwise, make sure the name is a valid identifier. + if (i == 0) + c = isalpha(c) ? c : '_'; + else if (i == 2 && str[0] == '_' && str[1] == 'm') + c = isalpha(c) ? c : '_'; + else + c = isalnum(c) ? c : '_'; + } + else + { + // _ variables are reserved by the internal implementation, + // otherwise, make sure the name is a valid identifier. + if (i == 0 || (str[0] == '_' && i == 1)) + c = isalpha(c) ? c : '_'; + else + c = isalnum(c) ? c : '_'; + } + } + return str; +} + +const string &ParsedIR::get_name(uint32_t id) const +{ + return meta[id].decoration.alias; +} + +const string &ParsedIR::get_member_name(uint32_t id, uint32_t index) const +{ + auto &m = meta[id]; + if (index >= m.members.size()) + { + static string empty; + return empty; + } + + return m.members[index].alias; +} + +void ParsedIR::set_name(uint32_t id, const string &name) +{ + auto &str = meta[id].decoration.alias; + str.clear(); + + if (name.empty()) + return; + + // glslang uses identifiers to pass along meaningful information + // about HLSL reflection. + // FIXME: This should be deprecated eventually. + auto &m = meta[id]; + if (source.hlsl && name.size() >= 6 && name.find("@count") == name.size() - 6) + { + m.hlsl_magic_counter_buffer_candidate = true; + m.hlsl_magic_counter_buffer_name = name.substr(0, name.find("@count")); + } + else + { + m.hlsl_magic_counter_buffer_candidate = false; + m.hlsl_magic_counter_buffer_name.clear(); + } + + // Reserved for temporaries. + if (name[0] == '_' && name.size() >= 2 && isdigit(name[1])) + return; + + str = ensure_valid_identifier(name, false); +} + +void ParsedIR::set_member_name(uint32_t id, uint32_t index, const string &name) +{ + meta[id].members.resize(max(meta[id].members.size(), size_t(index) + 1)); + + auto &str = meta[id].members[index].alias; + str.clear(); + if (name.empty()) + return; + + // Reserved for unnamed members. + if (name[0] == '_' && name.size() >= 3 && name[1] == 'm' && isdigit(name[2])) + return; + + str = ensure_valid_identifier(name, true); +} + +void ParsedIR::set_decoration_string(uint32_t id, Decoration decoration, const string &argument) +{ + auto &dec = meta[id].decoration; + dec.decoration_flags.set(decoration); + + switch (decoration) + { + case DecorationHlslSemanticGOOGLE: + dec.hlsl_semantic = argument; + break; + + default: + break; + } +} + +void ParsedIR::set_decoration(uint32_t id, Decoration decoration, uint32_t argument) +{ + auto &dec = meta[id].decoration; + dec.decoration_flags.set(decoration); + + switch (decoration) + { + case DecorationBuiltIn: + dec.builtin = true; + dec.builtin_type = static_cast(argument); + break; + + case DecorationLocation: + dec.location = argument; + break; + + case DecorationComponent: + dec.component = argument; + break; + + case DecorationOffset: + dec.offset = argument; + break; + + case DecorationArrayStride: + dec.array_stride = argument; + break; + + case DecorationMatrixStride: + dec.matrix_stride = argument; + break; + + case DecorationBinding: + dec.binding = argument; + break; + + case DecorationDescriptorSet: + dec.set = argument; + break; + + case DecorationInputAttachmentIndex: + dec.input_attachment = argument; + break; + + case DecorationSpecId: + dec.spec_id = argument; + break; + + case DecorationIndex: + dec.index = argument; + break; + + case DecorationHlslCounterBufferGOOGLE: + meta[id].hlsl_magic_counter_buffer = argument; + meta[id].hlsl_is_magic_counter_buffer = true; + break; + + default: + break; + } +} + +void ParsedIR::set_member_decoration(uint32_t id, uint32_t index, Decoration decoration, uint32_t argument) +{ + meta[id].members.resize(max(meta[id].members.size(), size_t(index) + 1)); + auto &dec = meta[id].members[index]; + dec.decoration_flags.set(decoration); + + switch (decoration) + { + case DecorationBuiltIn: + dec.builtin = true; + dec.builtin_type = static_cast(argument); + break; + + case DecorationLocation: + dec.location = argument; + break; + + case DecorationComponent: + dec.component = argument; + break; + + case DecorationBinding: + dec.binding = argument; + break; + + case DecorationOffset: + dec.offset = argument; + break; + + case DecorationSpecId: + dec.spec_id = argument; + break; + + case DecorationMatrixStride: + dec.matrix_stride = argument; + break; + + case DecorationIndex: + dec.index = argument; + break; + + default: + break; + } +} + +// Recursively marks any constants referenced by the specified constant instruction as being used +// as an array length. The id must be a constant instruction (SPIRConstant or SPIRConstantOp). +void ParsedIR::mark_used_as_array_length(uint32_t id) +{ + switch (ids[id].get_type()) + { + case TypeConstant: + get(id).is_used_as_array_length = true; + break; + + case TypeConstantOp: + { + auto &cop = get(id); + for (uint32_t arg_id : cop.arguments) + mark_used_as_array_length(arg_id); + break; + } + + case TypeUndef: + break; + + default: + assert(0); + } +} + +Bitset ParsedIR::get_buffer_block_flags(const SPIRVariable &var) const +{ + auto &type = get(var.basetype); + assert(type.basetype == SPIRType::Struct); + + // Some flags like non-writable, non-readable are actually found + // as member decorations. If all members have a decoration set, propagate + // the decoration up as a regular variable decoration. + Bitset base_flags = meta[var.self].decoration.decoration_flags; + + if (type.member_types.empty()) + return base_flags; + + Bitset all_members_flags = get_member_decoration_bitset(type.self, 0); + for (uint32_t i = 1; i < uint32_t(type.member_types.size()); i++) + all_members_flags.merge_and(get_member_decoration_bitset(type.self, i)); + + base_flags.merge_or(all_members_flags); + return base_flags; +} + +const Bitset &ParsedIR::get_member_decoration_bitset(uint32_t id, uint32_t index) const +{ + auto &m = meta[id]; + if (index >= m.members.size()) + { + static const Bitset cleared = {}; + return cleared; + } + + return m.members[index].decoration_flags; +} + +bool ParsedIR::has_decoration(uint32_t id, Decoration decoration) const +{ + return get_decoration_bitset(id).get(decoration); +} + +uint32_t ParsedIR::get_decoration(uint32_t id, Decoration decoration) const +{ + auto &dec = meta[id].decoration; + if (!dec.decoration_flags.get(decoration)) + return 0; + + switch (decoration) + { + case DecorationBuiltIn: + return dec.builtin_type; + case DecorationLocation: + return dec.location; + case DecorationComponent: + return dec.component; + case DecorationOffset: + return dec.offset; + case DecorationBinding: + return dec.binding; + case DecorationDescriptorSet: + return dec.set; + case DecorationInputAttachmentIndex: + return dec.input_attachment; + case DecorationSpecId: + return dec.spec_id; + case DecorationArrayStride: + return dec.array_stride; + case DecorationMatrixStride: + return dec.matrix_stride; + case DecorationIndex: + return dec.index; + default: + return 1; + } +} + +const string &ParsedIR::get_decoration_string(uint32_t id, Decoration decoration) const +{ + auto &dec = meta[id].decoration; + static const string empty; + + if (!dec.decoration_flags.get(decoration)) + return empty; + + switch (decoration) + { + case DecorationHlslSemanticGOOGLE: + return dec.hlsl_semantic; + + default: + return empty; + } +} + +void ParsedIR::unset_decoration(uint32_t id, Decoration decoration) +{ + auto &dec = meta[id].decoration; + dec.decoration_flags.clear(decoration); + switch (decoration) + { + case DecorationBuiltIn: + dec.builtin = false; + break; + + case DecorationLocation: + dec.location = 0; + break; + + case DecorationComponent: + dec.component = 0; + break; + + case DecorationOffset: + dec.offset = 0; + break; + + case DecorationBinding: + dec.binding = 0; + break; + + case DecorationDescriptorSet: + dec.set = 0; + break; + + case DecorationInputAttachmentIndex: + dec.input_attachment = 0; + break; + + case DecorationSpecId: + dec.spec_id = 0; + break; + + case DecorationHlslSemanticGOOGLE: + dec.hlsl_semantic.clear(); + break; + + case DecorationHlslCounterBufferGOOGLE: + { + auto &counter = meta[id].hlsl_magic_counter_buffer; + if (counter) + { + meta[counter].hlsl_is_magic_counter_buffer = false; + counter = 0; + } + break; + } + + default: + break; + } +} + +bool ParsedIR::has_member_decoration(uint32_t id, uint32_t index, Decoration decoration) const +{ + return get_member_decoration_bitset(id, index).get(decoration); +} + +uint32_t ParsedIR::get_member_decoration(uint32_t id, uint32_t index, Decoration decoration) const +{ + auto &m = meta[id]; + if (index >= m.members.size()) + return 0; + + auto &dec = m.members[index]; + if (!dec.decoration_flags.get(decoration)) + return 0; + + switch (decoration) + { + case DecorationBuiltIn: + return dec.builtin_type; + case DecorationLocation: + return dec.location; + case DecorationComponent: + return dec.component; + case DecorationBinding: + return dec.binding; + case DecorationOffset: + return dec.offset; + case DecorationSpecId: + return dec.spec_id; + case DecorationIndex: + return dec.index; + default: + return 1; + } +} + +const Bitset &ParsedIR::get_decoration_bitset(uint32_t id) const +{ + auto &dec = meta[id].decoration; + return dec.decoration_flags; +} + +void ParsedIR::set_member_decoration_string(uint32_t id, uint32_t index, Decoration decoration, const string &argument) +{ + meta[id].members.resize(max(meta[id].members.size(), size_t(index) + 1)); + auto &dec = meta[id].members[index]; + dec.decoration_flags.set(decoration); + + switch (decoration) + { + case DecorationHlslSemanticGOOGLE: + dec.hlsl_semantic = argument; + break; + + default: + break; + } +} + +const string &ParsedIR::get_member_decoration_string(uint32_t id, uint32_t index, Decoration decoration) const +{ + static const string empty; + auto &m = meta[id]; + + if (!has_member_decoration(id, index, decoration)) + return empty; + + auto &dec = m.members[index]; + + switch (decoration) + { + case DecorationHlslSemanticGOOGLE: + return dec.hlsl_semantic; + + default: + return empty; + } +} + +void ParsedIR::unset_member_decoration(uint32_t id, uint32_t index, Decoration decoration) +{ + auto &m = meta[id]; + if (index >= m.members.size()) + return; + + auto &dec = m.members[index]; + + dec.decoration_flags.clear(decoration); + switch (decoration) + { + case DecorationBuiltIn: + dec.builtin = false; + break; + + case DecorationLocation: + dec.location = 0; + break; + + case DecorationComponent: + dec.component = 0; + break; + + case DecorationOffset: + dec.offset = 0; + break; + + case DecorationSpecId: + dec.spec_id = 0; + break; + + case DecorationHlslSemanticGOOGLE: + dec.hlsl_semantic.clear(); + break; + + default: + break; + } +} + +uint32_t ParsedIR::increase_bound_by(uint32_t incr_amount) +{ + auto curr_bound = ids.size(); + auto new_bound = curr_bound + incr_amount; + ids.resize(new_bound); + meta.resize(new_bound); + block_meta.resize(new_bound); + return uint32_t(curr_bound); +} + +} // namespace spirv_cross \ No newline at end of file diff --git a/spirv_cross_parsed_ir.hpp b/spirv_cross_parsed_ir.hpp new file mode 100644 index 000000000..80283cfdc --- /dev/null +++ b/spirv_cross_parsed_ir.hpp @@ -0,0 +1,129 @@ +/* + * Copyright 2018 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. + */ + +#ifndef SPIRV_CROSS_PARSED_IR_HPP +#define SPIRV_CROSS_PARSED_IR_HPP + +#include "spirv_common.hpp" +#include +#include +#include + +namespace spirv_cross +{ + +// This data structure holds all information needed to perform cross-compilation and reflection. +// It is the output of the Parser, but any implementation could create this structure. +// It is intentionally very "open" and struct-like with some helper functions to deal with decorations. +// Parser is the reference implementation of how this data structure should be filled in. + +class ParsedIR +{ +public: + // Resizes ids, meta and block_meta. + void set_id_bounds(uint32_t bounds); + + // The raw SPIR-V, instructions and opcodes refer to this by offset + count. + std::vector spirv; + + // Holds various data structures which inherit from IVariant. + std::vector ids; + + // Various meta data for IDs, decorations, names, etc. + std::vector meta; + + // Declared capabilities and extensions in the SPIR-V module. + // Not really used except for reflection at the moment. + std::vector declared_capabilities; + std::vector declared_extensions; + + // Meta data about blocks. The cross-compiler needs to query if a block is either of these types. + // It is a bitset as there can be more than one tag per block. + enum BlockMetaFlagBits + { + BLOCK_META_LOOP_HEADER_BIT = 1 << 0, + BLOCK_META_CONTINUE_BIT = 1 << 1, + BLOCK_META_LOOP_MERGE_BIT = 1 << 2, + BLOCK_META_SELECTION_MERGE_BIT = 1 << 3, + BLOCK_META_MULTISELECT_MERGE_BIT = 1 << 4 + }; + using BlockMetaFlags = uint8_t; + std::vector block_meta; + std::unordered_map continue_block_to_loop_header; + + // 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; + uint32_t default_entry_point = 0; + + struct Source + { + uint32_t version = 0; + bool es = false; + bool known = false; + bool hlsl = false; + + Source() = default; + }; + + Source source; + + // Decoration handling methods. + // Can be useful for simple "raw" reflection. + // However, most members are here because the Parser needs most of these, + // and might as well just have the whole suite of decoration/name handling in one place. + void set_name(uint32_t id, const std::string &name); + const std::string &get_name(uint32_t id) const; + void set_decoration(uint32_t id, spv::Decoration decoration, uint32_t argument = 0); + void set_decoration_string(uint32_t id, spv::Decoration decoration, const std::string &argument); + bool has_decoration(uint32_t id, spv::Decoration decoration) const; + uint32_t get_decoration(uint32_t id, spv::Decoration decoration) const; + const std::string &get_decoration_string(uint32_t id, spv::Decoration decoration) const; + const Bitset &get_decoration_bitset(uint32_t id) const; + void unset_decoration(uint32_t id, spv::Decoration decoration); + + // Decoration handling methods (for members of a struct). + void set_member_name(uint32_t id, uint32_t index, const std::string &name); + const std::string &get_member_name(uint32_t id, uint32_t index) const; + void set_member_decoration(uint32_t id, uint32_t index, spv::Decoration decoration, uint32_t argument = 0); + void set_member_decoration_string(uint32_t id, uint32_t index, spv::Decoration decoration, + const std::string &argument); + uint32_t get_member_decoration(uint32_t id, uint32_t index, spv::Decoration decoration) const; + const std::string &get_member_decoration_string(uint32_t id, uint32_t index, spv::Decoration decoration) const; + bool has_member_decoration(uint32_t id, uint32_t index, spv::Decoration decoration) const; + const Bitset &get_member_decoration_bitset(uint32_t id, uint32_t index) const; + void unset_member_decoration(uint32_t id, uint32_t index, spv::Decoration decoration); + + void mark_used_as_array_length(uint32_t id); + uint32_t increase_bound_by(uint32_t count); + Bitset get_buffer_block_flags(const SPIRVariable &var) const; + +private: + template + T &get(uint32_t id) + { + return variant_get(ids[id]); + } + + template + const T &get(uint32_t id) const + { + return variant_get(ids[id]); + } +}; +} // namespace spirv_cross + +#endif \ No newline at end of file diff --git a/spirv_glsl.cpp b/spirv_glsl.cpp index 04bef7924..113dc0f1c 100644 --- a/spirv_glsl.cpp +++ b/spirv_glsl.cpp @@ -290,7 +290,7 @@ void CompilerGLSL::reset() block_ssbo_names.clear(); function_overloads.clear(); - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { @@ -343,7 +343,7 @@ void CompilerGLSL::remap_pls_variables() void CompilerGLSL::find_static_extensions() { - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeType) { @@ -443,7 +443,7 @@ string CompilerGLSL::compile() emit_header(); emit_resources(); - emit_function(get(entry_point), Bitset()); + emit_function(get(ir.default_entry_point), Bitset()); pass_count++; } while (force_recompile); @@ -738,12 +738,12 @@ string CompilerGLSL::layout_for_member(const SPIRType &type, uint32_t index) if (is_legacy()) return ""; - bool is_block = meta[type.self].decoration.decoration_flags.get(DecorationBlock) || - meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); + bool is_block = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock) || + ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); if (!is_block) return ""; - auto &memb = meta[type.self].members; + auto &memb = ir.meta[type.self].members; if (index >= memb.size()) return ""; auto &dec = memb[index]; @@ -938,7 +938,7 @@ uint32_t CompilerGLSL::type_to_packed_alignment(const SPIRType &type, const Bits uint32_t alignment = 0; for (uint32_t i = 0; i < type.member_types.size(); i++) { - auto member_flags = meta[type.self].members.at(i).decoration_flags; + auto member_flags = ir.meta[type.self].members.at(i).decoration_flags; alignment = max(alignment, type_to_packed_alignment(get(type.member_types[i]), member_flags, packing)); } @@ -1043,7 +1043,7 @@ uint32_t CompilerGLSL::type_to_packed_size(const SPIRType &type, const Bitset &f for (uint32_t i = 0; i < type.member_types.size(); i++) { - auto member_flags = meta[type.self].members.at(i).decoration_flags; + auto member_flags = ir.meta[type.self].members.at(i).decoration_flags; auto &member_type = get(type.member_types[i]); uint32_t packed_alignment = type_to_packed_alignment(member_type, member_flags, packing); @@ -1113,7 +1113,7 @@ bool CompilerGLSL::buffer_is_packing_standard(const SPIRType &type, BufferPackin for (uint32_t i = 0; i < type.member_types.size(); i++) { auto &memb_type = get(type.member_types[i]); - auto member_flags = meta[type.self].members.at(i).decoration_flags; + auto member_flags = ir.meta[type.self].members.at(i).decoration_flags; // Verify alignment rules. uint32_t packed_alignment = type_to_packed_alignment(memb_type, member_flags, packing); @@ -1221,10 +1221,10 @@ string CompilerGLSL::layout_for_variable(const SPIRVariable &var) vector attr; - auto &dec = meta[var.self].decoration; + auto &dec = ir.meta[var.self].decoration; auto &type = get(var.basetype); auto flags = dec.decoration_flags; - auto typeflags = meta[type.self].decoration.decoration_flags; + auto typeflags = ir.meta[type.self].decoration.decoration_flags; if (options.vulkan_semantics && var.storage == StorageClassPushConstant) attr.push_back("push_constant"); @@ -1244,7 +1244,7 @@ string CompilerGLSL::layout_for_variable(const SPIRVariable &var) if (flags.get(DecorationLocation) && can_use_io_location(var.storage, is_block)) { Bitset combined_decoration; - for (uint32_t i = 0; i < meta[type.self].members.size(); i++) + for (uint32_t i = 0; i < ir.meta[type.self].members.size(); i++) combined_decoration.merge_or(combined_decoration_for_member(type, i)); // If our members have location decorations, we don't need to @@ -1411,7 +1411,7 @@ void CompilerGLSL::emit_push_constant_block_glsl(const SPIRVariable &var) // OpenGL has no concept of push constant blocks, implement it as a uniform struct. auto &type = get(var.basetype); - auto &flags = meta[var.self].decoration.decoration_flags; + auto &flags = ir.meta[var.self].decoration.decoration_flags; flags.clear(DecorationBinding); flags.clear(DecorationDescriptorSet); @@ -1423,7 +1423,7 @@ void CompilerGLSL::emit_push_constant_block_glsl(const SPIRVariable &var) // We're emitting the push constant block as a regular struct, so disable the block qualifier temporarily. // Otherwise, we will end up emitting layout() qualifiers on naked structs which is not allowed. - auto &block_flags = meta[type.self].decoration.decoration_flags; + auto &block_flags = ir.meta[type.self].decoration.decoration_flags; bool block_flag = block_flags.get(DecorationBlock); block_flags.clear(DecorationBlock); @@ -1450,13 +1450,13 @@ void CompilerGLSL::emit_buffer_block_legacy(const SPIRVariable &var) { auto &type = get(var.basetype); bool ssbo = var.storage == StorageClassStorageBuffer || - meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); + ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); if (ssbo) SPIRV_CROSS_THROW("SSBOs not supported in legacy targets."); // We're emitting the push constant block as a regular struct, so disable the block qualifier temporarily. // Otherwise, we will end up emitting layout() qualifiers on naked structs which is not allowed. - auto &block_flags = meta[type.self].decoration.decoration_flags; + auto &block_flags = ir.meta[type.self].decoration.decoration_flags; bool block_flag = block_flags.get(DecorationBlock); block_flags.clear(DecorationBlock); emit_struct(type); @@ -1470,9 +1470,9 @@ void CompilerGLSL::emit_buffer_block_native(const SPIRVariable &var) { auto &type = get(var.basetype); - Bitset flags = get_buffer_block_flags(var); + Bitset flags = ir.get_buffer_block_flags(var); bool ssbo = var.storage == StorageClassStorageBuffer || - meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); + ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); bool is_restrict = ssbo && flags.get(DecorationRestrict); bool is_writeonly = ssbo && flags.get(DecorationNonReadable); bool is_readonly = ssbo && flags.get(DecorationNonWritable); @@ -1485,7 +1485,7 @@ void CompilerGLSL::emit_buffer_block_native(const SPIRVariable &var) // Shaders never use the block by interface name, so we don't // have to track this other than updating name caches. - if (meta[type.self].decoration.alias.empty() || block_namespace.find(buffer_name) != end(block_namespace)) + if (ir.meta[type.self].decoration.alias.empty() || block_namespace.find(buffer_name) != end(block_namespace)) buffer_name = get_block_fallback_name(var.self); // Make sure we get something unique. @@ -1540,7 +1540,7 @@ void CompilerGLSL::emit_buffer_block_flattened(const SPIRVariable &var) if (basic_type != SPIRType::Float && basic_type != SPIRType::Int && basic_type != SPIRType::UInt) SPIRV_CROSS_THROW("Basic types in a flattened UBO must be float, int or uint."); - auto flags = get_buffer_block_flags(var); + auto flags = ir.get_buffer_block_flags(var); statement("uniform ", flags_to_precision_qualifiers_glsl(tmp, flags), type_to_glsl(tmp), " ", buffer_name, "[", buffer_size, "];"); } @@ -1576,9 +1576,9 @@ void CompilerGLSL::emit_flattened_io_block(const SPIRVariable &var, const char * if (!type.array.empty()) SPIRV_CROSS_THROW("Array of varying structs cannot be flattened to legacy-compatible varyings."); - auto old_flags = meta[type.self].decoration.decoration_flags; + auto old_flags = ir.meta[type.self].decoration.decoration_flags; // Emit the members as if they are part of a block to get all qualifiers. - meta[type.self].decoration.decoration_flags.set(DecorationBlock); + ir.meta[type.self].decoration.decoration_flags.set(DecorationBlock); type.member_name_cache.clear(); @@ -1604,7 +1604,7 @@ void CompilerGLSL::emit_flattened_io_block(const SPIRVariable &var, const char * i++; } - meta[type.self].decoration.decoration_flags = old_flags; + ir.meta[type.self].decoration.decoration_flags = old_flags; // Treat this variable as flattened from now on. flattened_structs.insert(var.self); @@ -1615,7 +1615,7 @@ void CompilerGLSL::emit_interface_block(const SPIRVariable &var) auto &type = get(var.basetype); // Either make it plain in/out or in/out blocks depending on what shader is doing ... - bool block = meta[type.self].decoration.decoration_flags.get(DecorationBlock); + bool block = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock); const char *qual = to_storage_qualifiers_glsl(var); if (block) @@ -1800,14 +1800,14 @@ void CompilerGLSL::replace_illegal_names() }; // clang-format on - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { auto &var = id.get(); if (!is_hidden_variable(var)) { - auto &m = meta[var.self].decoration; + auto &m = ir.meta[var.self].decoration; if (m.alias.compare(0, 3, "gl_") == 0 || keywords.find(m.alias) != end(keywords)) m.alias = join("_", m.alias); } @@ -1817,7 +1817,7 @@ void CompilerGLSL::replace_illegal_names() void CompilerGLSL::replace_fragment_output(SPIRVariable &var) { - auto &m = meta[var.self].decoration; + auto &m = ir.meta[var.self].decoration; uint32_t location = 0; if (m.decoration_flags.get(DecorationLocation)) location = m.location; @@ -1855,7 +1855,7 @@ void CompilerGLSL::replace_fragment_output(SPIRVariable &var) void CompilerGLSL::replace_fragment_outputs() { - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { @@ -1925,7 +1925,7 @@ void CompilerGLSL::emit_pls() void CompilerGLSL::fixup_image_load_store_access() { - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() != TypeVariable) continue; @@ -1938,7 +1938,7 @@ void CompilerGLSL::fixup_image_load_store_access() // Solve this by making the image access as restricted as possible and loosen up if we need to. // If any no-read/no-write flags are actually set, assume that the compiler knows what it's doing. - auto &flags = meta.at(var).decoration.decoration_flags; + auto &flags = ir.meta.at(var).decoration.decoration_flags; if (!flags.get(DecorationNonWritable) && !flags.get(DecorationNonReadable)) { flags.set(DecorationNonWritable); @@ -1961,7 +1961,7 @@ void CompilerGLSL::emit_declared_builtin_block(StorageClass storage, ExecutionMo uint32_t cull_distance_size = 0; uint32_t clip_distance_size = 0; - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() != TypeVariable) continue; @@ -1974,7 +1974,7 @@ void CompilerGLSL::emit_declared_builtin_block(StorageClass storage, ExecutionMo if (var.storage == storage && block && is_builtin_variable(var)) { uint32_t index = 0; - for (auto &m : meta[type.self].members) + for (auto &m : ir.meta[type.self].members) { if (m.builtin) { @@ -1990,7 +1990,7 @@ void CompilerGLSL::emit_declared_builtin_block(StorageClass storage, ExecutionMo else if (var.storage == storage && !block && is_builtin_variable(var)) { // While we're at it, collect all declared global builtins (HLSL mostly ...). - auto &m = meta[var.self].decoration; + auto &m = ir.meta[var.self].decoration; if (m.builtin) { global_builtins.set(m.builtin_type); @@ -2062,7 +2062,7 @@ void CompilerGLSL::emit_declared_builtin_block(StorageClass storage, ExecutionMo void CompilerGLSL::declare_undefined_values() { bool emitted = false; - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() != TypeUndef) continue; @@ -2138,7 +2138,7 @@ 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. - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeConstant) { @@ -2165,14 +2165,14 @@ void CompilerGLSL::emit_resources() // Output all basic struct types which are not Block or BufferBlock as these are declared inplace // when such variables are instantiated. - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeType) { auto &type = id.get(); if (type.basetype == SPIRType::Struct && type.array.empty() && !type.pointer && - (!meta[type.self].decoration.decoration_flags.get(DecorationBlock) && - !meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock))) + (!ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock) && + !ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock))) { emit_struct(type); } @@ -2180,7 +2180,7 @@ void CompilerGLSL::emit_resources() } // Output UBOs and SSBOs - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { @@ -2188,8 +2188,8 @@ void CompilerGLSL::emit_resources() auto &type = get(var.basetype); bool is_block_storage = type.storage == StorageClassStorageBuffer || type.storage == StorageClassUniform; - bool has_block_flags = meta[type.self].decoration.decoration_flags.get(DecorationBlock) || - meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); + bool has_block_flags = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock) || + ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); if (var.storage != StorageClassFunction && type.pointer && is_block_storage && !is_hidden_variable(var) && has_block_flags) @@ -2200,7 +2200,7 @@ void CompilerGLSL::emit_resources() } // Output push constant blocks - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { @@ -2217,7 +2217,7 @@ void CompilerGLSL::emit_resources() bool skip_separate_image_sampler = !combined_image_samplers.empty() || !options.vulkan_semantics; // Output Uniform Constants (values, samplers, images, etc). - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { @@ -2250,7 +2250,7 @@ void CompilerGLSL::emit_resources() emitted = false; // Output in/out interfaces. - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { @@ -2269,7 +2269,7 @@ void CompilerGLSL::emit_resources() // For gl_InstanceIndex emulation on GLES, the API user needs to // supply this uniform. if (options.vertex.support_nonzero_base_instance && - meta[var.self].decoration.builtin_type == BuiltInInstanceIndex && !options.vulkan_semantics) + ir.meta[var.self].decoration.builtin_type == BuiltInInstanceIndex && !options.vulkan_semantics) { statement("uniform int SPIRV_Cross_BaseInstance;"); emitted = true; @@ -2438,7 +2438,7 @@ string CompilerGLSL::to_expression(uint32_t id) if (itr != end(invalid_expressions)) handle_invalid_expression(id); - if (ids[id].get_type() == TypeExpression) + if (ir.ids[id].get_type() == TypeExpression) { // We might have a more complex chain of dependencies. // A possible scenario is that we @@ -2459,7 +2459,7 @@ string CompilerGLSL::to_expression(uint32_t id) track_expression_read(id); - switch (ids[id].get_type()) + switch (ir.ids[id].get_type()) { case TypeExpression: { @@ -2491,7 +2491,7 @@ string CompilerGLSL::to_expression(uint32_t id) auto &type = get(c.constant_type); // WorkGroupSize may be a constant. - auto &dec = meta[c.self].decoration; + auto &dec = ir.meta[c.self].decoration; if (dec.builtin) return builtin_to_glsl(dec.builtin_type, StorageClassGeneric); else if (c.specialization && options.vulkan_semantics) @@ -2530,7 +2530,7 @@ string CompilerGLSL::to_expression(uint32_t id) } else { - auto &dec = meta[var.self].decoration; + auto &dec = ir.meta[var.self].decoration; if (dec.builtin) return builtin_to_glsl(dec.builtin_type, var.storage); else @@ -3270,7 +3270,7 @@ string CompilerGLSL::constant_expression_vector(const SPIRConstant &c, uint32_t string CompilerGLSL::declare_temporary(uint32_t result_type, uint32_t result_id) { auto &type = get(result_type); - auto flags = meta[result_id].decoration.decoration_flags; + auto flags = ir.meta[result_id].decoration.decoration_flags; // If we're declaring temporaries inside continue blocks, // we must declare the temporary in the loop header so that the continue block can avoid declaring new variables. @@ -3863,13 +3863,10 @@ void CompilerGLSL::emit_sampled_image_op(uint32_t result_type, uint32_t result_i void CompilerGLSL::emit_texture_op(const Instruction &i) { - auto ops = stream(i); + auto *ops = stream(i); auto op = static_cast(i.op); uint32_t length = i.length; - if (i.offset + length > spirv.size()) - SPIRV_CROSS_THROW("Compiler::parse() opcode out of range."); - vector inherited_expressions; uint32_t result_type = ops[0]; @@ -4365,7 +4362,7 @@ void CompilerGLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop, { forced_temporaries.insert(id); auto &type = get(result_type); - auto flags = meta[id].decoration.decoration_flags; + auto flags = ir.meta[id].decoration.decoration_flags; statement(flags_to_precision_qualifiers_glsl(type, flags), variable_decl(type, to_name(id)), ";"); set(id, to_name(id), result_type, true); @@ -4488,7 +4485,7 @@ void CompilerGLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop, { forced_temporaries.insert(id); auto &type = get(result_type); - auto flags = meta[id].decoration.decoration_flags; + auto flags = ir.meta[id].decoration.decoration_flags; statement(flags_to_precision_qualifiers_glsl(type, flags), variable_decl(type, to_name(id)), ";"); set(id, to_name(id), result_type, true); @@ -5306,7 +5303,7 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice // but HLSL seems to just emit straight arrays here. // We must pretend this access goes through gl_in/gl_out arrays // to be able to access certain builtins as arrays. - auto builtin = meta[base].decoration.builtin_type; + auto builtin = ir.meta[base].decoration.builtin_type; switch (builtin) { // case BuiltInCullDistance: // These are already arrays, need to figure out rules for these in tess/geom. @@ -5426,7 +5423,7 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice expr += "."; expr += index_to_swizzle(index); } - else if (ids[index].get_type() == TypeConstant && !is_packed) + else if (ir.ids[index].get_type() == TypeConstant && !is_packed) { auto &c = get(index); expr += "."; @@ -6486,7 +6483,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) splat = in_type->vecsize == 1 && in_type->columns == 1 && !composite && backend.use_constructor_splatting; swizzle_splat = in_type->vecsize == 1 && in_type->columns == 1 && backend.can_swizzle_scalar; - if (ids[elems[0]].get_type() == TypeConstant && !type_is_floating_point(*in_type)) + if (ir.ids[elems[0]].get_type() == TypeConstant && !type_is_floating_point(*in_type)) { // Cannot swizzle literal integers as a special case. swizzle_splat = false; @@ -6519,7 +6516,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) // We cannot construct array of arrays because we cannot treat the inputs // as value types. Need to declare the array-of-arrays, and copy in elements one by one. forced_temporaries.insert(id); - auto flags = meta[id].decoration.decoration_flags; + auto flags = ir.meta[id].decoration.decoration_flags; statement(flags_to_precision_qualifiers_glsl(out_type, flags), variable_decl(out_type, to_name(id)), ";"); set(id, to_name(id), result_type, true); for (uint32_t i = 0; i < length; i++) @@ -6906,7 +6903,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) case OpBitwiseXor: { auto type = get(ops[0]).basetype; - GLSL_BOP_CAST (^, type); + GLSL_BOP_CAST(^, type); break; } @@ -7580,7 +7577,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) auto *var = maybe_get_backing_variable(ops[2]); if (var) { - auto &flags = meta.at(var->self).decoration.decoration_flags; + auto &flags = ir.meta.at(var->self).decoration.decoration_flags; if (flags.get(DecorationNonReadable)) { flags.clear(DecorationNonReadable); @@ -7728,7 +7725,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) auto *var = maybe_get_backing_variable(ops[0]); if (var) { - auto &flags = meta.at(var->self).decoration.decoration_flags; + auto &flags = ir.meta.at(var->self).decoration.decoration_flags; if (flags.get(DecorationNonWritable)) { flags.clear(DecorationNonWritable); @@ -8225,7 +8222,7 @@ void CompilerGLSL::append_global_func_args(const SPIRFunction &func, uint32_t in string CompilerGLSL::to_member_name(const SPIRType &type, uint32_t index) { - auto &memb = meta[type.self].members; + auto &memb = ir.meta[type.self].members; if (index < memb.size() && !memb[index].alias.empty()) return memb[index].alias; else @@ -8239,7 +8236,7 @@ string CompilerGLSL::to_member_reference(const SPIRVariable *, const SPIRType &t void CompilerGLSL::add_member_name(SPIRType &type, uint32_t index) { - auto &memb = meta[type.self].members; + auto &memb = ir.meta[type.self].members; if (index < memb.size() && !memb[index].alias.empty()) { auto &name = memb[index].alias; @@ -8266,7 +8263,7 @@ bool CompilerGLSL::is_non_native_row_major_matrix(uint32_t id) return false; // Non-matrix or column-major matrix types do not need to be converted. - if (!meta[id].decoration.decoration_flags.get(DecorationRowMajor)) + if (!ir.meta[id].decoration.decoration_flags.get(DecorationRowMajor)) return false; // Only square row-major matrices can be converted at this time. @@ -8332,13 +8329,13 @@ void CompilerGLSL::emit_struct_member(const SPIRType &type, uint32_t member_type auto &membertype = get(member_type_id); Bitset memberflags; - auto &memb = meta[type.self].members; + auto &memb = ir.meta[type.self].members; if (index < memb.size()) memberflags = memb[index].decoration_flags; string qualifiers; - bool is_block = meta[type.self].decoration.decoration_flags.get(DecorationBlock) || - meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); + bool is_block = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock) || + ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); if (is_block) qualifiers = to_interpolation_qualifiers(memberflags); @@ -8402,12 +8399,12 @@ const char *CompilerGLSL::flags_to_precision_qualifiers_glsl(const SPIRType &typ const char *CompilerGLSL::to_precision_qualifiers_glsl(uint32_t id) { - return flags_to_precision_qualifiers_glsl(expression_type(id), meta[id].decoration.decoration_flags); + return flags_to_precision_qualifiers_glsl(expression_type(id), ir.meta[id].decoration.decoration_flags); } string CompilerGLSL::to_qualifiers_glsl(uint32_t id) { - auto flags = meta[id].decoration.decoration_flags; + auto flags = ir.meta[id].decoration.decoration_flags; string res; auto *var = maybe_get(id); @@ -8469,13 +8466,13 @@ string CompilerGLSL::variable_decl(const SPIRVariable &variable) if (variable.loop_variable && variable.static_expression) { uint32_t expr = variable.static_expression; - if (ids[expr].get_type() != TypeUndef) + if (ir.ids[expr].get_type() != TypeUndef) res += join(" = ", to_expression(variable.static_expression)); } else if (variable.initializer) { uint32_t expr = variable.initializer; - if (ids[expr].get_type() != TypeUndef) + if (ir.ids[expr].get_type() != TypeUndef) res += join(" = ", to_initializer_expression(variable)); } return res; @@ -8483,7 +8480,7 @@ string CompilerGLSL::variable_decl(const SPIRVariable &variable) const char *CompilerGLSL::to_pls_qualifiers_glsl(const SPIRVariable &variable) { - auto flags = meta[variable.self].decoration.decoration_flags; + auto flags = ir.meta[variable.self].decoration.decoration_flags; if (flags.get(DecorationRelaxedPrecision)) return "mediump "; else @@ -8838,7 +8835,7 @@ void CompilerGLSL::add_variable(unordered_set &variables, string &name) void CompilerGLSL::add_variable(unordered_set &variables, uint32_t id) { - auto &name = meta[id].decoration.alias; + auto &name = ir.meta[id].decoration.alias; add_variable(variables, name); } @@ -8883,7 +8880,7 @@ void CompilerGLSL::flatten_buffer_block(uint32_t id) auto &var = get(id); auto &type = get(var.basetype); auto name = to_name(type.self, false); - auto flags = meta.at(type.self).decoration.decoration_flags; + auto flags = ir.meta.at(type.self).decoration.decoration_flags; if (!type.array.empty()) SPIRV_CROSS_THROW(name + " is an array of UBOs."); @@ -8908,7 +8905,7 @@ bool CompilerGLSL::check_atomic_image(uint32_t id) auto *var = maybe_get_backing_variable(id); if (var) { - auto &flags = meta.at(var->self).decoration.decoration_flags; + auto &flags = ir.meta.at(var->self).decoration.decoration_flags; if (flags.get(DecorationNonWritable) || flags.get(DecorationNonReadable)) { flags.clear(DecorationNonWritable); @@ -8978,7 +8975,7 @@ void CompilerGLSL::add_function_overload(const SPIRFunction &func) void CompilerGLSL::emit_function_prototype(SPIRFunction &func, const Bitset &return_flags) { - if (func.self != entry_point) + if (func.self != ir.default_entry_point) add_function_overload(func); // Avoid shadow declarations. @@ -8992,7 +8989,7 @@ void CompilerGLSL::emit_function_prototype(SPIRFunction &func, const Bitset &ret decl += type_to_array_glsl(type); decl += " "; - if (func.self == entry_point) + if (func.self == ir.default_entry_point) { decl += "main"; processing_entry_point = true; @@ -9064,7 +9061,7 @@ void CompilerGLSL::emit_function(SPIRFunction &func, const Bitset &return_flags) { // Recursively emit functions which are called. uint32_t id = ops[2]; - emit_function(get(id), meta[ops[1]].decoration.decoration_flags); + emit_function(get(id), ir.meta[ops[1]].decoration.decoration_flags); } } } @@ -9072,7 +9069,7 @@ void CompilerGLSL::emit_function(SPIRFunction &func, const Bitset &return_flags) emit_function_prototype(func, return_flags); begin_scope(); - if (func.self == entry_point) + if (func.self == ir.default_entry_point) emit_entry_point_declarations(); current_function = &func; @@ -9272,7 +9269,7 @@ void CompilerGLSL::branch(uint32_t from, uint32_t to) flush_all_active_variables(); // This is only a continue if we branch to our loop dominator. - if (loop_blocks.find(to) != end(loop_blocks) && get(from).loop_dominator == to) + if ((ir.block_meta[to] & ParsedIR::BLOCK_META_LOOP_HEADER_BIT) != 0 && get(from).loop_dominator == to) { // This can happen if we had a complex continue block which was emitted. // Once the continue block tries to branch to the loop header, just emit continue; @@ -9426,7 +9423,7 @@ string CompilerGLSL::emit_continue_block(uint32_t continue_block) redirect_statement = &statements; // Stamp out all blocks one after each other. - while (loop_blocks.find(block->self) == end(loop_blocks)) + while ((ir.block_meta[block->self] & ParsedIR::BLOCK_META_LOOP_HEADER_BIT) == 0) { propagate_loop_dominators(*block); // Write out all instructions we have in this block. @@ -9488,7 +9485,7 @@ string CompilerGLSL::emit_for_loop_initializers(const SPIRBlock &block) // 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) + if (expr == 0 || ir.ids[expr].get_type() == TypeUndef) missing_initializers++; } @@ -9511,7 +9508,7 @@ string CompilerGLSL::emit_for_loop_initializers(const SPIRBlock &block) for (auto &loop_var : block.loop_variables) { uint32_t static_expr = get(loop_var).static_expression; - if (static_expr == 0 || ids[static_expr].get_type() == TypeUndef) + if (static_expr == 0 || ir.ids[static_expr].get_type() == TypeUndef) { statement(variable_decl(get(loop_var)), ";"); } @@ -9546,7 +9543,7 @@ bool CompilerGLSL::for_loop_initializers_are_same_type(const SPIRBlock &block) { // Don't care about uninitialized variables as they will not be part of the initializers. uint32_t expr = get(var).static_expression; - if (expr == 0 || ids[expr].get_type() == TypeUndef) + if (expr == 0 || ir.ids[expr].get_type() == TypeUndef) continue; if (expected == 0) @@ -9711,7 +9708,7 @@ void CompilerGLSL::emit_hoisted_temporaries(vector> &te for (auto &tmp : temporaries) { add_local_variable_name(tmp.second); - auto flags = meta[tmp.second].decoration.decoration_flags; + auto flags = ir.meta[tmp.second].decoration.decoration_flags; auto &type = get(tmp.first); statement(flags_to_precision_qualifiers_glsl(type, flags), variable_decl(type, to_name(tmp.second)), ";"); @@ -9971,7 +9968,7 @@ void CompilerGLSL::emit_block_chain(SPIRBlock &block) { // If we cannot return arrays, we will have a special out argument we can write to instead. // The backend is responsible for setting this up, and redirection the return values as appropriate. - if (ids.at(block.return_value).get_type() != TypeUndef) + if (ir.ids.at(block.return_value).get_type() != TypeUndef) emit_array_copy("SPIRV_Cross_return_value", block.return_value); if (!block_is_outside_flow_control_from_block(get(current_function->entry_block), block) || @@ -9983,7 +9980,7 @@ void CompilerGLSL::emit_block_chain(SPIRBlock &block) else { // OpReturnValue can return Undef, so don't emit anything for this case. - if (ids.at(block.return_value).get_type() != TypeUndef) + if (ir.ids.at(block.return_value).get_type() != TypeUndef) statement("return ", to_expression(block.return_value), ";"); } } diff --git a/spirv_glsl.hpp b/spirv_glsl.hpp index 7fd1938e2..c2d7f6707 100644 --- a/spirv_glsl.hpp +++ b/spirv_glsl.hpp @@ -127,14 +127,26 @@ class CompilerGLSL : public Compiler remap_pls_variables(); } - CompilerGLSL(std::vector spirv_) + explicit CompilerGLSL(std::vector spirv_) : Compiler(move(spirv_)) { init(); } - CompilerGLSL(const uint32_t *ir, size_t word_count) - : Compiler(ir, word_count) + CompilerGLSL(const uint32_t *ir_, size_t word_count) + : Compiler(ir_, word_count) + { + init(); + } + + explicit CompilerGLSL(const ParsedIR &ir_) + : Compiler(ir_) + { + init(); + } + + explicit CompilerGLSL(ParsedIR &&ir_) + : Compiler(std::move(ir_)) { init(); } @@ -586,10 +598,10 @@ class CompilerGLSL : public Compiler private: void init() { - if (source.known) + if (ir.source.known) { - options.es = source.es; - options.version = source.version; + options.es = ir.source.es; + options.version = ir.source.version; } } }; diff --git a/spirv_hlsl.cpp b/spirv_hlsl.cpp index ddad0bc64..712082484 100644 --- a/spirv_hlsl.cpp +++ b/spirv_hlsl.cpp @@ -472,7 +472,7 @@ void CompilerHLSL::emit_interface_block_globally(const SPIRVariable &var) // The global copies of I/O variables should not contain interpolation qualifiers. // These are emitted inside the interface structs. - auto &flags = meta[var.self].decoration.decoration_flags; + auto &flags = ir.meta[var.self].decoration.decoration_flags; auto old_flags = flags; flags.reset(); statement("static ", variable_decl(var), ";"); @@ -806,7 +806,7 @@ void CompilerHLSL::emit_interface_block_in_struct(const SPIRVariable &var, unord bool need_matrix_unroll = var.storage == StorageClassInput && execution.model == ExecutionModelVertex; - auto &m = meta[var.self].decoration; + auto &m = ir.meta[var.self].decoration; auto name = to_name(var.self); if (use_location_number) { @@ -992,7 +992,7 @@ void CompilerHLSL::emit_composite_constants() // global constants directly. bool emitted = false; - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeConstant) { @@ -1020,7 +1020,7 @@ void CompilerHLSL::emit_specialization_constants() SpecializationConstant wg_x, wg_y, wg_z; uint32_t workgroup_size_id = get_work_group_size_specialization_constants(wg_x, wg_y, wg_z); - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeConstant) { @@ -1063,14 +1063,14 @@ void CompilerHLSL::replace_illegal_names() "line", "linear", "matrix", "point", "row_major", "sampler", }; - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { auto &var = id.get(); if (!is_hidden_variable(var)) { - auto &m = meta[var.self].decoration; + auto &m = ir.meta[var.self].decoration; if (keywords.find(m.alias) != end(keywords)) m.alias = join("_", m.alias); } @@ -1090,14 +1090,14 @@ void CompilerHLSL::emit_resources() // Output all basic struct types which are not Block or BufferBlock as these are declared inplace // when such variables are instantiated. - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeType) { auto &type = id.get(); if (type.basetype == SPIRType::Struct && type.array.empty() && !type.pointer && - (!meta[type.self].decoration.decoration_flags.get(DecorationBlock) && - !meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock))) + (!ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock) && + !ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock))) { emit_struct(type); } @@ -1109,7 +1109,7 @@ void CompilerHLSL::emit_resources() bool emitted = false; // Output UBOs and SSBOs - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { @@ -1117,8 +1117,8 @@ void CompilerHLSL::emit_resources() auto &type = get(var.basetype); bool is_block_storage = type.storage == StorageClassStorageBuffer || type.storage == StorageClassUniform; - bool has_block_flags = meta[type.self].decoration.decoration_flags.get(DecorationBlock) || - meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); + bool has_block_flags = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock) || + ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); if (var.storage != StorageClassFunction && type.pointer && is_block_storage && !is_hidden_variable(var) && has_block_flags) @@ -1130,7 +1130,7 @@ void CompilerHLSL::emit_resources() } // Output push constant blocks - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { @@ -1152,7 +1152,7 @@ void CompilerHLSL::emit_resources() } // Output Uniform Constants (values, samplers, images, etc). - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { @@ -1176,13 +1176,13 @@ void CompilerHLSL::emit_resources() // Emit builtin input and output variables here. emit_builtin_variables(); - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { auto &var = id.get(); auto &type = get(var.basetype); - bool block = meta[type.self].decoration.decoration_flags.get(DecorationBlock); + bool block = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock); // Do not emit I/O blocks here. // I/O blocks can be arrayed, so we must deal with them separately to support geometry shaders @@ -1208,13 +1208,13 @@ void CompilerHLSL::emit_resources() unordered_set active_outputs; vector input_variables; vector output_variables; - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { auto &var = id.get(); auto &type = get(var.basetype); - bool block = meta[type.self].decoration.decoration_flags.get(DecorationBlock); + bool block = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock); if (var.storage != StorageClassInput && var.storage != StorageClassOutput) continue; @@ -1801,13 +1801,13 @@ void CompilerHLSL::emit_struct_member(const SPIRType &type, uint32_t member_type auto &membertype = get(member_type_id); Bitset memberflags; - auto &memb = meta[type.self].members; + auto &memb = ir.meta[type.self].members; if (index < memb.size()) memberflags = memb[index].decoration_flags; string qualifiers; - bool is_block = meta[type.self].decoration.decoration_flags.get(DecorationBlock) || - meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); + bool is_block = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock) || + ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); if (is_block) qualifiers = to_interpolation_qualifiers(memberflags); @@ -1838,7 +1838,7 @@ void CompilerHLSL::emit_buffer_block(const SPIRVariable &var) if (is_uav) { - Bitset flags = get_buffer_block_flags(var); + Bitset flags = ir.get_buffer_block_flags(var); bool is_readonly = flags.get(DecorationNonWritable); bool is_coherent = flags.get(DecorationCoherent); add_resource_name(var.self); @@ -1918,7 +1918,7 @@ void CompilerHLSL::emit_push_constant_block(const SPIRVariable &var) flattened_structs.insert(var.self); type.member_name_cache.clear(); add_resource_name(var.self); - auto &memb = meta[type.self].members; + auto &memb = ir.meta[type.self].members; statement("cbuffer SPIRV_CROSS_RootConstant_", to_name(var.self), to_resource_register('b', layout.binding, layout.space)); @@ -1994,7 +1994,7 @@ string CompilerHLSL::to_func_call_arg(uint32_t id) void CompilerHLSL::emit_function_prototype(SPIRFunction &func, const Bitset &return_flags) { - if (func.self != entry_point) + if (func.self != ir.default_entry_point) add_function_overload(func); auto &execution = get_entry_point(); @@ -2016,7 +2016,7 @@ void CompilerHLSL::emit_function_prototype(SPIRFunction &func, const Bitset &ret decl = "void "; } - if (func.self == entry_point) + if (func.self == ir.default_entry_point) { if (execution.model == ExecutionModelVertex) decl += "vert_main"; @@ -2087,13 +2087,13 @@ void CompilerHLSL::emit_hlsl_entry_point() arguments.push_back("SPIRV_Cross_Input stage_input"); // Add I/O blocks as separate arguments with appropriate storage qualifier. - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { auto &var = id.get(); auto &type = get(var.basetype); - bool block = meta[type.self].decoration.decoration_flags.get(DecorationBlock); + bool block = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock); if (var.storage != StorageClassInput && var.storage != StorageClassOutput) continue; @@ -2257,13 +2257,13 @@ void CompilerHLSL::emit_hlsl_entry_point() }); // Copy from stage input struct to globals. - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { auto &var = id.get(); auto &type = get(var.basetype); - bool block = meta[type.self].decoration.decoration_flags.get(DecorationBlock); + bool block = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock); if (var.storage != StorageClassInput) continue; @@ -2307,13 +2307,13 @@ void CompilerHLSL::emit_hlsl_entry_point() SPIRV_CROSS_THROW("Unsupported shader stage."); // Copy block outputs. - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { auto &var = id.get(); auto &type = get(var.basetype); - bool block = meta[type.self].decoration.decoration_flags.get(DecorationBlock); + bool block = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock); if (var.storage != StorageClassOutput) continue; @@ -2361,13 +2361,13 @@ void CompilerHLSL::emit_hlsl_entry_point() } }); - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { auto &var = id.get(); auto &type = get(var.basetype); - bool block = meta[type.self].decoration.decoration_flags.get(DecorationBlock); + bool block = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock); if (var.storage != StorageClassOutput) continue; @@ -2421,13 +2421,10 @@ void CompilerHLSL::emit_fixup() void CompilerHLSL::emit_texture_op(const Instruction &i) { - auto ops = stream(i); + auto *ops = stream(i); auto op = static_cast(i.op); uint32_t length = i.length; - if (i.offset + length > spirv.size()) - SPIRV_CROSS_THROW("Compiler::parse() opcode out of range."); - vector inherited_expressions; uint32_t result_type = ops[0]; @@ -2886,7 +2883,7 @@ string CompilerHLSL::to_resource_binding(const SPIRVariable &var) { if (has_decoration(type.self, DecorationBufferBlock)) { - Bitset flags = get_buffer_block_flags(var); + Bitset flags = ir.get_buffer_block_flags(var); bool is_readonly = flags.get(DecorationNonWritable); space = is_readonly ? 't' : 'u'; // UAV } @@ -4550,7 +4547,7 @@ uint32_t CompilerHLSL::remap_num_workgroups_builtin() return 0; // Create a new, fake UBO. - uint32_t offset = increase_bound_by(4); + uint32_t offset = ir.increase_bound_by(4); uint32_t uint_type_id = offset; uint32_t block_type_id = offset + 1; @@ -4582,7 +4579,7 @@ uint32_t CompilerHLSL::remap_num_workgroups_builtin() ptr_type.self = block_type_id; set(variable_id, block_pointer_type_id, StorageClassUniform); - meta[variable_id].decoration.alias = "SPIRV_Cross_NumWorkgroups"; + ir.meta[variable_id].decoration.alias = "SPIRV_Cross_NumWorkgroups"; num_workgroups_builtin = variable_id; return variable_id; @@ -4635,7 +4632,7 @@ string CompilerHLSL::compile() emit_header(); emit_resources(); - emit_function(get(entry_point), Bitset()); + emit_function(get(ir.default_entry_point), Bitset()); emit_hlsl_entry_point(); pass_count++; diff --git a/spirv_hlsl.hpp b/spirv_hlsl.hpp index bcc82b104..1cab2cb4c 100644 --- a/spirv_hlsl.hpp +++ b/spirv_hlsl.hpp @@ -56,13 +56,23 @@ class CompilerHLSL : public CompilerGLSL bool point_coord_compat = false; }; - CompilerHLSL(std::vector spirv_) + explicit CompilerHLSL(std::vector spirv_) : CompilerGLSL(move(spirv_)) { } - CompilerHLSL(const uint32_t *ir, size_t size) - : CompilerGLSL(ir, size) + CompilerHLSL(const uint32_t *ir_, size_t size) + : CompilerGLSL(ir_, size) + { + } + + explicit CompilerHLSL(const ParsedIR &ir_) + : CompilerGLSL(ir_) + { + } + + explicit CompilerHLSL(ParsedIR &&ir_) + : CompilerGLSL(std::move(ir_)) { } diff --git a/spirv_msl.cpp b/spirv_msl.cpp index 8e2e5a922..608119cd9 100644 --- a/spirv_msl.cpp +++ b/spirv_msl.cpp @@ -43,9 +43,35 @@ CompilerMSL::CompilerMSL(vector spirv_, vector *p_vtx_a resource_bindings.push_back(&rb); } -CompilerMSL::CompilerMSL(const uint32_t *ir, size_t word_count, MSLVertexAttr *p_vtx_attrs, size_t vtx_attrs_count, +CompilerMSL::CompilerMSL(const uint32_t *ir_, size_t word_count, MSLVertexAttr *p_vtx_attrs, size_t vtx_attrs_count, MSLResourceBinding *p_res_bindings, size_t res_bindings_count) - : CompilerGLSL(ir, word_count) + : CompilerGLSL(ir_, word_count) +{ + if (p_vtx_attrs) + for (size_t i = 0; i < vtx_attrs_count; i++) + vtx_attrs_by_location[p_vtx_attrs[i].location] = &p_vtx_attrs[i]; + + if (p_res_bindings) + for (size_t i = 0; i < res_bindings_count; i++) + resource_bindings.push_back(&p_res_bindings[i]); +} + +CompilerMSL::CompilerMSL(const ParsedIR &ir_, MSLVertexAttr *p_vtx_attrs, size_t vtx_attrs_count, + MSLResourceBinding *p_res_bindings, size_t res_bindings_count) + : CompilerGLSL(ir_) +{ + if (p_vtx_attrs) + for (size_t i = 0; i < vtx_attrs_count; i++) + vtx_attrs_by_location[p_vtx_attrs[i].location] = &p_vtx_attrs[i]; + + if (p_res_bindings) + for (size_t i = 0; i < res_bindings_count; i++) + resource_bindings.push_back(&p_res_bindings[i]); +} + +CompilerMSL::CompilerMSL(ParsedIR &&ir_, MSLVertexAttr *p_vtx_attrs, size_t vtx_attrs_count, + MSLResourceBinding *p_res_bindings, size_t res_bindings_count) + : CompilerGLSL(std::move(ir_)) { if (p_vtx_attrs) for (size_t i = 0; i < vtx_attrs_count; i++) @@ -64,23 +90,23 @@ void CompilerMSL::build_implicit_builtins() bool has_frag_coord = false; bool has_sample_id = false; - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() != TypeVariable) continue; auto &var = id.get(); - if (need_subpass_input && var.storage == StorageClassInput && meta[var.self].decoration.builtin && - meta[var.self].decoration.builtin_type == BuiltInFragCoord) + if (need_subpass_input && var.storage == StorageClassInput && ir.meta[var.self].decoration.builtin && + ir.meta[var.self].decoration.builtin_type == BuiltInFragCoord) { builtin_frag_coord_id = var.self; has_frag_coord = true; break; } - if (need_sample_pos && var.storage == StorageClassInput && meta[var.self].decoration.builtin && - meta[var.self].decoration.builtin_type == BuiltInSampleId) + if (need_sample_pos && var.storage == StorageClassInput && ir.meta[var.self].decoration.builtin && + ir.meta[var.self].decoration.builtin_type == BuiltInSampleId) { builtin_sample_id_id = var.self; has_sample_id = true; @@ -90,7 +116,7 @@ void CompilerMSL::build_implicit_builtins() if (!has_frag_coord && need_subpass_input) { - uint32_t offset = increase_bound_by(3); + uint32_t offset = ir.increase_bound_by(3); uint32_t type_id = offset; uint32_t type_ptr_id = offset + 1; uint32_t var_id = offset + 2; @@ -117,7 +143,7 @@ void CompilerMSL::build_implicit_builtins() if (!has_sample_id && need_sample_pos) { - uint32_t offset = increase_bound_by(3); + uint32_t offset = ir.increase_bound_by(3); uint32_t type_id = offset; uint32_t type_ptr_id = offset + 1; uint32_t var_id = offset + 2; @@ -144,7 +170,7 @@ void CompilerMSL::build_implicit_builtins() if (msl_options.swizzle_texture_samples && has_sampled_images) { - uint32_t offset = increase_bound_by(5); + uint32_t offset = ir.increase_bound_by(5); uint32_t type_id = offset; uint32_t type_arr_id = offset + 1; uint32_t struct_id = offset + 2; @@ -439,7 +465,7 @@ string CompilerMSL::compile() emit_specialization_constants(); emit_resources(); emit_custom_functions(); - emit_function(get(entry_point), Bitset()); + emit_function(get(ir.default_entry_point), Bitset()); pass_count++; } while (force_recompile); @@ -477,7 +503,7 @@ string CompilerMSL::compile(MSLConfiguration &msl_cfg, vector *p_ void CompilerMSL::preprocess_op_codes() { OpCodePreprocessor preproc(*this); - traverse_all_reachable_opcodes(get(entry_point), preproc); + traverse_all_reachable_opcodes(get(ir.default_entry_point), preproc); if (preproc.suppress_missing_prototypes) add_pragma_line("#pragma clang diagnostic ignored \"-Wmissing-prototypes\""); @@ -497,7 +523,7 @@ void CompilerMSL::preprocess_op_codes() // Non-constant variables cannot have global scope in Metal. void CompilerMSL::localize_global_variables() { - auto &entry_func = get(entry_point); + auto &entry_func = get(ir.default_entry_point); auto iter = global_variables.begin(); while (iter != global_variables.end()) { @@ -517,7 +543,7 @@ void CompilerMSL::localize_global_variables() // Turn off specialization of any constants that are used for array lengths. void CompilerMSL::resolve_specialized_array_lengths() { - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeConstant) { @@ -534,7 +560,7 @@ void CompilerMSL::extract_global_variables_from_functions() { // Uniforms unordered_set global_var_ids; - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { @@ -549,14 +575,14 @@ void CompilerMSL::extract_global_variables_from_functions() } // Local vars that are declared in the main function and accessed directly by a function - auto &entry_func = get(entry_point); + auto &entry_func = get(ir.default_entry_point); for (auto &var : entry_func.local_variables) if (get(var).storage != StorageClassFunction) global_var_ids.insert(var); std::set added_arg_ids; unordered_set processed_func_ids; - extract_global_variables_from_function(entry_point, added_arg_ids, global_var_ids, processed_func_ids); + extract_global_variables_from_function(ir.default_entry_point, added_arg_ids, global_var_ids, processed_func_ids); } // MSL does not support the use of global variables for shader input content. @@ -655,7 +681,7 @@ void CompilerMSL::extract_global_variables_from_function(uint32_t func_id, std:: function_global_vars[func_id] = added_arg_ids; // Add the global variables as arguments to the function - if (func_id != entry_point) + if (func_id != ir.default_entry_point) { for (uint32_t arg_id : added_arg_ids) { @@ -677,7 +703,7 @@ void CompilerMSL::extract_global_variables_from_function(uint32_t func_id, std:: if (is_builtin && has_active_builtin(builtin, var.storage)) { // Add a arg variable with the same type and decorations as the member - uint32_t next_ids = increase_bound_by(2); + uint32_t next_ids = ir.increase_bound_by(2); uint32_t ptr_type_id = next_ids + 0; uint32_t var_id = next_ids + 1; @@ -690,20 +716,20 @@ void CompilerMSL::extract_global_variables_from_function(uint32_t func_id, std:: func.add_parameter(mbr_type_id, var_id, true); set(var_id, ptr_type_id, StorageClassFunction); - meta[var_id].decoration = meta[type_id].members[mbr_idx]; + ir.meta[var_id].decoration = ir.meta[type_id].members[mbr_idx]; } mbr_idx++; } } else { - uint32_t next_id = increase_bound_by(1); + uint32_t next_id = ir.increase_bound_by(1); func.add_parameter(type_id, next_id, true); set(next_id, type_id, StorageClassFunction, 0, arg_id); // Ensure the existing variable has a valid name and the new variable has all the same meta info set_name(arg_id, ensure_valid_name(to_name(arg_id), "v")); - meta[next_id] = meta[arg_id]; + ir.meta[next_id] = ir.meta[arg_id]; } } } @@ -713,7 +739,7 @@ void CompilerMSL::extract_global_variables_from_function(uint32_t func_id, std:: // that are recursively contained within the type referenced by that variable should be packed tightly. void CompilerMSL::mark_packable_structs() { - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { @@ -778,7 +804,7 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage) // Accumulate the variables that should appear in the interface struct vector vars; bool incl_builtins = (storage == StorageClassOutput); - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { @@ -799,7 +825,7 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage) // Add a new typed variable for this interface structure. // The initializer expression is allocated here, but populated when the function // declaraion is emitted, because it is cleared after each compilation pass. - uint32_t next_id = increase_bound_by(3); + uint32_t next_id = ir.increase_bound_by(3); uint32_t ib_type_id = next_id++; auto &ib_type = set(ib_type_id); ib_type.basetype = SPIRType::Struct; @@ -827,7 +853,7 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage) // Indicate the output var requires early initialization. bool ep_should_return_output = !get_is_rasterization_disabled(); uint32_t rtn_id = ep_should_return_output ? ib_var_id : 0; - auto &entry_func = get(entry_point); + auto &entry_func = get(ir.default_entry_point); entry_func.add_local_variable(ib_var_id); for (auto &blk_id : entry_func.blocks) { @@ -843,10 +869,10 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage) break; } - set_name(ib_type_id, to_name(entry_point) + "_" + ib_var_ref); + set_name(ib_type_id, to_name(ir.default_entry_point) + "_" + ib_var_ref); set_name(ib_var_id, ib_var_ref); - auto &entry_func = get(entry_point); + auto &entry_func = get(ir.default_entry_point); for (auto p_var : vars) { @@ -1083,7 +1109,7 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage) // Update the original variable reference to include the structure reference string qual_var_name = ib_var_ref + "." + mbr_name; - meta[p_var->self].decoration.qualified_alias = qual_var_name; + ir.meta[p_var->self].decoration.qualified_alias = qual_var_name; // Copy the variable location from the original variable to the member if (get_decoration_bitset(p_var->self).get(DecorationLocation)) @@ -1128,7 +1154,7 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage) } // Sort the members of the structure by their locations. - MemberSorter member_sorter(ib_type, meta[ib_type_id], MemberSorter::Location); + MemberSorter member_sorter(ib_type, ir.meta[ib_type_id], MemberSorter::Location); member_sorter.sort(); return ib_var_id; @@ -1144,7 +1170,7 @@ uint32_t CompilerMSL::ensure_correct_builtin_type(uint32_t type_id, BuiltIn buil if ((builtin == BuiltInSampleMask && is_array(type)) || ((builtin == BuiltInLayer || builtin == BuiltInViewportIndex) && type.basetype != SPIRType::UInt)) { - uint32_t next_id = increase_bound_by(type.pointer ? 2 : 1); + uint32_t next_id = ir.increase_bound_by(type.pointer ? 2 : 1); uint32_t base_type_id = next_id++; auto &base_type = set(base_type_id); base_type.basetype = SPIRType::UInt; @@ -1175,7 +1201,7 @@ void CompilerMSL::align_struct(SPIRType &ib_type) // Sort the members of the interface structure by their offset. // They should already be sorted per SPIR-V spec anyway. - MemberSorter member_sorter(ib_type, meta[ib_type_id], MemberSorter::Offset); + MemberSorter member_sorter(ib_type, ir.meta[ib_type_id], MemberSorter::Offset); member_sorter.sort(); uint32_t curr_offset; @@ -1800,7 +1826,7 @@ void CompilerMSL::emit_custom_functions() void CompilerMSL::declare_undefined_values() { bool emitted = false; - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeUndef) { @@ -1821,7 +1847,7 @@ void CompilerMSL::declare_constant_arrays() // global constants directly, so we are able to use constants as variable expressions. bool emitted = false; - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeConstant) { @@ -1848,7 +1874,7 @@ void CompilerMSL::emit_resources() // Output non-builtin interface structs. These include local function structs // and structs nested within uniform and read-write buffers. unordered_set declared_structs; - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeType) { @@ -1896,7 +1922,7 @@ void CompilerMSL::emit_specialization_constants() uint32_t workgroup_size_id = get_work_group_size_specialization_constants(wg_x, wg_y, wg_z); bool emitted = false; - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeConstant) { @@ -2520,12 +2546,12 @@ void CompilerMSL::emit_array_copy(const string &lhs, uint32_t rhs_id) // Unfortunately, we cannot template on address space in MSL, // so explicit address space redirection it is ... bool is_constant = false; - if (ids[rhs_id].get_type() == TypeConstant) + if (ir.ids[rhs_id].get_type() == TypeConstant) { is_constant = true; } else if (var && var->remapped_variable && var->statically_assigned && - ids[var->static_expression].get_type() == TypeConstant) + ir.ids[var->static_expression].get_type() == TypeConstant) { is_constant = true; } @@ -2551,7 +2577,7 @@ bool CompilerMSL::maybe_emit_array_assignment(uint32_t id_lhs, uint32_t id_rhs) if (var && var->remapped_variable && var->statically_assigned) return true; - if (ids[id_rhs].get_type() == TypeConstant && var && var->deferred_declaration) + if (ir.ids[id_rhs].get_type() == TypeConstant && var && var->deferred_declaration) { // Special case, if we end up declaring a variable when assigning the constant array, // we can avoid the copy by directly assigning the constant expression. @@ -2818,13 +2844,13 @@ void CompilerMSL::emit_interface_block(uint32_t ib_var_id) // If this is the entry point function, Metal-specific return value and function arguments are added. void CompilerMSL::emit_function_prototype(SPIRFunction &func, const Bitset &) { - if (func.self != entry_point) + if (func.self != ir.default_entry_point) add_function_overload(func); local_variable_names = resource_names; string decl; - processing_entry_point = (func.self == entry_point); + processing_entry_point = (func.self == ir.default_entry_point); auto &type = get(func.return_type); @@ -2865,7 +2891,7 @@ void CompilerMSL::emit_function_prototype(SPIRFunction &func, const Bitset &) { auto &ed_var = get(var_id); if (!ed_var.initializer) - ed_var.initializer = increase_bound_by(1); + ed_var.initializer = ir.increase_bound_by(1); set(ed_var.initializer, "{}", ed_var.basetype, true); } @@ -3271,7 +3297,7 @@ string CompilerMSL::round_fp_tex_coords(string tex_coords, bool coord_is_fp) // The ID must be a scalar constant. string CompilerMSL::to_component_argument(uint32_t id) { - if (ids[id].get_type() != TypeConstant) + if (ir.ids[id].get_type() != TypeConstant) { SPIRV_CROSS_THROW("ID " + to_string(id) + " is not an OpConstant."); return "component::x"; @@ -3347,7 +3373,7 @@ bool CompilerMSL::is_non_native_row_major_matrix(uint32_t id) return false; // Non-matrix or column-major matrix types do not need to be converted. - if (!meta[id].decoration.decoration_flags.get(DecorationRowMajor)) + if (!ir.meta[id].decoration.decoration_flags.get(DecorationRowMajor)) return false; // Generate a function that will swap matrix elements from row-major to column-major. @@ -3682,7 +3708,7 @@ string CompilerMSL::member_attribute_qualifier(const SPIRType &type, uint32_t in // index as the location. uint32_t CompilerMSL::get_ordered_member_location(uint32_t type_id, uint32_t index, uint32_t *comp) { - auto &m = meta.at(type_id); + auto &m = ir.meta.at(type_id); if (index < m.members.size()) { auto &dec = m.members[index]; @@ -3754,7 +3780,7 @@ string CompilerMSL::get_argument_address_space(const SPIRVariable &argument) case StorageClassStorageBuffer: { - bool readonly = get_buffer_block_flags(argument).get(DecorationNonWritable); + bool readonly = ir.get_buffer_block_flags(argument).get(DecorationNonWritable); return readonly ? "const device" : "device"; } @@ -3766,7 +3792,7 @@ string CompilerMSL::get_argument_address_space(const SPIRVariable &argument) bool ssbo = has_decoration(type.self, DecorationBufferBlock); if (ssbo) { - bool readonly = get_buffer_block_flags(argument).get(DecorationNonWritable); + bool readonly = ir.get_buffer_block_flags(argument).get(DecorationNonWritable); return readonly ? "const device" : "device"; } else @@ -3816,7 +3842,7 @@ string CompilerMSL::entry_point_args(bool append_comma) vector resources; - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { @@ -3865,7 +3891,7 @@ string CompilerMSL::entry_point_args(bool append_comma) { case SPIRType::Struct: { - auto &m = meta.at(type.self); + auto &m = ir.meta.at(type.self); if (m.members.size() == 0) break; if (!type.array.empty()) @@ -3920,14 +3946,14 @@ string CompilerMSL::entry_point_args(bool append_comma) } // Builtin variables - for (auto &id : ids) + for (auto &id : ir.ids) { if (id.get_type() == TypeVariable) { auto &var = id.get(); uint32_t var_id = var.self; - BuiltIn bi_type = meta[var_id].decoration.builtin_type; + BuiltIn bi_type = ir.meta[var_id].decoration.builtin_type; // Don't emit SamplePosition as a separate parameter. In the entry // point, we get that by calling get_sample_position() on the sample ID. @@ -3943,7 +3969,7 @@ string CompilerMSL::entry_point_args(bool append_comma) } else { - auto &entry_func = get(entry_point); + auto &entry_func = get(ir.default_entry_point); entry_func.fixup_hooks_in.push_back([=]() { statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = get_sample_position(", to_expression(builtin_sample_id_id), ");"); @@ -3970,7 +3996,7 @@ string CompilerMSL::entry_point_args(bool append_comma) uint32_t CompilerMSL::get_metal_resource_index(SPIRVariable &var, SPIRType::BaseType basetype) { auto &execution = get_entry_point(); - auto &var_dec = meta[var.self].decoration; + auto &var_dec = ir.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; @@ -4100,9 +4126,9 @@ string CompilerMSL::argument_decl(const SPIRFunction::Parameter &arg) // has a qualified name, use it, otherwise use the standard name. string CompilerMSL::to_name(uint32_t id, bool allow_alias) const { - if (current_function && (current_function->self == entry_point)) + if (current_function && (current_function->self == ir.default_entry_point)) { - string qual_name = meta.at(id).decoration.qualified_alias; + string qual_name = ir.meta.at(id).decoration.qualified_alias; if (!qual_name.empty()) return qual_name; } @@ -4145,13 +4171,13 @@ void CompilerMSL::replace_illegal_names() "saturate", }; - for (auto &id : ids) + for (auto &id : ir.ids) { switch (id.get_type()) { case TypeVariable: { - auto &dec = meta[id.get_id()].decoration; + auto &dec = ir.meta[id.get_id()].decoration; if (keywords.find(dec.alias) != end(keywords)) dec.alias += "0"; @@ -4160,7 +4186,7 @@ void CompilerMSL::replace_illegal_names() case TypeFunction: { - auto &dec = meta[id.get_id()].decoration; + auto &dec = ir.meta[id.get_id()].decoration; if (illegal_func_names.find(dec.alias) != end(illegal_func_names)) dec.alias += "0"; @@ -4169,7 +4195,7 @@ void CompilerMSL::replace_illegal_names() case TypeType: { - for (auto &mbr_dec : meta[id.get_id()].members) + for (auto &mbr_dec : ir.meta[id.get_id()].members) if (keywords.find(mbr_dec.alias) != end(keywords)) mbr_dec.alias += "0"; @@ -4181,7 +4207,7 @@ void CompilerMSL::replace_illegal_names() } } - for (auto &entry : entry_points) + for (auto &entry : ir.entry_points) { // Change both the entry point name and the alias, to keep them synced. string &ep_name = entry.second.name; @@ -4189,7 +4215,7 @@ void CompilerMSL::replace_illegal_names() ep_name += "0"; // Always write this because entry point might have been renamed earlier. - meta[entry.first].decoration.alias = ep_name; + ir.meta[entry.first].decoration.alias = ep_name; } CompilerGLSL::replace_illegal_names(); @@ -4528,7 +4554,7 @@ string CompilerMSL::builtin_to_glsl(BuiltIn builtin, StorageClass storage) case BuiltInLayer: case BuiltInFragDepth: case BuiltInSampleMask: - if (storage != StorageClassInput && current_function && (current_function->self == entry_point)) + if (storage != StorageClassInput && current_function && (current_function->self == ir.default_entry_point)) return stage_out_var_name + "." + CompilerGLSL::builtin_to_glsl(builtin, storage); break; @@ -4787,7 +4813,7 @@ void CompilerMSL::analyze_sampled_image_usage() if (msl_options.swizzle_texture_samples) { SampledImageScanner scanner(*this); - traverse_all_reachable_opcodes(get(entry_point), scanner); + traverse_all_reachable_opcodes(get(ir.default_entry_point), scanner); } } @@ -4935,7 +4961,7 @@ CompilerMSL::SPVFuncImpl CompilerMSL::OpCodePreprocessor::get_spv_func_impl(Op o uint32_t id_rhs = args[1]; const SPIRType *type = nullptr; - if (compiler.ids[id_rhs].get_type() != TypeNone) + if (compiler.ir.ids[id_rhs].get_type() != TypeNone) { // Could be a constant, or similar. type = &compiler.expression_type(id_rhs); @@ -5187,7 +5213,8 @@ std::string CompilerMSL::to_initializer_expression(const SPIRVariable &var) // FIXME: We cannot handle non-constant arrays being initialized. // We will need to inject spvArrayCopy here somehow ... auto &type = get(var.basetype); - if (ids[var.initializer].get_type() == TypeConstant && (!type.array.empty() || type.basetype == SPIRType::Struct)) + if (ir.ids[var.initializer].get_type() == TypeConstant && + (!type.array.empty() || type.basetype == SPIRType::Struct)) return constant_expression(get(var.initializer)); else return CompilerGLSL::to_initializer_expression(var); diff --git a/spirv_msl.hpp b/spirv_msl.hpp index 5784f9560..604ee261f 100644 --- a/spirv_msl.hpp +++ b/spirv_msl.hpp @@ -271,6 +271,13 @@ class CompilerMSL : public CompilerGLSL CompilerMSL(const uint32_t *ir, size_t word_count, MSLVertexAttr *p_vtx_attrs = nullptr, size_t vtx_attrs_count = 0, MSLResourceBinding *p_res_bindings = nullptr, size_t res_bindings_count = 0); + // Alternate constructors taking pre-parsed IR directly. + CompilerMSL(const ParsedIR &ir, MSLVertexAttr *p_vtx_attrs = nullptr, size_t vtx_attrs_count = 0, + MSLResourceBinding *p_res_bindings = nullptr, size_t res_bindings_count = 0); + + CompilerMSL(ParsedIR &&ir, MSLVertexAttr *p_vtx_attrs = nullptr, size_t vtx_attrs_count = 0, + MSLResourceBinding *p_res_bindings = nullptr, size_t res_bindings_count = 0); + // Compiles the SPIR-V code into Metal Shading Language. std::string compile() override; diff --git a/spirv_parser.cpp b/spirv_parser.cpp new file mode 100644 index 000000000..9231f65f6 --- /dev/null +++ b/spirv_parser.cpp @@ -0,0 +1,1025 @@ +/* + * Copyright 2018 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_parser.hpp" +#include + +using namespace std; +using namespace spv; + +namespace spirv_cross +{ +Parser::Parser(std::vector spirv) +{ + ir.spirv = move(spirv); +} + +Parser::Parser(const uint32_t *spirv_data, size_t word_count) +{ + ir.spirv = vector(spirv_data, spirv_data + word_count); +} + +static inline uint32_t swap_endian(uint32_t v) +{ + return ((v >> 24) & 0x000000ffu) | ((v >> 8) & 0x0000ff00u) | ((v << 8) & 0x00ff0000u) | ((v << 24) & 0xff000000u); +} + +static bool is_valid_spirv_version(uint32_t version) +{ + switch (version) + { + // Allow v99 since it tends to just work. + case 99: + case 0x10000: // SPIR-V 1.0 + case 0x10100: // SPIR-V 1.1 + case 0x10200: // SPIR-V 1.2 + case 0x10300: // SPIR-V 1.3 + return true; + + default: + return false; + } +} + +void Parser::parse() +{ + auto &spirv = ir.spirv; + + auto len = spirv.size(); + if (len < 5) + SPIRV_CROSS_THROW("SPIRV file too small."); + + auto s = spirv.data(); + + // Endian-swap if we need to. + if (s[0] == swap_endian(MagicNumber)) + transform(begin(spirv), end(spirv), begin(spirv), [](uint32_t c) { return swap_endian(c); }); + + if (s[0] != MagicNumber || !is_valid_spirv_version(s[1])) + SPIRV_CROSS_THROW("Invalid SPIRV format."); + + uint32_t bound = s[3]; + ir.set_id_bounds(bound); + + uint32_t offset = 5; + + vector instructions; + while (offset < len) + { + Instruction instr = {}; + instr.op = spirv[offset] & 0xffff; + instr.count = (spirv[offset] >> 16) & 0xffff; + + if (instr.count == 0) + SPIRV_CROSS_THROW("SPIR-V instructions cannot consume 0 words. Invalid SPIR-V file."); + + instr.offset = offset + 1; + instr.length = instr.count - 1; + + offset += instr.count; + + if (offset > spirv.size()) + SPIRV_CROSS_THROW("SPIR-V instruction goes out of bounds."); + + instructions.push_back(instr); + } + + for (auto &i : instructions) + parse(i); + + if (current_function) + SPIRV_CROSS_THROW("Function was not terminated."); + if (current_block) + SPIRV_CROSS_THROW("Block was not terminated."); +} + +const uint32_t *Parser::stream(const Instruction &instr) const +{ + // If we're not going to use any arguments, just return nullptr. + // We want to avoid case where we return an out of range pointer + // that trips debug assertions on some platforms. + if (!instr.length) + return nullptr; + + if (instr.offset + instr.length > ir.spirv.size()) + SPIRV_CROSS_THROW("Compiler::stream() out of range."); + return &ir.spirv[instr.offset]; +} + +static string extract_string(const vector &spirv, uint32_t offset) +{ + string ret; + for (uint32_t i = offset; i < spirv.size(); i++) + { + uint32_t w = spirv[i]; + + for (uint32_t j = 0; j < 4; j++, w >>= 8) + { + char c = w & 0xff; + if (c == '\0') + return ret; + ret += c; + } + } + + SPIRV_CROSS_THROW("String was not terminated before EOF"); +} + +void Parser::parse(const Instruction &instruction) +{ + auto *ops = stream(instruction); + auto op = static_cast(instruction.op); + uint32_t length = instruction.length; + + switch (op) + { + case OpMemoryModel: + case OpSourceExtension: + case OpNop: + case OpLine: + case OpNoLine: + case OpString: + break; + + case OpSource: + { + auto lang = static_cast(ops[0]); + switch (lang) + { + case SourceLanguageESSL: + ir.source.es = true; + ir.source.version = ops[1]; + ir.source.known = true; + ir.source.hlsl = false; + break; + + case SourceLanguageGLSL: + ir.source.es = false; + ir.source.version = ops[1]; + ir.source.known = true; + ir.source.hlsl = false; + break; + + case SourceLanguageHLSL: + // For purposes of cross-compiling, this is GLSL 450. + ir.source.es = false; + ir.source.version = 450; + ir.source.known = true; + ir.source.hlsl = true; + break; + + default: + ir.source.known = false; + break; + } + break; + } + + case OpUndef: + { + uint32_t result_type = ops[0]; + uint32_t id = ops[1]; + set(id, result_type); + break; + } + + case OpCapability: + { + uint32_t cap = ops[0]; + if (cap == CapabilityKernel) + SPIRV_CROSS_THROW("Kernel capability not supported."); + + ir.declared_capabilities.push_back(static_cast(ops[0])); + break; + } + + case OpExtension: + { + auto ext = extract_string(ir.spirv, instruction.offset); + ir.declared_extensions.push_back(move(ext)); + break; + } + + case OpExtInstImport: + { + uint32_t id = ops[0]; + auto ext = extract_string(ir.spirv, instruction.offset + 1); + if (ext == "GLSL.std.450") + set(id, SPIRExtension::GLSL); + else if (ext == "SPV_AMD_shader_ballot") + set(id, SPIRExtension::SPV_AMD_shader_ballot); + else if (ext == "SPV_AMD_shader_explicit_vertex_parameter") + set(id, SPIRExtension::SPV_AMD_shader_explicit_vertex_parameter); + else if (ext == "SPV_AMD_shader_trinary_minmax") + set(id, SPIRExtension::SPV_AMD_shader_trinary_minmax); + else if (ext == "SPV_AMD_gcn_shader") + set(id, SPIRExtension::SPV_AMD_gcn_shader); + else + set(id, SPIRExtension::Unsupported); + + // Other SPIR-V extensions which have ExtInstrs are currently not supported. + + break; + } + + case OpEntryPoint: + { + auto itr = + ir.entry_points.insert(make_pair(ops[1], SPIREntryPoint(ops[1], static_cast(ops[0]), + extract_string(ir.spirv, instruction.offset + 2)))); + auto &e = itr.first->second; + + // Strings need nul-terminator and consume the whole word. + uint32_t strlen_words = uint32_t((e.name.size() + 1 + 3) >> 2); + e.interface_variables.insert(end(e.interface_variables), ops + strlen_words + 2, ops + instruction.length); + + // Set the name of the entry point in case OpName is not provided later. + ir.set_name(ops[1], e.name); + + // If we don't have an entry, make the first one our "default". + if (!ir.default_entry_point) + ir.default_entry_point = ops[1]; + break; + } + + case OpExecutionMode: + { + auto &execution = ir.entry_points[ops[0]]; + auto mode = static_cast(ops[1]); + execution.flags.set(mode); + + switch (mode) + { + case ExecutionModeInvocations: + execution.invocations = ops[2]; + break; + + case ExecutionModeLocalSize: + execution.workgroup_size.x = ops[2]; + execution.workgroup_size.y = ops[3]; + execution.workgroup_size.z = ops[4]; + break; + + case ExecutionModeOutputVertices: + execution.output_vertices = ops[2]; + break; + + default: + break; + } + break; + } + + case OpName: + { + uint32_t id = ops[0]; + ir.set_name(id, extract_string(ir.spirv, instruction.offset + 1)); + break; + } + + case OpMemberName: + { + uint32_t id = ops[0]; + uint32_t member = ops[1]; + ir.set_member_name(id, member, extract_string(ir.spirv, instruction.offset + 2)); + break; + } + + case OpDecorate: + case OpDecorateId: + { + uint32_t id = ops[0]; + + auto decoration = static_cast(ops[1]); + if (length >= 3) + { + ir.meta[id].decoration_word_offset[decoration] = uint32_t(&ops[2] - ir.spirv.data()); + ir.set_decoration(id, decoration, ops[2]); + } + else + ir.set_decoration(id, decoration); + + break; + } + + case OpDecorateStringGOOGLE: + { + uint32_t id = ops[0]; + auto decoration = static_cast(ops[1]); + ir.set_decoration_string(id, decoration, extract_string(ir.spirv, instruction.offset + 2)); + break; + } + + case OpMemberDecorate: + { + uint32_t id = ops[0]; + uint32_t member = ops[1]; + auto decoration = static_cast(ops[2]); + if (length >= 4) + ir.set_member_decoration(id, member, decoration, ops[3]); + else + ir.set_member_decoration(id, member, decoration); + break; + } + + case OpMemberDecorateStringGOOGLE: + { + uint32_t id = ops[0]; + uint32_t member = ops[1]; + auto decoration = static_cast(ops[2]); + ir.set_member_decoration_string(id, member, decoration, extract_string(ir.spirv, instruction.offset + 3)); + break; + } + + // Build up basic types. + case OpTypeVoid: + { + uint32_t id = ops[0]; + auto &type = set(id); + type.basetype = SPIRType::Void; + break; + } + + case OpTypeBool: + { + uint32_t id = ops[0]; + auto &type = set(id); + type.basetype = SPIRType::Boolean; + type.width = 1; + break; + } + + case OpTypeFloat: + { + uint32_t id = ops[0]; + uint32_t width = ops[1]; + auto &type = set(id); + if (width == 64) + type.basetype = SPIRType::Double; + else if (width == 32) + type.basetype = SPIRType::Float; + else if (width == 16) + type.basetype = SPIRType::Half; + else + SPIRV_CROSS_THROW("Unrecognized bit-width of floating point type."); + type.width = width; + break; + } + + case OpTypeInt: + { + uint32_t id = ops[0]; + uint32_t width = ops[1]; + auto &type = set(id); + type.basetype = + ops[2] ? (width > 32 ? SPIRType::Int64 : SPIRType::Int) : (width > 32 ? SPIRType::UInt64 : SPIRType::UInt); + type.width = width; + break; + } + + // Build composite types by "inheriting". + // NOTE: The self member is also copied! For pointers and array modifiers this is a good thing + // since we can refer to decorations on pointee classes which is needed for UBO/SSBO, I/O blocks in geometry/tess etc. + case OpTypeVector: + { + uint32_t id = ops[0]; + uint32_t vecsize = ops[2]; + + auto &base = get(ops[1]); + auto &vecbase = set(id); + + vecbase = base; + vecbase.vecsize = vecsize; + vecbase.self = id; + vecbase.parent_type = ops[1]; + break; + } + + case OpTypeMatrix: + { + uint32_t id = ops[0]; + uint32_t colcount = ops[2]; + + auto &base = get(ops[1]); + auto &matrixbase = set(id); + + matrixbase = base; + matrixbase.columns = colcount; + matrixbase.self = id; + matrixbase.parent_type = ops[1]; + break; + } + + case OpTypeArray: + { + uint32_t id = ops[0]; + auto &arraybase = set(id); + + uint32_t tid = ops[1]; + auto &base = get(tid); + + arraybase = base; + arraybase.parent_type = tid; + + uint32_t cid = ops[2]; + ir.mark_used_as_array_length(cid); + auto *c = maybe_get(cid); + bool literal = c && !c->specialization; + + arraybase.array_size_literal.push_back(literal); + arraybase.array.push_back(literal ? c->scalar() : cid); + // Do NOT set arraybase.self! + break; + } + + case OpTypeRuntimeArray: + { + uint32_t id = ops[0]; + + auto &base = get(ops[1]); + auto &arraybase = set(id); + + arraybase = base; + arraybase.array.push_back(0); + arraybase.array_size_literal.push_back(true); + arraybase.parent_type = ops[1]; + // Do NOT set arraybase.self! + break; + } + + case OpTypeImage: + { + uint32_t id = ops[0]; + auto &type = set(id); + type.basetype = SPIRType::Image; + type.image.type = ops[1]; + type.image.dim = static_cast(ops[2]); + type.image.depth = ops[3] == 1; + type.image.arrayed = ops[4] != 0; + type.image.ms = ops[5] != 0; + type.image.sampled = ops[6]; + type.image.format = static_cast(ops[7]); + type.image.access = (length >= 9) ? static_cast(ops[8]) : AccessQualifierMax; + + if (type.image.sampled == 0) + SPIRV_CROSS_THROW("OpTypeImage Sampled parameter must not be zero."); + + break; + } + + case OpTypeSampledImage: + { + uint32_t id = ops[0]; + uint32_t imagetype = ops[1]; + auto &type = set(id); + type = get(imagetype); + type.basetype = SPIRType::SampledImage; + type.self = id; + break; + } + + case OpTypeSampler: + { + uint32_t id = ops[0]; + auto &type = set(id); + type.basetype = SPIRType::Sampler; + break; + } + + case OpTypePointer: + { + uint32_t id = ops[0]; + + auto &base = get(ops[2]); + auto &ptrbase = set(id); + + ptrbase = base; + if (ptrbase.pointer) + SPIRV_CROSS_THROW("Cannot make pointer-to-pointer type."); + ptrbase.pointer = true; + ptrbase.storage = static_cast(ops[1]); + + if (ptrbase.storage == StorageClassAtomicCounter) + ptrbase.basetype = SPIRType::AtomicCounter; + + ptrbase.parent_type = ops[2]; + + // Do NOT set ptrbase.self! + break; + } + + case OpTypeStruct: + { + uint32_t id = ops[0]; + auto &type = set(id); + type.basetype = SPIRType::Struct; + for (uint32_t i = 1; i < length; i++) + type.member_types.push_back(ops[i]); + + // Check if we have seen this struct type before, with just different + // decorations. + // + // Add workaround for issue #17 as well by looking at OpName for the struct + // types, which we shouldn't normally do. + // We should not normally have to consider type aliases like this to begin with + // however ... glslang issues #304, #307 cover this. + + // For stripped names, never consider struct type aliasing. + // We risk declaring the same struct multiple times, but type-punning is not allowed + // so this is safe. + bool consider_aliasing = !ir.get_name(type.self).empty(); + if (consider_aliasing) + { + for (auto &other : global_struct_cache) + { + if (ir.get_name(type.self) == ir.get_name(other) && + types_are_logically_equivalent(type, get(other))) + { + type.type_alias = other; + break; + } + } + + if (type.type_alias == 0) + global_struct_cache.push_back(id); + } + break; + } + + case OpTypeFunction: + { + uint32_t id = ops[0]; + uint32_t ret = ops[1]; + + auto &func = set(id, ret); + for (uint32_t i = 2; i < length; i++) + func.parameter_types.push_back(ops[i]); + break; + } + + // Variable declaration + // All variables are essentially pointers with a storage qualifier. + case OpVariable: + { + uint32_t type = ops[0]; + uint32_t id = ops[1]; + auto storage = static_cast(ops[2]); + uint32_t initializer = length == 4 ? ops[3] : 0; + + if (storage == StorageClassFunction) + { + if (!current_function) + SPIRV_CROSS_THROW("No function currently in scope"); + current_function->add_local_variable(id); + } + + set(id, type, storage, initializer); + + // hlsl based shaders don't have those decorations. force them and then reset when reading/writing images + auto &ttype = get(type); + if (ttype.basetype == SPIRType::BaseType::Image) + { + ir.set_decoration(id, DecorationNonWritable); + ir.set_decoration(id, DecorationNonReadable); + } + + break; + } + + // OpPhi + // OpPhi is a fairly magical opcode. + // It selects temporary variables based on which parent block we *came from*. + // In high-level languages we can "de-SSA" by creating a function local, and flush out temporaries to this function-local + // variable to emulate SSA Phi. + case OpPhi: + { + if (!current_function) + SPIRV_CROSS_THROW("No function currently in scope"); + if (!current_block) + SPIRV_CROSS_THROW("No block currently in scope"); + + uint32_t result_type = ops[0]; + uint32_t id = ops[1]; + + // Instead of a temporary, create a new function-wide temporary with this ID instead. + auto &var = set(id, result_type, spv::StorageClassFunction); + var.phi_variable = true; + + current_function->add_local_variable(id); + + for (uint32_t i = 2; i + 2 <= length; i += 2) + current_block->phi_variables.push_back({ ops[i], ops[i + 1], id }); + break; + } + + // Constants + case OpSpecConstant: + case OpConstant: + { + uint32_t id = ops[1]; + auto &type = get(ops[0]); + + if (type.width > 32) + set(id, ops[0], ops[2] | (uint64_t(ops[3]) << 32), op == OpSpecConstant); + else + set(id, ops[0], ops[2], op == OpSpecConstant); + break; + } + + case OpSpecConstantFalse: + case OpConstantFalse: + { + uint32_t id = ops[1]; + set(id, ops[0], uint32_t(0), op == OpSpecConstantFalse); + break; + } + + case OpSpecConstantTrue: + case OpConstantTrue: + { + uint32_t id = ops[1]; + set(id, ops[0], uint32_t(1), op == OpSpecConstantTrue); + break; + } + + case OpConstantNull: + { + uint32_t id = ops[1]; + uint32_t type = ops[0]; + make_constant_null(id, type); + break; + } + + case OpSpecConstantComposite: + case OpConstantComposite: + { + uint32_t id = ops[1]; + uint32_t type = ops[0]; + + auto &ctype = get(type); + + // We can have constants which are structs and arrays. + // In this case, our SPIRConstant will be a list of other SPIRConstant ids which we + // can refer to. + if (ctype.basetype == SPIRType::Struct || !ctype.array.empty()) + { + set(id, type, ops + 2, length - 2, op == OpSpecConstantComposite); + } + else + { + uint32_t elements = length - 2; + if (elements > 4) + SPIRV_CROSS_THROW("OpConstantComposite only supports 1, 2, 3 and 4 elements."); + + SPIRConstant remapped_constant_ops[4]; + const SPIRConstant *c[4]; + for (uint32_t i = 0; i < elements; i++) + { + // Specialization constants operations can also be part of this. + // We do not know their value, so any attempt to query SPIRConstant later + // will fail. We can only propagate the ID of the expression and use to_expression on it. + auto *constant_op = maybe_get(ops[2 + i]); + if (constant_op) + { + if (op == OpConstantComposite) + SPIRV_CROSS_THROW("Specialization constant operation used in OpConstantComposite."); + + remapped_constant_ops[i].make_null(get(constant_op->basetype)); + remapped_constant_ops[i].self = constant_op->self; + remapped_constant_ops[i].constant_type = constant_op->basetype; + remapped_constant_ops[i].specialization = true; + c[i] = &remapped_constant_ops[i]; + } + else + c[i] = &get(ops[2 + i]); + } + set(id, type, c, elements, op == OpSpecConstantComposite); + } + break; + } + + // Functions + case OpFunction: + { + uint32_t res = ops[0]; + uint32_t id = ops[1]; + // Control + uint32_t type = ops[3]; + + if (current_function) + SPIRV_CROSS_THROW("Must end a function before starting a new one!"); + + current_function = &set(id, res, type); + break; + } + + case OpFunctionParameter: + { + uint32_t type = ops[0]; + uint32_t id = ops[1]; + + if (!current_function) + SPIRV_CROSS_THROW("Must be in a function!"); + + current_function->add_parameter(type, id); + set(id, type, StorageClassFunction); + break; + } + + case OpFunctionEnd: + { + if (current_block) + { + // Very specific error message, but seems to come up quite often. + SPIRV_CROSS_THROW( + "Cannot end a function before ending the current block.\n" + "Likely cause: If this SPIR-V was created from glslang HLSL, make sure the entry point is valid."); + } + current_function = nullptr; + break; + } + + // Blocks + case OpLabel: + { + // OpLabel always starts a block. + if (!current_function) + SPIRV_CROSS_THROW("Blocks cannot exist outside functions!"); + + uint32_t id = ops[0]; + + current_function->blocks.push_back(id); + if (!current_function->entry_block) + current_function->entry_block = id; + + if (current_block) + SPIRV_CROSS_THROW("Cannot start a block before ending the current block."); + + current_block = &set(id); + break; + } + + // Branch instructions end blocks. + case OpBranch: + { + if (!current_block) + SPIRV_CROSS_THROW("Trying to end a non-existing block."); + + uint32_t target = ops[0]; + current_block->terminator = SPIRBlock::Direct; + current_block->next_block = target; + current_block = nullptr; + break; + } + + case OpBranchConditional: + { + if (!current_block) + SPIRV_CROSS_THROW("Trying to end a non-existing block."); + + current_block->condition = ops[0]; + current_block->true_block = ops[1]; + current_block->false_block = ops[2]; + + current_block->terminator = SPIRBlock::Select; + current_block = nullptr; + break; + } + + case OpSwitch: + { + if (!current_block) + SPIRV_CROSS_THROW("Trying to end a non-existing block."); + + if (current_block->merge == SPIRBlock::MergeNone) + SPIRV_CROSS_THROW("Switch statement is not structured"); + + current_block->terminator = SPIRBlock::MultiSelect; + + current_block->condition = ops[0]; + current_block->default_block = ops[1]; + + for (uint32_t i = 2; i + 2 <= length; i += 2) + current_block->cases.push_back({ ops[i], ops[i + 1] }); + + // If we jump to next block, make it break instead since we're inside a switch case block at that point. + ir.block_meta[current_block->next_block] |= ParsedIR::BLOCK_META_MULTISELECT_MERGE_BIT; + + current_block = nullptr; + break; + } + + case OpKill: + { + if (!current_block) + SPIRV_CROSS_THROW("Trying to end a non-existing block."); + current_block->terminator = SPIRBlock::Kill; + current_block = nullptr; + break; + } + + case OpReturn: + { + if (!current_block) + SPIRV_CROSS_THROW("Trying to end a non-existing block."); + current_block->terminator = SPIRBlock::Return; + current_block = nullptr; + break; + } + + case OpReturnValue: + { + if (!current_block) + SPIRV_CROSS_THROW("Trying to end a non-existing block."); + current_block->terminator = SPIRBlock::Return; + current_block->return_value = ops[0]; + current_block = nullptr; + break; + } + + case OpUnreachable: + { + if (!current_block) + SPIRV_CROSS_THROW("Trying to end a non-existing block."); + current_block->terminator = SPIRBlock::Unreachable; + current_block = nullptr; + break; + } + + case OpSelectionMerge: + { + if (!current_block) + SPIRV_CROSS_THROW("Trying to modify a non-existing block."); + + current_block->next_block = ops[0]; + current_block->merge = SPIRBlock::MergeSelection; + ir.block_meta[current_block->next_block] |= ParsedIR::BLOCK_META_SELECTION_MERGE_BIT; + + if (length >= 2) + { + if (ops[1] & SelectionControlFlattenMask) + current_block->hint = SPIRBlock::HintFlatten; + else if (ops[1] & SelectionControlDontFlattenMask) + current_block->hint = SPIRBlock::HintDontFlatten; + } + break; + } + + case OpLoopMerge: + { + if (!current_block) + SPIRV_CROSS_THROW("Trying to modify a non-existing block."); + + current_block->merge_block = ops[0]; + current_block->continue_block = ops[1]; + current_block->merge = SPIRBlock::MergeLoop; + + ir.block_meta[current_block->self] |= ParsedIR::BLOCK_META_LOOP_HEADER_BIT; + ir.block_meta[current_block->merge_block] |= ParsedIR::BLOCK_META_LOOP_MERGE_BIT; + + ir.continue_block_to_loop_header[current_block->continue_block] = current_block->self; + + // Don't add loop headers to continue blocks, + // which would make it impossible branch into the loop header since + // they are treated as continues. + if (current_block->continue_block != current_block->self) + ir.block_meta[current_block->continue_block] |= ParsedIR::BLOCK_META_CONTINUE_BIT; + + if (length >= 3) + { + if (ops[2] & LoopControlUnrollMask) + current_block->hint = SPIRBlock::HintUnroll; + else if (ops[2] & LoopControlDontUnrollMask) + current_block->hint = SPIRBlock::HintDontUnroll; + } + break; + } + + case OpSpecConstantOp: + { + if (length < 3) + SPIRV_CROSS_THROW("OpSpecConstantOp not enough arguments."); + + uint32_t result_type = ops[0]; + uint32_t id = ops[1]; + auto spec_op = static_cast(ops[2]); + + set(id, result_type, spec_op, ops + 3, length - 3); + break; + } + + // Actual opcodes. + default: + { + if (!current_block) + SPIRV_CROSS_THROW("Currently no block to insert opcode."); + + current_block->ops.push_back(instruction); + break; + } + } +} + +bool Parser::types_are_logically_equivalent(const SPIRType &a, const SPIRType &b) const +{ + if (a.basetype != b.basetype) + return false; + if (a.width != b.width) + return false; + if (a.vecsize != b.vecsize) + return false; + if (a.columns != b.columns) + return false; + if (a.array.size() != b.array.size()) + return false; + + size_t array_count = a.array.size(); + if (array_count && memcmp(a.array.data(), b.array.data(), array_count * sizeof(uint32_t)) != 0) + return false; + + if (a.basetype == SPIRType::Image || a.basetype == SPIRType::SampledImage) + { + if (memcmp(&a.image, &b.image, sizeof(SPIRType::Image)) != 0) + return false; + } + + if (a.member_types.size() != b.member_types.size()) + return false; + + size_t member_types = a.member_types.size(); + for (size_t i = 0; i < member_types; i++) + { + if (!types_are_logically_equivalent(get(a.member_types[i]), get(b.member_types[i]))) + return false; + } + + return true; +} + +bool Parser::variable_storage_is_aliased(const SPIRVariable &v) const +{ + auto &type = get(v.basetype); + bool ssbo = v.storage == StorageClassStorageBuffer || + ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); + bool image = type.basetype == SPIRType::Image; + bool counter = type.basetype == SPIRType::AtomicCounter; + + bool is_restrict; + if (ssbo) + is_restrict = ir.get_buffer_block_flags(v).get(DecorationRestrict); + else + is_restrict = ir.has_decoration(v.self, DecorationRestrict); + + return !is_restrict && (ssbo || image || counter); +} + +void Parser::make_constant_null(uint32_t id, uint32_t type) +{ + auto &constant_type = get(type); + + if (!constant_type.array.empty()) + { + assert(constant_type.parent_type); + uint32_t parent_id = ir.increase_bound_by(1); + make_constant_null(parent_id, constant_type.parent_type); + + if (!constant_type.array_size_literal.back()) + SPIRV_CROSS_THROW("Array size of OpConstantNull must be a literal."); + + vector elements(constant_type.array.back()); + for (uint32_t i = 0; i < constant_type.array.back(); i++) + elements[i] = parent_id; + set(id, type, elements.data(), uint32_t(elements.size()), false); + } + else if (!constant_type.member_types.empty()) + { + uint32_t member_ids = ir.increase_bound_by(uint32_t(constant_type.member_types.size())); + vector elements(constant_type.member_types.size()); + for (uint32_t i = 0; i < constant_type.member_types.size(); i++) + { + make_constant_null(member_ids + i, constant_type.member_types[i]); + elements[i] = member_ids + i; + } + set(id, type, elements.data(), uint32_t(elements.size()), false); + } + else + { + auto &constant = set(id, type); + constant.make_null(constant_type); + } +} + +} // namespace spirv_cross \ No newline at end of file diff --git a/spirv_parser.hpp b/spirv_parser.hpp new file mode 100644 index 000000000..6d29f50b7 --- /dev/null +++ b/spirv_parser.hpp @@ -0,0 +1,94 @@ +/* + * Copyright 2018 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. + */ + +#ifndef SPIRV_CROSS_PARSER_HPP +#define SPIRV_CROSS_PARSER_HPP + +#include "spirv_cross_parsed_ir.hpp" +#include +#include + +namespace spirv_cross +{ +class Parser +{ +public: + Parser(const uint32_t *spirv_data, size_t word_count); + Parser(std::vector spirv); + + void parse(); + + ParsedIR &get_parsed_ir() + { + return ir; + } + +private: + ParsedIR ir; + SPIRFunction *current_function = nullptr; + SPIRBlock *current_block = nullptr; + + void parse(const Instruction &instr); + const uint32_t *stream(const Instruction &instr) const; + + template + T &set(uint32_t id, P &&... args) + { + auto &var = variant_set(ir.ids.at(id), std::forward

(args)...); + var.self = id; + return var; + } + + template + T &get(uint32_t id) + { + return variant_get(ir.ids.at(id)); + } + + template + T *maybe_get(uint32_t id) + { + if (ir.ids.at(id).get_type() == T::type) + return &get(id); + else + return nullptr; + } + + template + const T &get(uint32_t id) const + { + return variant_get(ir.ids.at(id)); + } + + template + const T *maybe_get(uint32_t id) const + { + if (ir.ids.at(id).get_type() == T::type) + return &get(id); + else + return nullptr; + } + + // This must be an ordered data structure so we always pick the same type aliases. + std::vector global_struct_cache; + + bool types_are_logically_equivalent(const SPIRType &a, const SPIRType &b) const; + bool variable_storage_is_aliased(const SPIRVariable &v) const; + void make_constant_null(uint32_t id, uint32_t type); +}; +} // namespace spirv_cross + +#endif diff --git a/spirv_reflect.cpp b/spirv_reflect.cpp index c322d972b..c555e25de 100644 --- a/spirv_reflect.cpp +++ b/spirv_reflect.cpp @@ -264,7 +264,7 @@ string CompilerReflection::compile() void CompilerReflection::emit_types() { bool emitted_open_tag = false; - for (auto &id : ids) + for (auto &id : ir.ids) { auto idType = id.get_type(); if (idType == TypeType) @@ -360,7 +360,7 @@ void CompilerReflection::emit_type_member_qualifiers(const SPIRType &type, uint3 auto &membertype = get(type.member_types[index]); emit_type_array(membertype); - auto &memb = meta[type.self].members; + auto &memb = ir.meta[type.self].members; if (index < memb.size()) { auto &dec = memb[index]; @@ -437,7 +437,7 @@ void CompilerReflection::emit_resources(const char *tag, const vector for (auto &res : resources) { auto &type = get_type(res.type_id); - auto typeflags = meta[type.self].decoration.decoration_flags; + auto typeflags = ir.meta[type.self].decoration.decoration_flags; auto &mask = get_decoration_bitset(res.id); // If we don't have a name, use the fallback for the type instead of the variable @@ -565,7 +565,7 @@ void CompilerReflection::emit_specialization_constants() string CompilerReflection::to_member_name(const SPIRType &type, uint32_t index) const { - auto &memb = meta[type.self].members; + auto &memb = ir.meta[type.self].members; if (index < memb.size() && !memb[index].alias.empty()) return memb[index].alias; else diff --git a/spirv_reflect.hpp b/spirv_reflect.hpp index e402fa8fb..fbe7d1f42 100644 --- a/spirv_reflect.hpp +++ b/spirv_reflect.hpp @@ -33,14 +33,26 @@ class CompilerReflection : public CompilerGLSL using Parent = CompilerGLSL; public: - CompilerReflection(std::vector spirv_) + explicit CompilerReflection(std::vector spirv_) : Parent(move(spirv_)) { options.vulkan_semantics = true; } - CompilerReflection(const uint32_t *ir, size_t word_count) - : Parent(ir, word_count) + CompilerReflection(const uint32_t *ir_, size_t word_count) + : Parent(ir_, word_count) + { + options.vulkan_semantics = true; + } + + explicit CompilerReflection(const ParsedIR &ir_) + : CompilerGLSL(ir_) + { + options.vulkan_semantics = true; + } + + explicit CompilerReflection(ParsedIR &&ir_) + : CompilerGLSL(std::move(ir_)) { options.vulkan_semantics = true; }