diff --git a/3rdparty/spirv-cross/main.cpp b/3rdparty/spirv-cross/main.cpp index 6cb539cf8..9124ecaf9 100644 --- a/3rdparty/spirv-cross/main.cpp +++ b/3rdparty/spirv-cross/main.cpp @@ -77,7 +77,7 @@ struct CLICallbacks struct CLIParser { CLIParser(CLICallbacks cbs_, int argc_, char *argv_[]) - : cbs(move(cbs_)) + : cbs(std::move(cbs_)) , argc(argc_) , argv(argv_) { @@ -669,6 +669,8 @@ struct CLIArguments bool emit_line_directives = false; bool enable_storage_image_qualifier_deduction = true; bool force_zero_initialized_variables = false; + bool relax_nan_checks = false; + uint32_t force_recompile_max_debug_iterations = 3; SmallVector msl_discrete_descriptor_sets; SmallVector msl_device_argument_buffers; SmallVector> msl_dynamic_buffers; @@ -790,6 +792,9 @@ static void print_help_hlsl() // clang-format off fprintf(stderr, "\nHLSL options:\n" "\t[--shader-model]:\n\t\tEnables a specific shader model, e.g. --shader-model 50 for SM 5.0.\n" + "\t[--flatten-ubo]:\n\t\tEmit UBOs as plain uniform arrays.\n" + "\t\tE.g.: uniform MyUBO { vec4 a; float b, c, d, e; }; will be emitted as uniform float4 MyUBO[2];\n" + "\t\tCaveat: You cannot mix and match floating-point and integer in the same UBO with this option.\n" "\t[--hlsl-enable-compat]:\n\t\tAllow point size and point coord to be used, even if they won't work as expected.\n" "\t\tPointSize is ignored, and PointCoord returns (0.5, 0.5).\n" "\t[--hlsl-support-nonzero-basevertex-baseinstance]:\n\t\tSupport base vertex and base instance by emitting a special cbuffer declared as:\n" @@ -915,6 +920,7 @@ static void print_help_common() "\t[--mask-stage-output-builtin ]:\n" "\t\tIf a stage output variable with matching builtin is active, " "optimize away the variable if it can affect cross-stage linking correctness.\n" + "\t[--relax-nan-checks]:\n\t\tRelax NaN checks for N{Clamp,Min,Max} and ordered vs. unordered compare instructions.\n" ); // clang-format on } @@ -932,6 +938,8 @@ static void print_help_obscure() "\t\tdo not attempt to analyze usage, and always emit read/write state.\n" "\t[--flatten-multidimensional-arrays]:\n\t\tDo not support multi-dimensional arrays and flatten them to one dimension.\n" "\t[--cpp-interface-name ]:\n\t\tEmit a specific class name in C++ codegen.\n" + "\t[--force-recompile-max-debug-iterations ]:\n\t\tAllow compilation loop to run for N loops.\n" + "\t\tCan be used to triage workarounds, but should not be used as a crutch, since it masks an implementation bug.\n" ); // clang-format on } @@ -1083,7 +1091,7 @@ static HLSLBindingFlags hlsl_resource_type_to_flag(const std::string &arg) static string compile_iteration(const CLIArguments &args, std::vector spirv_file) { - Parser spirv_parser(move(spirv_file)); + Parser spirv_parser(std::move(spirv_file)); spirv_parser.parse(); unique_ptr compiler; @@ -1092,13 +1100,13 @@ static string compile_iteration(const CLIArguments &args, std::vector if (args.cpp) { - compiler.reset(new CompilerCPP(move(spirv_parser.get_parsed_ir()))); + compiler.reset(new CompilerCPP(std::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.reset(new CompilerMSL(move(spirv_parser.get_parsed_ir()))); + compiler.reset(new CompilerMSL(std::move(spirv_parser.get_parsed_ir()))); auto *msl_comp = static_cast(compiler.get()); auto msl_opts = msl_comp->get_msl_options(); @@ -1156,13 +1164,13 @@ static string compile_iteration(const CLIArguments &args, std::vector msl_comp->set_combined_sampler_suffix(args.msl_combined_sampler_suffix); } else if (args.hlsl) - compiler.reset(new CompilerHLSL(move(spirv_parser.get_parsed_ir()))); + compiler.reset(new CompilerHLSL(std::move(spirv_parser.get_parsed_ir()))); else { combined_image_samplers = !args.vulkan_semantics; if (!args.vulkan_semantics || args.vulkan_glsl_disable_ext_samplerless_texture_functions) build_dummy_sampler = true; - compiler.reset(new CompilerGLSL(move(spirv_parser.get_parsed_ir()))); + compiler.reset(new CompilerGLSL(std::move(spirv_parser.get_parsed_ir()))); } if (!args.variable_type_remaps.empty()) @@ -1173,7 +1181,7 @@ static string compile_iteration(const CLIArguments &args, std::vector out = remap.new_variable_type; }; - compiler->set_variable_type_remap_callback(move(remap_cb)); + compiler->set_variable_type_remap_callback(std::move(remap_cb)); } for (auto &masked : args.masked_stage_outputs) @@ -1286,6 +1294,8 @@ static string compile_iteration(const CLIArguments &args, std::vector opts.emit_line_directives = args.emit_line_directives; opts.enable_storage_image_qualifier_deduction = args.enable_storage_image_qualifier_deduction; opts.force_zero_initialized_variables = args.force_zero_initialized_variables; + opts.relax_nan_checks = args.relax_nan_checks; + opts.force_recompile_max_debug_iterations = args.force_recompile_max_debug_iterations; compiler->set_common_options(opts); for (auto &fetch : args.glsl_ext_framebuffer_fetch) @@ -1345,7 +1355,7 @@ static string compile_iteration(const CLIArguments &args, std::vector { auto active = compiler->get_active_interface_variables(); res = compiler->get_shader_resources(active); - compiler->set_enabled_interface_variables(move(active)); + compiler->set_enabled_interface_variables(std::move(active)); } else res = compiler->get_shader_resources(); @@ -1360,7 +1370,7 @@ static string compile_iteration(const CLIArguments &args, std::vector auto pls_inputs = remap_pls(args.pls_in, res.stage_inputs, &res.subpass_inputs); auto pls_outputs = remap_pls(args.pls_out, res.stage_outputs, nullptr); - compiler->remap_pixel_local_storage(move(pls_inputs), move(pls_outputs)); + compiler->remap_pixel_local_storage(std::move(pls_inputs), std::move(pls_outputs)); for (auto &ext : args.extensions) compiler->require_extension(ext); @@ -1589,7 +1599,7 @@ static int main_inner(int argc, char *argv[]) auto old_name = parser.next_string(); auto new_name = parser.next_string(); auto model = stage_to_execution_model(parser.next_string()); - args.entry_point_rename.push_back({ old_name, new_name, move(model) }); + args.entry_point_rename.push_back({ old_name, new_name, std::move(model) }); }); cbs.add("--entry", [&args](CLIParser &parser) { args.entry = parser.next_string(); }); cbs.add("--stage", [&args](CLIParser &parser) { args.entry_stage = parser.next_string(); }); @@ -1598,20 +1608,20 @@ static int main_inner(int argc, char *argv[]) HLSLVertexAttributeRemap remap; remap.location = parser.next_uint(); remap.semantic = parser.next_string(); - args.hlsl_attr_remap.push_back(move(remap)); + args.hlsl_attr_remap.push_back(std::move(remap)); }); cbs.add("--remap", [&args](CLIParser &parser) { string src = parser.next_string(); string dst = parser.next_string(); uint32_t components = parser.next_uint(); - args.remaps.push_back({ move(src), move(dst), components }); + args.remaps.push_back({ std::move(src), std::move(dst), components }); }); cbs.add("--remap-variable-type", [&args](CLIParser &parser) { string var_name = parser.next_string(); string new_type = parser.next_string(); - args.variable_type_remaps.push_back({ move(var_name), move(new_type) }); + args.variable_type_remaps.push_back({ std::move(var_name), std::move(new_type) }); }); cbs.add("--rename-interface-variable", [&args](CLIParser &parser) { @@ -1624,18 +1634,18 @@ static int main_inner(int argc, char *argv[]) uint32_t loc = parser.next_uint(); string var_name = parser.next_string(); - args.interface_variable_renames.push_back({ cls, loc, move(var_name) }); + args.interface_variable_renames.push_back({ cls, loc, std::move(var_name) }); }); cbs.add("--pls-in", [&args](CLIParser &parser) { auto fmt = pls_format(parser.next_string()); auto name = parser.next_string(); - args.pls_in.push_back({ move(fmt), move(name) }); + args.pls_in.push_back({ std::move(fmt), std::move(name) }); }); cbs.add("--pls-out", [&args](CLIParser &parser) { auto fmt = pls_format(parser.next_string()); auto name = parser.next_string(); - args.pls_out.push_back({ move(fmt), move(name) }); + args.pls_out.push_back({ std::move(fmt), std::move(name) }); }); cbs.add("--shader-model", [&args](CLIParser &parser) { args.shader_model = parser.next_uint(); @@ -1678,11 +1688,17 @@ static int main_inner(int argc, char *argv[]) args.masked_stage_builtins.push_back(masked_builtin); }); + cbs.add("--force-recompile-max-debug-iterations", [&](CLIParser &parser) { + args.force_recompile_max_debug_iterations = parser.next_uint(); + }); + + cbs.add("--relax-nan-checks", [&](CLIParser &) { args.relax_nan_checks = true; }); + cbs.default_handler = [&args](const char *value) { args.input = value; }; cbs.add("-", [&args](CLIParser &) { args.input = "-"; }); cbs.error_handler = [] { print_help(); }; - CLIParser parser{ move(cbs), argc - 1, argv + 1 }; + CLIParser parser{ std::move(cbs), argc - 1, argv + 1 }; if (!parser.parse()) return EXIT_FAILURE; else if (parser.ended_state) @@ -1702,10 +1718,10 @@ static int main_inner(int argc, char *argv[]) // Special case reflection because it has little to do with the path followed by code-outputting compilers if (!args.reflect.empty()) { - Parser spirv_parser(move(spirv_file)); + Parser spirv_parser(std::move(spirv_file)); spirv_parser.parse(); - CompilerReflection compiler(move(spirv_parser.get_parsed_ir())); + CompilerReflection compiler(std::move(spirv_parser.get_parsed_ir())); compiler.set_format(args.reflect); auto json = compiler.compile(); if (args.output) @@ -1718,7 +1734,7 @@ static int main_inner(int argc, char *argv[]) string compiled_output; if (args.iterations == 1) - compiled_output = compile_iteration(args, move(spirv_file)); + compiled_output = compile_iteration(args, std::move(spirv_file)); else { for (unsigned i = 0; i < args.iterations; i++) diff --git a/3rdparty/spirv-cross/spirv_common.hpp b/3rdparty/spirv-cross/spirv_common.hpp index e656a9a07..4aaa71488 100644 --- a/3rdparty/spirv-cross/spirv_common.hpp +++ b/3rdparty/spirv-cross/spirv_common.hpp @@ -690,7 +690,7 @@ struct SPIRExpression : IVariant // Only created by the backend target to avoid creating tons of temporaries. SPIRExpression(std::string expr, TypeID expression_type_, bool immutable_) - : expression(move(expr)) + : expression(std::move(expr)) , expression_type(expression_type_) , immutable(immutable_) { @@ -1072,7 +1072,6 @@ struct SPIRVariable : IVariant // Temporaries which can remain forwarded as long as this variable is not modified. SmallVector dependees; - bool forwardable = true; bool deferred_declaration = false; bool phi_variable = false; diff --git a/3rdparty/spirv-cross/spirv_cross.cpp b/3rdparty/spirv-cross/spirv_cross.cpp index cf60197e6..a68ef7578 100644 --- a/3rdparty/spirv-cross/spirv_cross.cpp +++ b/3rdparty/spirv-cross/spirv_cross.cpp @@ -36,16 +36,16 @@ using namespace SPIRV_CROSS_NAMESPACE; Compiler::Compiler(vector ir_) { - Parser parser(move(ir_)); + Parser parser(std::move(ir_)); parser.parse(); - set_ir(move(parser.get_parsed_ir())); + set_ir(std::move(parser.get_parsed_ir())); } Compiler::Compiler(const uint32_t *ir_, size_t word_count) { Parser parser(ir_, word_count); parser.parse(); - set_ir(move(parser.get_parsed_ir())); + set_ir(std::move(parser.get_parsed_ir())); } Compiler::Compiler(const ParsedIR &ir_) @@ -55,12 +55,12 @@ Compiler::Compiler(const ParsedIR &ir_) Compiler::Compiler(ParsedIR &&ir_) { - set_ir(move(ir_)); + set_ir(std::move(ir_)); } void Compiler::set_ir(ParsedIR &&ir_) { - ir = move(ir_); + ir = std::move(ir_); parse_fixup(); } @@ -852,7 +852,7 @@ unordered_set Compiler::get_active_interface_variables() const void Compiler::set_enabled_interface_variables(std::unordered_set active_variables) { - active_interface_variables = move(active_variables); + active_interface_variables = std::move(active_variables); check_active_interface_variables = true; } @@ -1677,6 +1677,11 @@ const SmallVector &Compiler::get_case_list(const SPIRBlock &blo const auto &type = get(var->basetype); width = type.width; } + else if (const auto *undef = maybe_get(block.condition)) + { + const auto &type = get(undef->basetype); + width = type.width; + } else { auto search = ir.load_type_width.find(block.condition); @@ -2508,7 +2513,7 @@ void Compiler::CombinedImageSamplerHandler::push_remap_parameters(const SPIRFunc unordered_map remapping; for (uint32_t i = 0; i < length; i++) remapping[func.arguments[i].id] = remap_parameter(args[i]); - parameter_remapping.push(move(remapping)); + parameter_remapping.push(std::move(remapping)); } void Compiler::CombinedImageSamplerHandler::pop_remap_parameters() @@ -4362,7 +4367,7 @@ void Compiler::analyze_image_and_sampler_usage() handler.dependency_hierarchy.clear(); traverse_all_reachable_opcodes(get(ir.default_entry_point), handler); - comparison_ids = move(handler.comparison_ids); + comparison_ids = std::move(handler.comparison_ids); need_subpass_input = handler.need_subpass_input; // Forward information from separate images and samplers into combined image samplers. @@ -4415,7 +4420,7 @@ void Compiler::build_function_control_flow_graphs_and_analyze() CFGBuilder handler(*this); 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); + function_cfgs = std::move(handler.function_cfgs); bool single_function = function_cfgs.size() <= 1; for (auto &f : function_cfgs) @@ -5024,7 +5029,7 @@ void Compiler::analyze_non_block_pointer_types() for (auto type : handler.non_block_types) physical_storage_non_block_pointer_types.push_back(type); sort(begin(physical_storage_non_block_pointer_types), end(physical_storage_non_block_pointer_types)); - physical_storage_type_to_alignment = move(handler.physical_block_type_meta); + physical_storage_type_to_alignment = std::move(handler.physical_block_type_meta); } bool Compiler::InterlockedResourceAccessPrepassHandler::handle(Op op, const uint32_t *, uint32_t) diff --git a/3rdparty/spirv-cross/spirv_cross_c.cpp b/3rdparty/spirv-cross/spirv_cross_c.cpp index 4d5615404..4a62b635c 100644 --- a/3rdparty/spirv-cross/spirv_cross_c.cpp +++ b/3rdparty/spirv-cross/spirv_cross_c.cpp @@ -251,7 +251,7 @@ spvc_result spvc_context_parse_spirv(spvc_context context, const SpvId *spirv, s pir->context = context; Parser parser(spirv, word_count); parser.parse(); - pir->parsed = move(parser.get_parsed_ir()); + pir->parsed = std::move(parser.get_parsed_ir()); *parsed_ir = pir.get(); context->allocations.push_back(std::move(pir)); } @@ -283,7 +283,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back { case SPVC_BACKEND_NONE: if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP) - comp->compiler.reset(new Compiler(move(parsed_ir->parsed))); + comp->compiler.reset(new Compiler(std::move(parsed_ir->parsed))); else if (mode == SPVC_CAPTURE_MODE_COPY) comp->compiler.reset(new Compiler(parsed_ir->parsed)); break; @@ -291,7 +291,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back #if SPIRV_CROSS_C_API_GLSL case SPVC_BACKEND_GLSL: if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP) - comp->compiler.reset(new CompilerGLSL(move(parsed_ir->parsed))); + comp->compiler.reset(new CompilerGLSL(std::move(parsed_ir->parsed))); else if (mode == SPVC_CAPTURE_MODE_COPY) comp->compiler.reset(new CompilerGLSL(parsed_ir->parsed)); break; @@ -300,7 +300,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back #if SPIRV_CROSS_C_API_HLSL case SPVC_BACKEND_HLSL: if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP) - comp->compiler.reset(new CompilerHLSL(move(parsed_ir->parsed))); + comp->compiler.reset(new CompilerHLSL(std::move(parsed_ir->parsed))); else if (mode == SPVC_CAPTURE_MODE_COPY) comp->compiler.reset(new CompilerHLSL(parsed_ir->parsed)); break; @@ -309,7 +309,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back #if SPIRV_CROSS_C_API_MSL case SPVC_BACKEND_MSL: if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP) - comp->compiler.reset(new CompilerMSL(move(parsed_ir->parsed))); + comp->compiler.reset(new CompilerMSL(std::move(parsed_ir->parsed))); else if (mode == SPVC_CAPTURE_MODE_COPY) comp->compiler.reset(new CompilerMSL(parsed_ir->parsed)); break; @@ -318,7 +318,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back #if SPIRV_CROSS_C_API_CPP case SPVC_BACKEND_CPP: if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP) - comp->compiler.reset(new CompilerCPP(move(parsed_ir->parsed))); + comp->compiler.reset(new CompilerCPP(std::move(parsed_ir->parsed))); else if (mode == SPVC_CAPTURE_MODE_COPY) comp->compiler.reset(new CompilerCPP(parsed_ir->parsed)); break; @@ -327,7 +327,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back #if SPIRV_CROSS_C_API_REFLECT case SPVC_BACKEND_JSON: if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP) - comp->compiler.reset(new CompilerReflection(move(parsed_ir->parsed))); + comp->compiler.reset(new CompilerReflection(std::move(parsed_ir->parsed))); else if (mode == SPVC_CAPTURE_MODE_COPY) comp->compiler.reset(new CompilerReflection(parsed_ir->parsed)); break; @@ -475,6 +475,9 @@ spvc_result spvc_compiler_options_set_uint(spvc_compiler_options options, spvc_c case SPVC_COMPILER_OPTION_GLSL_OVR_MULTIVIEW_VIEW_COUNT: options->glsl.ovr_multiview_view_count = value; break; + case SPVC_COMPILER_OPTION_RELAX_NAN_CHECKS: + options->glsl.relax_nan_checks = value != 0; + break; #endif #if SPIRV_CROSS_C_API_HLSL diff --git a/3rdparty/spirv-cross/spirv_cross_c.h b/3rdparty/spirv-cross/spirv_cross_c.h index a590c805a..a35a5d651 100644 --- a/3rdparty/spirv-cross/spirv_cross_c.h +++ b/3rdparty/spirv-cross/spirv_cross_c.h @@ -40,7 +40,7 @@ extern "C" { /* Bumped if ABI or API breaks backwards compatibility. */ #define SPVC_C_API_VERSION_MAJOR 0 /* Bumped if APIs or enumerations are added in a backwards compatible way. */ -#define SPVC_C_API_VERSION_MINOR 48 +#define SPVC_C_API_VERSION_MINOR 49 /* Bumped if internal implementation details change. */ #define SPVC_C_API_VERSION_PATCH 0 @@ -677,6 +677,8 @@ typedef enum spvc_compiler_option SPVC_COMPILER_OPTION_GLSL_OVR_MULTIVIEW_VIEW_COUNT = 77 | SPVC_COMPILER_OPTION_GLSL_BIT, + SPVC_COMPILER_OPTION_RELAX_NAN_CHECKS = 78 | SPVC_COMPILER_OPTION_COMMON_BIT, + SPVC_COMPILER_OPTION_INT_MAX = 0x7fffffff } spvc_compiler_option; diff --git a/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp b/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp index e7fcdff04..ce2f98f13 100644 --- a/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp +++ b/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp @@ -54,26 +54,26 @@ ParsedIR::ParsedIR() // Should have been default-implemented, but need this on MSVC 2013. ParsedIR::ParsedIR(ParsedIR &&other) SPIRV_CROSS_NOEXCEPT { - *this = move(other); + *this = std::move(other); } ParsedIR &ParsedIR::operator=(ParsedIR &&other) SPIRV_CROSS_NOEXCEPT { if (this != &other) { - pool_group = move(other.pool_group); - spirv = move(other.spirv); - meta = move(other.meta); + pool_group = std::move(other.pool_group); + spirv = std::move(other.spirv); + meta = std::move(other.meta); for (int i = 0; i < TypeCount; i++) - ids_for_type[i] = move(other.ids_for_type[i]); - ids_for_constant_or_type = move(other.ids_for_constant_or_type); - ids_for_constant_or_variable = move(other.ids_for_constant_or_variable); - declared_capabilities = move(other.declared_capabilities); - declared_extensions = move(other.declared_extensions); - block_meta = move(other.block_meta); - continue_block_to_loop_header = move(other.continue_block_to_loop_header); - entry_points = move(other.entry_points); - ids = move(other.ids); + ids_for_type[i] = std::move(other.ids_for_type[i]); + ids_for_constant_or_type = std::move(other.ids_for_constant_or_type); + ids_for_constant_or_variable = std::move(other.ids_for_constant_or_variable); + declared_capabilities = std::move(other.declared_capabilities); + declared_extensions = std::move(other.declared_extensions); + block_meta = std::move(other.block_meta); + continue_block_to_loop_header = std::move(other.continue_block_to_loop_header); + entry_points = std::move(other.entry_points); + ids = std::move(other.ids); addressing_model = other.addressing_model; memory_model = other.memory_model; @@ -999,7 +999,7 @@ ParsedIR::LoopLock::LoopLock(uint32_t *lock_) ParsedIR::LoopLock::LoopLock(LoopLock &&other) SPIRV_CROSS_NOEXCEPT { - *this = move(other); + *this = std::move(other); } ParsedIR::LoopLock &ParsedIR::LoopLock::operator=(LoopLock &&other) SPIRV_CROSS_NOEXCEPT diff --git a/3rdparty/spirv-cross/spirv_glsl.cpp b/3rdparty/spirv-cross/spirv_glsl.cpp index 3cdb742f0..52cef4ec2 100644 --- a/3rdparty/spirv-cross/spirv_glsl.cpp +++ b/3rdparty/spirv-cross/spirv_glsl.cpp @@ -306,8 +306,8 @@ void CompilerGLSL::reset(uint32_t iteration_count) // It is highly context-sensitive when we need to force recompilation, // and it is not practical with the current architecture // to resolve everything up front. - if (iteration_count >= 3 && !is_force_recompile_forward_progress) - SPIRV_CROSS_THROW("Over 3 compilation loops detected and no forward progress was made. Must be a bug!"); + if (iteration_count >= options.force_recompile_max_debug_iterations && !is_force_recompile_forward_progress) + SPIRV_CROSS_THROW("Maximum compilation loops detected and no forward progress was made. Must be a SPIRV-Cross bug!"); // We do some speculative optimizations which should pretty much always work out, // but just in case the SPIR-V is rather weird, recompile until it's happy. @@ -4298,10 +4298,8 @@ string CompilerGLSL::to_func_call_arg(const SPIRFunction::Parameter &, uint32_t return to_expression(name_id); } -void CompilerGLSL::handle_invalid_expression(uint32_t id) +void CompilerGLSL::force_temporary_and_recompile(uint32_t id) { - // We tried to read an invalidated expression. - // This means we need another pass at compilation, but next time, force temporary variables so that they cannot be invalidated. auto res = forced_temporaries.insert(id); // Forcing new temporaries guarantees forward progress. @@ -4311,6 +4309,14 @@ void CompilerGLSL::handle_invalid_expression(uint32_t id) force_recompile(); } +void CompilerGLSL::handle_invalid_expression(uint32_t id) +{ + // We tried to read an invalidated expression. + // This means we need another pass at compilation, but next time, + // force temporary variables so that they cannot be invalidated. + force_temporary_and_recompile(id); +} + // Converts the format of the current expression from packed to unpacked, // by wrapping the expression in a constructor of the appropriate type. // GLSL does not support packed formats, so simply return the expression. @@ -7237,18 +7243,11 @@ string CompilerGLSL::to_function_args(const TextureFunctionArguments &args, bool forward = forward && should_forward(args.lod); farg_str += ", "; - auto &lod_expr_type = expression_type(args.lod); - // Lod expression for TexelFetch in GLSL must be int, and only int. - if (args.base.is_fetch && imgtype.image.dim != DimBuffer && !imgtype.image.ms && - lod_expr_type.basetype != SPIRType::Int) - { - farg_str += join("int(", to_expression(args.lod), ")"); - } + if (args.base.is_fetch && imgtype.image.dim != DimBuffer && !imgtype.image.ms) + farg_str += bitcast_expression(SPIRType::Int, args.lod); else - { farg_str += to_expression(args.lod); - } } } else if (args.base.is_fetch && imgtype.image.dim != DimBuffer && !imgtype.image.ms) @@ -7261,19 +7260,19 @@ string CompilerGLSL::to_function_args(const TextureFunctionArguments &args, bool { forward = forward && should_forward(args.coffset); farg_str += ", "; - farg_str += to_expression(args.coffset); + farg_str += bitcast_expression(SPIRType::Int, args.coffset); } else if (args.offset) { forward = forward && should_forward(args.offset); farg_str += ", "; - farg_str += to_expression(args.offset); + farg_str += bitcast_expression(SPIRType::Int, args.offset); } if (args.sample) { farg_str += ", "; - farg_str += to_expression(args.sample); + farg_str += bitcast_expression(SPIRType::Int, args.sample); } if (args.min_lod) @@ -7300,11 +7299,7 @@ string CompilerGLSL::to_function_args(const TextureFunctionArguments &args, bool { forward = forward && should_forward(args.component); farg_str += ", "; - auto &component_type = expression_type(args.component); - if (component_type.basetype == SPIRType::Int) - farg_str += to_expression(args.component); - else - farg_str += join("int(", to_expression(args.component), ")"); + farg_str += bitcast_expression(SPIRType::Int, args.component); } *p_forward = forward; @@ -7312,6 +7307,63 @@ string CompilerGLSL::to_function_args(const TextureFunctionArguments &args, bool return farg_str; } +Op CompilerGLSL::get_remapped_spirv_op(Op op) const +{ + if (options.relax_nan_checks) + { + switch (op) + { + case OpFUnordLessThan: + op = OpFOrdLessThan; + break; + case OpFUnordLessThanEqual: + op = OpFOrdLessThanEqual; + break; + case OpFUnordGreaterThan: + op = OpFOrdGreaterThan; + break; + case OpFUnordGreaterThanEqual: + op = OpFOrdGreaterThanEqual; + break; + case OpFUnordEqual: + op = OpFOrdEqual; + break; + case OpFOrdNotEqual: + op = OpFUnordNotEqual; + break; + + default: + break; + } + } + + return op; +} + +GLSLstd450 CompilerGLSL::get_remapped_glsl_op(GLSLstd450 std450_op) const +{ + // Relax to non-NaN aware opcodes. + if (options.relax_nan_checks) + { + switch (std450_op) + { + case GLSLstd450NClamp: + std450_op = GLSLstd450FClamp; + break; + case GLSLstd450NMin: + std450_op = GLSLstd450FMin; + break; + case GLSLstd450NMax: + std450_op = GLSLstd450FMax; + break; + default: + break; + } + } + + return std450_op; +} + void CompilerGLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop, const uint32_t *args, uint32_t length) { auto op = static_cast(eop); @@ -7324,6 +7376,8 @@ void CompilerGLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop, auto int_type = to_signed_basetype(integer_width); auto uint_type = to_unsigned_basetype(integer_width); + op = get_remapped_glsl_op(op); + switch (op) { // FP fiddling @@ -9189,8 +9243,13 @@ std::string CompilerGLSL::flattened_access_chain_struct(uint32_t base, const uin { std::string expr; - expr += type_to_glsl_constructor(target_type); - expr += "("; + if (backend.can_declare_struct_inline) + { + expr += type_to_glsl_constructor(target_type); + expr += "("; + } + else + expr += "{"; for (uint32_t i = 0; i < uint32_t(target_type.member_types.size()); ++i) { @@ -9220,7 +9279,7 @@ std::string CompilerGLSL::flattened_access_chain_struct(uint32_t base, const uin expr += tmp; } - expr += ")"; + expr += backend.can_declare_struct_inline ? ")" : "}"; return expr; } @@ -9526,8 +9585,11 @@ bool CompilerGLSL::should_forward(uint32_t id) const // This is important because otherwise we'll get local sampler copies (highp sampler2D foo = bar) that are invalid in OpenGL GLSL auto *var = maybe_get(id); - if (var && var->forwardable) - return true; + if (var) + { + // Never forward volatile variables, e.g. SPIR-V 1.6 IsHelperInvocation. + return !has_decoration(id, DecorationVolatile); + } // For debugging emit temporary variables for all expressions if (options.force_temporary) @@ -9540,6 +9602,12 @@ bool CompilerGLSL::should_forward(uint32_t id) const if (expr && expr->expression_dependencies.size() >= max_expression_dependencies) return false; + if (expr && expr->loaded_from && has_decoration(expr->loaded_from, DecorationVolatile)) + { + // Never forward volatile variables. + return false; + } + // Immutable expression can always be forwarded. if (is_immutable(id)) return true; @@ -9596,9 +9664,8 @@ void CompilerGLSL::track_expression_read(uint32_t id) //if (v == 2) // fprintf(stderr, "ID %u was forced to temporary due to more than 1 expression use!\n", id); - forced_temporaries.insert(id); // Force a recompile after this pass to avoid forwarding this variable. - force_recompile(); + force_temporary_and_recompile(id); } } } @@ -9952,9 +10019,8 @@ void CompilerGLSL::disallow_forwarding_in_expression_chain(const SPIRExpression if (expression_is_forwarded(expr.self) && !expression_suppresses_usage_tracking(expr.self) && forced_invariant_temporaries.count(expr.self) == 0) { - forced_temporaries.insert(expr.self); + force_temporary_and_recompile(expr.self); forced_invariant_temporaries.insert(expr.self); - force_recompile(); for (auto &dependent : expr.expression_dependencies) disallow_forwarding_in_expression_chain(get(dependent)); @@ -10096,6 +10162,8 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) auto int_type = to_signed_basetype(integer_width); auto uint_type = to_unsigned_basetype(integer_width); + opcode = get_remapped_spirv_op(opcode); + switch (opcode) { // Dealing with memory @@ -10239,9 +10307,19 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) // If an expression is mutable and forwardable, we speculate that it is immutable. AccessChainMeta meta; bool ptr_chain = opcode == OpPtrAccessChain; - auto e = access_chain(ops[2], &ops[3], length - 3, get(ops[0]), &meta, ptr_chain); + auto &target_type = get(ops[0]); + auto e = access_chain(ops[2], &ops[3], length - 3, target_type, &meta, ptr_chain); - auto &expr = set(ops[1], move(e), ops[0], should_forward(ops[2])); + // If the base is flattened UBO of struct type, the expression has to be a composite. + // In that case, backends which do not support inline syntax need it to be bound to a temporary. + // Otherwise, invalid expressions like ({UBO[0].xyz, UBO[0].w, UBO[1]}).member are emitted. + bool requires_temporary = false; + if (flattened_buffer_blocks.count(ops[2]) && target_type.basetype == SPIRType::Struct) + requires_temporary = !backend.can_declare_struct_inline; + + auto &expr = requires_temporary ? + emit_op(ops[0], ops[1], std::move(e), false) : + set(ops[1], std::move(e), ops[0], should_forward(ops[2])); auto *backing_variable = maybe_get_backing_variable(ops[2]); expr.loaded_from = backing_variable ? backing_variable->self : ID(ops[2]); @@ -11212,7 +11290,11 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) case OpLogicalNotEqual: case OpFOrdNotEqual: + case OpFUnordNotEqual: { + // GLSL is fuzzy on what to do with ordered vs unordered not equal. + // glslang started emitting UnorderedNotEqual some time ago to harmonize with IEEE, + // but this means we have no easy way of implementing ordered not equal. if (expression_type(ops[2]).vecsize > 1) GLSL_BFOP(notEqual); else @@ -12016,7 +12098,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) pure = false; } - if (var && var->forwardable) + if (var) { bool forward = forced_temporaries.find(id) == end(forced_temporaries); auto &e = emit_op(result_type, id, imgexpr, forward); @@ -12532,7 +12614,6 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) break; case OpFUnordEqual: - case OpFUnordNotEqual: case OpFUnordLessThan: case OpFUnordGreaterThan: case OpFUnordLessThanEqual: @@ -12555,10 +12636,6 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) comp_op = "notEqual"; break; - case OpFUnordNotEqual: - comp_op = "equal"; - break; - case OpFUnordLessThan: comp_op = "greaterThanEqual"; break; @@ -12591,10 +12668,6 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) comp_op = " != "; break; - case OpFUnordNotEqual: - comp_op = " == "; - break; - case OpFUnordLessThan: comp_op = " >= "; break; @@ -12786,6 +12859,8 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) if (!options.vulkan_semantics) SPIRV_CROSS_THROW("GL_EXT_demote_to_helper_invocation is only supported in Vulkan GLSL."); require_extension_internal("GL_EXT_demote_to_helper_invocation"); + // Helper lane state with demote is volatile by nature. + // Do not forward this. emit_op(ops[0], ops[1], "helperInvocationEXT()", false); break; @@ -14637,9 +14712,16 @@ void CompilerGLSL::emit_hoisted_temporaries(SmallVector> &tempo for (auto &tmp : temporaries) { + auto &type = get(tmp.first); + + // There are some rare scenarios where we are asked to declare pointer types as hoisted temporaries. + // This should be ignored unless we're doing actual variable pointers and backend supports it. + // Access chains cannot normally be lowered to temporaries in GLSL and HLSL. + if (type.pointer && !backend.native_pointers) + continue; + add_local_variable_name(tmp.second); auto &flags = ir.meta[tmp.second].decoration.decoration_flags; - auto &type = get(tmp.first); // Not all targets support pointer literals, so don't bother with that case. string initializer; @@ -15502,7 +15584,7 @@ void CompilerGLSL::unroll_array_from_complex_load(uint32_t target_id, uint32_t s statement(new_expr, "[i] = ", expr, "[i];"); end_scope(); - expr = move(new_expr); + expr = std::move(new_expr); } } @@ -15837,7 +15919,7 @@ void CompilerGLSL::emit_copy_logical_type(uint32_t lhs_id, uint32_t lhs_type_id, rhs_id = id + 1; { - auto &lhs_expr = set(lhs_id, move(lhs), lhs_type_id, true); + auto &lhs_expr = set(lhs_id, std::move(lhs), lhs_type_id, true); lhs_expr.need_transpose = lhs_meta.need_transpose; if (lhs_meta.storage_is_packed) @@ -15850,7 +15932,7 @@ void CompilerGLSL::emit_copy_logical_type(uint32_t lhs_id, uint32_t lhs_type_id, } { - auto &rhs_expr = set(rhs_id, move(rhs), rhs_type_id, true); + auto &rhs_expr = set(rhs_id, std::move(rhs), rhs_type_id, true); rhs_expr.need_transpose = rhs_meta.need_transpose; if (rhs_meta.storage_is_packed) diff --git a/3rdparty/spirv-cross/spirv_glsl.hpp b/3rdparty/spirv-cross/spirv_glsl.hpp index 6bf9183f2..b892e0c3b 100644 --- a/3rdparty/spirv-cross/spirv_glsl.hpp +++ b/3rdparty/spirv-cross/spirv_glsl.hpp @@ -83,6 +83,11 @@ public: // Debug option to always emit temporary variables for all expressions. bool force_temporary = false; + // Debug option, can be increased in an attempt to workaround SPIRV-Cross bugs temporarily. + // If this limit has to be increased, it points to an implementation bug. + // In certain scenarios, the maximum number of debug iterations may increase beyond this limit + // as long as we can prove we're making certain kinds of forward progress. + uint32_t force_recompile_max_debug_iterations = 3; // If true, Vulkan GLSL features are used instead of GL-compatible features. // Mostly useful for debugging SPIR-V files. @@ -133,6 +138,13 @@ public: // what happens on legacy GLSL targets for blocks and structs. bool force_flattened_io_blocks = false; + // For opcodes where we have to perform explicit additional nan checks, very ugly code is generated. + // If we opt-in, ignore these requirements. + // In opcodes like NClamp/NMin/NMax and FP compare, ignore NaN behavior. + // Use FClamp/FMin/FMax semantics for clamps and lets implementation choose ordered or unordered + // compares. + bool relax_nan_checks = false; + // If non-zero, controls layout(num_views = N) in; in GL_OVR_multiview2. uint32_t ovr_multiview_view_count = 0; @@ -357,6 +369,11 @@ protected: virtual void emit_instruction(const Instruction &instr); void emit_block_instructions(SPIRBlock &block); + + // For relax_nan_checks. + GLSLstd450 get_remapped_glsl_op(GLSLstd450 std450_op) const; + spv::Op get_remapped_spirv_op(spv::Op op) const; + virtual void emit_glsl_op(uint32_t result_type, uint32_t result_id, uint32_t op, const uint32_t *args, uint32_t count); virtual void emit_spv_amd_shader_ballot_op(uint32_t result_type, uint32_t result_id, uint32_t op, @@ -881,6 +898,7 @@ protected: void check_function_call_constraints(const uint32_t *args, uint32_t length); void handle_invalid_expression(uint32_t id); + void force_temporary_and_recompile(uint32_t id); void find_static_extensions(); std::string emit_for_loop_initializers(const SPIRBlock &block); diff --git a/3rdparty/spirv-cross/spirv_hlsl.cpp b/3rdparty/spirv-cross/spirv_hlsl.cpp index 3d8347749..d0d0e15ac 100644 --- a/3rdparty/spirv-cross/spirv_hlsl.cpp +++ b/3rdparty/spirv-cross/spirv_hlsl.cpp @@ -728,6 +728,11 @@ void CompilerHLSL::emit_builtin_inputs_in_struct() // Handled specially. break; + case BuiltInHelperInvocation: + if (hlsl_options.shader_model < 50 || get_entry_point().model != ExecutionModelFragment) + SPIRV_CROSS_THROW("Helper Invocation input is only supported in PS 5.0 or higher."); + break; + case BuiltInClipDistance: // HLSL is a bit weird here, use SV_ClipDistance0, SV_ClipDistance1 and so on with vectors. for (uint32_t clip = 0; clip < clip_distance_count; clip += 4) @@ -984,6 +989,8 @@ std::string CompilerHLSL::builtin_to_glsl(spv::BuiltIn builtin, spv::StorageClas return "WaveGetLaneIndex()"; case BuiltInSubgroupSize: return "WaveGetLaneCount()"; + case BuiltInHelperInvocation: + return "IsHelperLane()"; default: return CompilerGLSL::builtin_to_glsl(builtin, storage); @@ -1103,6 +1110,11 @@ void CompilerHLSL::emit_builtin_variables() type = "uint4"; break; + case BuiltInHelperInvocation: + if (hlsl_options.shader_model < 50) + SPIRV_CROSS_THROW("Need SM 5.0 for Helper Invocation."); + break; + case BuiltInClipDistance: array_size = clip_distance_count; type = "float"; @@ -1335,7 +1347,8 @@ void CompilerHLSL::emit_resources() } }); - if (execution.model == ExecutionModelVertex && hlsl_options.shader_model <= 30) + if (execution.model == ExecutionModelVertex && hlsl_options.shader_model <= 30 && + active_output_builtins.get(BuiltInPosition)) { statement("uniform float4 gl_HalfPixel;"); emitted = true; @@ -2102,7 +2115,11 @@ void CompilerHLSL::emit_buffer_block(const SPIRVariable &var) bool is_uav = var.storage == StorageClassStorageBuffer || has_decoration(type.self, DecorationBufferBlock); - if (is_uav) + if (flattened_buffer_blocks.count(var.self)) + { + emit_buffer_block_flattened(var); + } + else if (is_uav) { Bitset flags = ir.get_buffer_block_flags(var); bool is_readonly = flags.get(DecorationNonWritable) && !is_hlsl_force_storage_buffer_as_uav(var.self); @@ -2207,7 +2224,11 @@ void CompilerHLSL::emit_buffer_block(const SPIRVariable &var) void CompilerHLSL::emit_push_constant_block(const SPIRVariable &var) { - if (root_constants_layout.empty()) + if (flattened_buffer_blocks.count(var.self)) + { + emit_buffer_block_flattened(var); + } + else if (root_constants_layout.empty()) { emit_buffer_block(var); } @@ -2365,7 +2386,7 @@ void CompilerHLSL::emit_function_prototype(SPIRFunction &func, const Bitset &ret out_argument += " "; out_argument += "spvReturnValue"; out_argument += type_to_array_glsl(type); - arglist.push_back(move(out_argument)); + arglist.push_back(std::move(out_argument)); } for (auto &arg : func.arguments) @@ -2512,6 +2533,7 @@ void CompilerHLSL::emit_hlsl_entry_point() case BuiltInPointCoord: case BuiltInSubgroupSize: case BuiltInSubgroupLocalInvocationId: + case BuiltInHelperInvocation: break; case BuiltInSubgroupEqMask: @@ -2731,7 +2753,7 @@ void CompilerHLSL::emit_hlsl_entry_point() void CompilerHLSL::emit_fixup() { - if (is_vertex_like_shader()) + if (is_vertex_like_shader() && active_output_builtins.get(BuiltInPosition)) { // Do various mangling on the gl_Position. if (hlsl_options.shader_model <= 30) @@ -3515,6 +3537,8 @@ void CompilerHLSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop, auto int_type = to_signed_basetype(integer_width); auto uint_type = to_unsigned_basetype(integer_width); + op = get_remapped_glsl_op(op); + switch (op) { case GLSLstd450InverseSqrt: @@ -3987,7 +4011,7 @@ void CompilerHLSL::read_access_chain(string *expr, const string &lhs, const SPIR if (lhs.empty()) { assert(expr); - *expr = move(load_expr); + *expr = std::move(load_expr); } else statement(lhs, " = ", load_expr, ";"); @@ -4783,6 +4807,8 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction) auto int_type = to_signed_basetype(integer_width); auto uint_type = to_unsigned_basetype(integer_width); + opcode = get_remapped_spirv_op(opcode); + switch (opcode) { case OpAccessChain: @@ -5333,7 +5359,7 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction) image_format_to_components(get(var->basetype).image.format), imgexpr); } - if (var && var->forwardable) + if (var) { bool forward = forced_temporaries.find(id) == end(forced_temporaries); auto &e = emit_op(result_type, id, imgexpr, forward); @@ -5577,7 +5603,12 @@ void CompilerHLSL::emit_instruction(const Instruction &instruction) } case OpIsHelperInvocationEXT: - SPIRV_CROSS_THROW("helperInvocationEXT() is not supported in HLSL."); + if (hlsl_options.shader_model < 50 || get_entry_point().model != ExecutionModelFragment) + SPIRV_CROSS_THROW("Helper Invocation input is only supported in PS 5.0 or higher."); + // Helper lane state with demote is volatile by nature. + // Do not forward this. + emit_op(ops[0], ops[1], "IsHelperLane()", false); + break; case OpBeginInvocationInterlockEXT: case OpEndInvocationInterlockEXT: @@ -5664,7 +5695,7 @@ void CompilerHLSL::require_texture_query_variant(uint32_t var_id) void CompilerHLSL::set_root_constant_layouts(std::vector layout) { - root_constants_layout = move(layout); + root_constants_layout = std::move(layout); } void CompilerHLSL::add_vertex_attribute_remap(const HLSLVertexAttributeRemap &vertex_attributes) diff --git a/3rdparty/spirv-cross/spirv_msl.cpp b/3rdparty/spirv-cross/spirv_msl.cpp index 52c232ab8..8ea27d5f6 100644 --- a/3rdparty/spirv-cross/spirv_msl.cpp +++ b/3rdparty/spirv-cross/spirv_msl.cpp @@ -37,7 +37,7 @@ static const uint32_t k_unknown_component = ~0u; static const char *force_inline = "static inline __attribute__((always_inline))"; CompilerMSL::CompilerMSL(std::vector spirv_) - : CompilerGLSL(move(spirv_)) + : CompilerGLSL(std::move(spirv_)) { } @@ -1548,6 +1548,14 @@ void CompilerMSL::extract_global_variables_from_functions() // Uniforms unordered_set global_var_ids; ir.for_each_typed_id([&](uint32_t, SPIRVariable &var) { + // Some builtins resolve directly to a function call which does not need any declared variables. + // Skip these. + if (var.storage == StorageClassInput && has_decoration(var.self, DecorationBuiltIn) && + BuiltIn(get_decoration(var.self, DecorationBuiltIn)) == BuiltInHelperInvocation) + { + return; + } + if (var.storage == StorageClassInput || var.storage == StorageClassOutput || var.storage == StorageClassUniform || var.storage == StorageClassUniformConstant || var.storage == StorageClassPushConstant || var.storage == StorageClassStorageBuffer) @@ -2508,12 +2516,15 @@ void CompilerMSL::add_composite_variable_to_interface_block(StorageClass storage } } -void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass storage, const string &ib_var_ref, - SPIRType &ib_type, SPIRVariable &var, - uint32_t mbr_idx, InterfaceBlockMeta &meta) +void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass storage, + const string &ib_var_ref, SPIRType &ib_type, + SPIRVariable &var, SPIRType &var_type, + uint32_t mbr_idx, InterfaceBlockMeta &meta, + const string &mbr_name_qual, + const string &var_chain_qual, + uint32_t &location, uint32_t &var_mbr_idx) { auto &entry_func = get(ir.default_entry_point); - auto &var_type = meta.strip_array ? get_variable_element_type(var) : get_variable_data_type(var); BuiltIn builtin = BuiltInMax; bool is_builtin = is_member_builtin(var_type, mbr_idx, &builtin); @@ -2528,8 +2539,8 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass uint32_t mbr_type_id = var_type.member_types[mbr_idx]; auto &mbr_type = get(mbr_type_id); - uint32_t elem_cnt = 0; + uint32_t elem_cnt = 1; if (is_matrix(mbr_type)) { if (is_array(mbr_type)) @@ -2572,6 +2583,27 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass return; } + // Recursively handle nested structures. + if (mbr_type.basetype == SPIRType::Struct) + { + for (uint32_t i = 0; i < elem_cnt; i++) + { + string mbr_name = append_member_name(mbr_name_qual, var_type, mbr_idx) + (elem_cnt == 1 ? "" : join("_", i)); + string var_chain = join(var_chain_qual, ".", to_member_name(var_type, mbr_idx), (elem_cnt == 1 ? "" : join("[", i, "]"))); + uint32_t sub_mbr_cnt = uint32_t(mbr_type.member_types.size()); + for (uint32_t sub_mbr_idx = 0; sub_mbr_idx < sub_mbr_cnt; sub_mbr_idx++) + { + add_composite_member_variable_to_interface_block(storage, ib_var_ref, ib_type, + var, mbr_type, sub_mbr_idx, + meta, mbr_name, var_chain, + location, var_mbr_idx); + // FIXME: Recursive structs and tessellation breaks here. + var_mbr_idx++; + } + } + return; + } + for (uint32_t i = 0; i < elem_cnt; i++) { // Add a reference to the variable type to the interface struct. @@ -2582,26 +2614,40 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass ib_type.member_types.push_back(usable_type->self); // Give the member a name - string mbr_name = ensure_valid_name(join(to_qualified_member_name(var_type, mbr_idx), "_", i), "m"); + string mbr_name = ensure_valid_name(append_member_name(mbr_name_qual, var_type, mbr_idx) + (elem_cnt == 1 ? "" : join("_", i)), "m"); set_member_name(ib_type.self, ib_mbr_idx, mbr_name); - if (has_member_decoration(var_type.self, mbr_idx, DecorationLocation)) + // Once we determine the location of the first member within nested structures, + // from a var of the topmost structure, the remaining flattened members of + // the nested structures will have consecutive location values. At this point, + // we've recursively tunnelled into structs, arrays, and matrices, and are + // down to a single location for each member now. + if (!is_builtin && location != UINT32_MAX) { - uint32_t locn = get_member_decoration(var_type.self, mbr_idx, DecorationLocation) + i; - set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn); - mark_location_as_used_by_shader(locn, *usable_type, storage); + set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location); + mark_location_as_used_by_shader(location, *usable_type, storage); + location++; + } + else if (has_member_decoration(var_type.self, mbr_idx, DecorationLocation)) + { + location = get_member_decoration(var_type.self, mbr_idx, DecorationLocation) + i; + set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location); + mark_location_as_used_by_shader(location, *usable_type, storage); + location++; } else if (has_decoration(var.self, DecorationLocation)) { - uint32_t locn = get_accumulated_member_location(var, mbr_idx, meta.strip_array) + i; - set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn); - mark_location_as_used_by_shader(locn, *usable_type, storage); + location = get_accumulated_member_location(var, mbr_idx, meta.strip_array) + i; + set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location); + mark_location_as_used_by_shader(location, *usable_type, storage); + location++; } else if (is_builtin && is_tessellation_shader() && storage == StorageClassInput && inputs_by_builtin.count(builtin)) { - uint32_t locn = inputs_by_builtin[builtin].location + i; - set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn); - mark_location_as_used_by_shader(locn, *usable_type, storage); + location = inputs_by_builtin[builtin].location + i; + set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location); + mark_location_as_used_by_shader(location, *usable_type, storage); + location++; } else if (is_builtin && (builtin == BuiltInClipDistance || builtin == BuiltInCullDistance)) { @@ -2611,7 +2657,7 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass } if (has_member_decoration(var_type.self, mbr_idx, DecorationComponent)) - SPIRV_CROSS_THROW("DecorationComponent on matrices and arrays make little sense."); + SPIRV_CROSS_THROW("DecorationComponent on matrices and arrays is not supported."); if (storage != StorageClassInput || !pull_model_inputs.count(var.self)) { @@ -2627,47 +2673,36 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass } set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceOrigID, var.self); - set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceMemberIndex, mbr_idx); + set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceMemberIndex, var_mbr_idx); // Unflatten or flatten from [[stage_in]] or [[stage_out]] as appropriate. if (!meta.strip_array && meta.allow_local_declaration) { + string var_chain = join(var_chain_qual, ".", to_member_name(var_type, mbr_idx), (elem_cnt == 1 ? "" : join("[", i, "]"))); switch (storage) { case StorageClassInput: - entry_func.fixup_hooks_in.push_back([=, &var, &var_type]() { + entry_func.fixup_hooks_in.push_back([=, &var]() { + string lerp_call; if (pull_model_inputs.count(var.self)) { - string lerp_call; if (is_centroid) lerp_call = ".interpolate_at_centroid()"; else if (is_sample) lerp_call = join(".interpolate_at_sample(", to_expression(builtin_sample_id_id), ")"); else lerp_call = ".interpolate_at_center()"; - statement(to_name(var.self), ".", to_member_name(var_type, mbr_idx), "[", i, "] = ", ib_var_ref, - ".", mbr_name, lerp_call, ";"); - } - else - { - statement(to_name(var.self), ".", to_member_name(var_type, mbr_idx), "[", i, "] = ", ib_var_ref, - ".", mbr_name, ";"); } + statement(var_chain, " = ", ib_var_ref, ".", mbr_name, lerp_call, ";"); }); break; case StorageClassOutput: - entry_func.fixup_hooks_out.push_back([=, &var, &var_type]() { + entry_func.fixup_hooks_out.push_back([=]() { if (flatten_from_ib_var) - { - statement(ib_var_ref, ".", mbr_name, " = ", ib_var_ref, ".", flatten_from_ib_mbr_name, "[", i, - "];"); - } + statement(ib_var_ref, ".", mbr_name, " = ", ib_var_ref, ".", flatten_from_ib_mbr_name, "[", i, "];"); else - { - statement(ib_var_ref, ".", mbr_name, " = ", to_name(var.self), ".", - to_member_name(var_type, mbr_idx), "[", i, "];"); - } + statement(ib_var_ref, ".", mbr_name, " = ", var_chain, ";"); }); break; @@ -2678,11 +2713,14 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass } } -void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass storage, const string &ib_var_ref, - SPIRType &ib_type, SPIRVariable &var, uint32_t mbr_idx, - InterfaceBlockMeta &meta) +void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass storage, + const string &ib_var_ref, SPIRType &ib_type, + SPIRVariable &var, SPIRType &var_type, + uint32_t mbr_idx, InterfaceBlockMeta &meta, + const string &mbr_name_qual, + const string &var_chain_qual, + uint32_t &location, uint32_t &var_mbr_idx) { - auto &var_type = meta.strip_array ? get_variable_element_type(var) : get_variable_data_type(var); auto &entry_func = get(ir.default_entry_point); BuiltIn builtin = BuiltInMax; @@ -2707,7 +2745,7 @@ void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass stor ib_type.member_types.push_back(mbr_type_id); // Give the member a name - string mbr_name = ensure_valid_name(to_qualified_member_name(var_type, mbr_idx), "m"); + string mbr_name = ensure_valid_name(append_member_name(mbr_name_qual, var_type, mbr_idx), "m"); set_member_name(ib_type.self, ib_mbr_idx, mbr_name); // Update the original variable reference to include the structure reference @@ -2724,7 +2762,7 @@ void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass stor } bool flatten_stage_out = false; - + string var_chain = var_chain_qual + "." + to_member_name(var_type, mbr_idx); if (is_builtin && !meta.strip_array) { // For the builtin gl_PerVertex, we cannot treat it as a block anyways, @@ -2737,15 +2775,15 @@ void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass stor switch (storage) { case StorageClassInput: - entry_func.fixup_hooks_in.push_back([=, &var, &var_type]() { - statement(to_name(var.self), ".", to_member_name(var_type, mbr_idx), " = ", qual_var_name, ";"); + entry_func.fixup_hooks_in.push_back([=]() { + statement(var_chain, " = ", qual_var_name, ";"); }); break; case StorageClassOutput: flatten_stage_out = true; - entry_func.fixup_hooks_out.push_back([=, &var, &var_type]() { - statement(qual_var_name, " = ", to_name(var.self), ".", to_member_name(var_type, mbr_idx), ";"); + entry_func.fixup_hooks_out.push_back([=]() { + statement(qual_var_name, " = ", var_chain, ";"); }); break; @@ -2754,48 +2792,56 @@ void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass stor } } - // Copy the variable location from the original variable to the member - if (has_member_decoration(var_type.self, mbr_idx, DecorationLocation)) + // Once we determine the location of the first member within nested structures, + // from a var of the topmost structure, the remaining flattened members of + // the nested structures will have consecutive location values. At this point, + // we've recursively tunnelled into structs, arrays, and matrices, and are + // down to a single location for each member now. + if (!is_builtin && location != UINT32_MAX) { - uint32_t locn = get_member_decoration(var_type.self, mbr_idx, DecorationLocation); + set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location); + mark_location_as_used_by_shader(location, get(mbr_type_id), storage); + location++; + } + else if (has_member_decoration(var_type.self, mbr_idx, DecorationLocation)) + { + location = get_member_decoration(var_type.self, mbr_idx, DecorationLocation); uint32_t comp = get_member_decoration(var_type.self, mbr_idx, DecorationComponent); if (storage == StorageClassInput) { - mbr_type_id = ensure_correct_input_type(mbr_type_id, locn, comp, 0, meta.strip_array); + mbr_type_id = ensure_correct_input_type(mbr_type_id, location, comp, 0, meta.strip_array); var_type.member_types[mbr_idx] = mbr_type_id; if (storage == StorageClassInput && pull_model_inputs.count(var.self)) ib_type.member_types[ib_mbr_idx] = build_msl_interpolant_type(mbr_type_id, is_noperspective); else ib_type.member_types[ib_mbr_idx] = mbr_type_id; } - set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn); - mark_location_as_used_by_shader(locn, get(mbr_type_id), storage); + set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location); + mark_location_as_used_by_shader(location, get(mbr_type_id), storage); + location++; } else if (has_decoration(var.self, DecorationLocation)) { - // The block itself might have a location and in this case, all members of the block - // receive incrementing locations. - uint32_t locn = get_accumulated_member_location(var, mbr_idx, meta.strip_array); + location = get_accumulated_member_location(var, mbr_idx, meta.strip_array); if (storage == StorageClassInput) { - mbr_type_id = ensure_correct_input_type(mbr_type_id, locn, 0, 0, meta.strip_array); + mbr_type_id = ensure_correct_input_type(mbr_type_id, location, 0, 0, meta.strip_array); var_type.member_types[mbr_idx] = mbr_type_id; if (storage == StorageClassInput && pull_model_inputs.count(var.self)) ib_type.member_types[ib_mbr_idx] = build_msl_interpolant_type(mbr_type_id, is_noperspective); else ib_type.member_types[ib_mbr_idx] = mbr_type_id; } - set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn); - mark_location_as_used_by_shader(locn, get(mbr_type_id), storage); + set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location); + mark_location_as_used_by_shader(location, get(mbr_type_id), storage); + location++; } else if (is_builtin && is_tessellation_shader() && storage == StorageClassInput && inputs_by_builtin.count(builtin)) { - uint32_t locn = 0; - auto builtin_itr = inputs_by_builtin.find(builtin); - if (builtin_itr != end(inputs_by_builtin)) - locn = builtin_itr->second.location; - set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn); - mark_location_as_used_by_shader(locn, get(mbr_type_id), storage); + location = inputs_by_builtin[builtin].location; + set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, location); + mark_location_as_used_by_shader(location, get(mbr_type_id), storage); + location++; } // Copy the component location, if present. @@ -2854,7 +2900,7 @@ void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass stor } set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceOrigID, var.self); - set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceMemberIndex, mbr_idx); + set_extended_member_decoration(ib_type.self, ib_mbr_idx, SPIRVCrossDecorationInterfaceMemberIndex, var_mbr_idx); } // In Metal, the tessellation levels are stored as tightly packed half-precision floating point values. @@ -3114,67 +3160,102 @@ void CompilerMSL::add_variable_to_interface_block(StorageClass storage, const st else { bool masked_block = false; - - // Flatten the struct members into the interface struct - for (uint32_t mbr_idx = 0; mbr_idx < uint32_t(var_type.member_types.size()); mbr_idx++) + uint32_t location = UINT32_MAX; + uint32_t var_mbr_idx = 0; + uint32_t elem_cnt = 1; + if (is_matrix(var_type)) { - builtin = BuiltInMax; - is_builtin = is_member_builtin(var_type, mbr_idx, &builtin); - auto &mbr_type = get(var_type.member_types[mbr_idx]); + if (is_array(var_type)) + SPIRV_CROSS_THROW("MSL cannot emit arrays-of-matrices in input and output variables."); - if (storage == StorageClassOutput && is_stage_output_block_member_masked(var, mbr_idx, meta.strip_array)) + elem_cnt = var_type.columns; + } + else if (is_array(var_type)) + { + if (var_type.array.size() != 1) + SPIRV_CROSS_THROW("MSL cannot emit arrays-of-arrays in input and output variables."); + + elem_cnt = to_array_size_literal(var_type); + } + + for (uint32_t elem_idx = 0; elem_idx < elem_cnt; elem_idx++) + { + // Flatten the struct members into the interface struct + for (uint32_t mbr_idx = 0; mbr_idx < uint32_t(var_type.member_types.size()); mbr_idx++) { - if (is_block) - masked_block = true; + builtin = BuiltInMax; + is_builtin = is_member_builtin(var_type, mbr_idx, &builtin); + auto &mbr_type = get(var_type.member_types[mbr_idx]); - // Non-builtin block output variables are just ignored, since they will still access - // the block variable as-is. They're just not flattened. - if (is_builtin && !meta.strip_array) + if (storage == StorageClassOutput && is_stage_output_block_member_masked(var, mbr_idx, meta.strip_array)) { - // Emit a fake variable instead. - uint32_t ids = ir.increase_bound_by(2); - uint32_t ptr_type_id = ids + 0; - uint32_t var_id = ids + 1; + location = UINT32_MAX; // Skip this member and resolve location again on next var member - auto ptr_type = mbr_type; - ptr_type.pointer = true; - ptr_type.pointer_depth++; - ptr_type.parent_type = var_type.member_types[mbr_idx]; - ptr_type.storage = StorageClassOutput; + if (is_block) + masked_block = true; - uint32_t initializer = 0; - if (var.initializer) - if (auto *c = maybe_get(var.initializer)) - initializer = c->subconstants[mbr_idx]; + // Non-builtin block output variables are just ignored, since they will still access + // the block variable as-is. They're just not flattened. + if (is_builtin && !meta.strip_array) + { + // Emit a fake variable instead. + uint32_t ids = ir.increase_bound_by(2); + uint32_t ptr_type_id = ids + 0; + uint32_t var_id = ids + 1; - set(ptr_type_id, ptr_type); - set(var_id, ptr_type_id, StorageClassOutput, initializer); - entry_func.add_local_variable(var_id); - vars_needing_early_declaration.push_back(var_id); - set_name(var_id, builtin_to_glsl(builtin, StorageClassOutput)); - set_decoration(var_id, DecorationBuiltIn, builtin); + auto ptr_type = mbr_type; + ptr_type.pointer = true; + ptr_type.pointer_depth++; + ptr_type.parent_type = var_type.member_types[mbr_idx]; + ptr_type.storage = StorageClassOutput; + + uint32_t initializer = 0; + if (var.initializer) + if (auto *c = maybe_get(var.initializer)) + initializer = c->subconstants[mbr_idx]; + + set(ptr_type_id, ptr_type); + set(var_id, ptr_type_id, StorageClassOutput, initializer); + entry_func.add_local_variable(var_id); + vars_needing_early_declaration.push_back(var_id); + set_name(var_id, builtin_to_glsl(builtin, StorageClassOutput)); + set_decoration(var_id, DecorationBuiltIn, builtin); + } } - } - else if (!is_builtin || has_active_builtin(builtin, storage)) - { - bool is_composite_type = is_matrix(mbr_type) || is_array(mbr_type); - bool attribute_load_store = - storage == StorageClassInput && get_execution_model() != ExecutionModelFragment; - bool storage_is_stage_io = variable_storage_requires_stage_io(storage); - - // Clip/CullDistance always need to be declared as user attributes. - if (builtin == BuiltInClipDistance || builtin == BuiltInCullDistance) - is_builtin = false; - - if ((!is_builtin || attribute_load_store) && storage_is_stage_io && is_composite_type) + else if (!is_builtin || has_active_builtin(builtin, storage)) { - add_composite_member_variable_to_interface_block(storage, ib_var_ref, ib_type, var, mbr_idx, - meta); - } - else - { - add_plain_member_variable_to_interface_block(storage, ib_var_ref, ib_type, var, mbr_idx, meta); + bool is_composite_type = is_matrix(mbr_type) || is_array(mbr_type) || mbr_type.basetype == SPIRType::Struct; + bool attribute_load_store = + storage == StorageClassInput && get_execution_model() != ExecutionModelFragment; + bool storage_is_stage_io = variable_storage_requires_stage_io(storage); + + // Clip/CullDistance always need to be declared as user attributes. + if (builtin == BuiltInClipDistance || builtin == BuiltInCullDistance) + is_builtin = false; + + string mbr_name_qual = to_name(var_type.self); + string var_chain_qual = to_name(var.self); + if (elem_cnt > 1) { + mbr_name_qual += join("_", elem_idx); + var_chain_qual += join("[", elem_idx, "]"); + } + + if ((!is_builtin || attribute_load_store) && storage_is_stage_io && is_composite_type) + { + add_composite_member_variable_to_interface_block(storage, ib_var_ref, ib_type, + var, var_type, mbr_idx, meta, + mbr_name_qual, var_chain_qual, + location, var_mbr_idx); + } + else + { + add_plain_member_variable_to_interface_block(storage, ib_var_ref, ib_type, + var, var_type, mbr_idx, meta, + mbr_name_qual, var_chain_qual, + location, var_mbr_idx); + } } + var_mbr_idx++; } } @@ -5370,7 +5451,7 @@ void CompilerMSL::emit_custom_functions() statement("template"); statement("inline T spvSubgroupBroadcast(T value, ushort lane)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return quad_broadcast(value, lane);"); else statement("return simd_broadcast(value, lane);"); @@ -5379,7 +5460,7 @@ void CompilerMSL::emit_custom_functions() statement("template<>"); statement("inline bool spvSubgroupBroadcast(bool value, ushort lane)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return !!quad_broadcast((ushort)value, lane);"); else statement("return !!simd_broadcast((ushort)value, lane);"); @@ -5388,7 +5469,7 @@ void CompilerMSL::emit_custom_functions() statement("template"); statement("inline vec spvSubgroupBroadcast(vec value, ushort lane)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return (vec)quad_broadcast((vec)value, lane);"); else statement("return (vec)simd_broadcast((vec)value, lane);"); @@ -5400,7 +5481,7 @@ void CompilerMSL::emit_custom_functions() statement("template"); statement("inline T spvSubgroupBroadcastFirst(T value)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return quad_broadcast_first(value);"); else statement("return simd_broadcast_first(value);"); @@ -5409,7 +5490,7 @@ void CompilerMSL::emit_custom_functions() statement("template<>"); statement("inline bool spvSubgroupBroadcastFirst(bool value)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return !!quad_broadcast_first((ushort)value);"); else statement("return !!simd_broadcast_first((ushort)value);"); @@ -5418,7 +5499,7 @@ void CompilerMSL::emit_custom_functions() statement("template"); statement("inline vec spvSubgroupBroadcastFirst(vec value)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return (vec)quad_broadcast_first((vec)value);"); else statement("return (vec)simd_broadcast_first((vec)value);"); @@ -5429,7 +5510,7 @@ void CompilerMSL::emit_custom_functions() case SPVFuncImplSubgroupBallot: statement("inline uint4 spvSubgroupBallot(bool value)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) { statement("return uint4((quad_vote::vote_t)quad_ballot(value), 0, 0, 0);"); } @@ -5557,7 +5638,7 @@ void CompilerMSL::emit_custom_functions() statement("template"); statement("inline bool spvSubgroupAllEqual(T value)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return quad_all(all(value == quad_broadcast_first(value)));"); else statement("return simd_all(all(value == simd_broadcast_first(value)));"); @@ -5566,7 +5647,7 @@ void CompilerMSL::emit_custom_functions() statement("template<>"); statement("inline bool spvSubgroupAllEqual(bool value)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return quad_all(value) || !quad_any(value);"); else statement("return simd_all(value) || !simd_any(value);"); @@ -5575,7 +5656,7 @@ void CompilerMSL::emit_custom_functions() statement("template"); statement("inline bool spvSubgroupAllEqual(vec value)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return quad_all(all(value == (vec)quad_broadcast_first((vec)value)));"); else statement("return simd_all(all(value == (vec)simd_broadcast_first((vec)value)));"); @@ -5587,7 +5668,7 @@ void CompilerMSL::emit_custom_functions() statement("template"); statement("inline T spvSubgroupShuffle(T value, ushort lane)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return quad_shuffle(value, lane);"); else statement("return simd_shuffle(value, lane);"); @@ -5596,7 +5677,7 @@ void CompilerMSL::emit_custom_functions() statement("template<>"); statement("inline bool spvSubgroupShuffle(bool value, ushort lane)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return !!quad_shuffle((ushort)value, lane);"); else statement("return !!simd_shuffle((ushort)value, lane);"); @@ -5605,7 +5686,7 @@ void CompilerMSL::emit_custom_functions() statement("template"); statement("inline vec spvSubgroupShuffle(vec value, ushort lane)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return (vec)quad_shuffle((vec)value, lane);"); else statement("return (vec)simd_shuffle((vec)value, lane);"); @@ -5617,7 +5698,7 @@ void CompilerMSL::emit_custom_functions() statement("template"); statement("inline T spvSubgroupShuffleXor(T value, ushort mask)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return quad_shuffle_xor(value, mask);"); else statement("return simd_shuffle_xor(value, mask);"); @@ -5626,7 +5707,7 @@ void CompilerMSL::emit_custom_functions() statement("template<>"); statement("inline bool spvSubgroupShuffleXor(bool value, ushort mask)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return !!quad_shuffle_xor((ushort)value, mask);"); else statement("return !!simd_shuffle_xor((ushort)value, mask);"); @@ -5635,7 +5716,7 @@ void CompilerMSL::emit_custom_functions() statement("template"); statement("inline vec spvSubgroupShuffleXor(vec value, ushort mask)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return (vec)quad_shuffle_xor((vec)value, mask);"); else statement("return (vec)simd_shuffle_xor((vec)value, mask);"); @@ -5647,7 +5728,7 @@ void CompilerMSL::emit_custom_functions() statement("template"); statement("inline T spvSubgroupShuffleUp(T value, ushort delta)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return quad_shuffle_up(value, delta);"); else statement("return simd_shuffle_up(value, delta);"); @@ -5656,7 +5737,7 @@ void CompilerMSL::emit_custom_functions() statement("template<>"); statement("inline bool spvSubgroupShuffleUp(bool value, ushort delta)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return !!quad_shuffle_up((ushort)value, delta);"); else statement("return !!simd_shuffle_up((ushort)value, delta);"); @@ -5665,7 +5746,7 @@ void CompilerMSL::emit_custom_functions() statement("template"); statement("inline vec spvSubgroupShuffleUp(vec value, ushort delta)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return (vec)quad_shuffle_up((vec)value, delta);"); else statement("return (vec)simd_shuffle_up((vec)value, delta);"); @@ -5677,7 +5758,7 @@ void CompilerMSL::emit_custom_functions() statement("template"); statement("inline T spvSubgroupShuffleDown(T value, ushort delta)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return quad_shuffle_down(value, delta);"); else statement("return simd_shuffle_down(value, delta);"); @@ -5686,7 +5767,7 @@ void CompilerMSL::emit_custom_functions() statement("template<>"); statement("inline bool spvSubgroupShuffleDown(bool value, ushort delta)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return !!quad_shuffle_down((ushort)value, delta);"); else statement("return !!simd_shuffle_down((ushort)value, delta);"); @@ -5695,7 +5776,7 @@ void CompilerMSL::emit_custom_functions() statement("template"); statement("inline vec spvSubgroupShuffleDown(vec value, ushort delta)"); begin_scope(); - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) statement("return (vec)quad_shuffle_down((vec)value, delta);"); else statement("return (vec)simd_shuffle_down((vec)value, delta);"); @@ -7368,7 +7449,7 @@ bool CompilerMSL::emit_tessellation_access_chain(const uint32_t *ops, uint32_t l expr_type->vecsize > result_ptr_type.vecsize) e += vector_swizzle(result_ptr_type.vecsize, 0); - auto &expr = set(ops[1], move(e), ops[0], should_forward(ops[2])); + auto &expr = set(ops[1], std::move(e), ops[0], should_forward(ops[2])); expr.loaded_from = var->self; expr.need_transpose = meta.need_transpose; expr.access_chain = true; @@ -7559,6 +7640,8 @@ void CompilerMSL::emit_instruction(const Instruction &instruction) auto ops = stream(instruction); auto opcode = static_cast(instruction.op); + opcode = get_remapped_spirv_op(opcode); + // If we need to do implicit bitcasts, make sure we do it with the correct type. uint32_t integer_width = get_integer_width_for_instruction(instruction); auto int_type = to_signed_basetype(integer_width); @@ -7601,6 +7684,10 @@ void CompilerMSL::emit_instruction(const Instruction &instruction) case OpLogicalNotEqual: case OpFOrdNotEqual: + // TODO: Should probably negate the == result here. + // Typically OrdNotEqual comes from GLSL which itself does not really specify what + // happens with NaN. + // Consider fixing this if we run into real issues. MSL_BOP(!=); break; @@ -7657,7 +7744,9 @@ void CompilerMSL::emit_instruction(const Instruction &instruction) break; case OpFUnordNotEqual: - MSL_UNORD_BOP(!=); + // not equal in MSL generates une opcodes to begin with. + // Since unordered not equal is how it works in C, just inherit that behavior. + MSL_BOP(!=); break; case OpFUnordGreaterThan: @@ -8920,6 +9009,8 @@ void CompilerMSL::emit_glsl_op(uint32_t result_type, uint32_t id, uint32_t eop, auto int_type = to_signed_basetype(integer_width); auto uint_type = to_unsigned_basetype(integer_width); + op = get_remapped_glsl_op(op); + switch (op) { case GLSLstd450Sinh: @@ -9335,6 +9426,9 @@ void CompilerMSL::emit_function_prototype(SPIRFunction &func, const Bitset &) else decl += entry_point_args_classic(!func.arguments.empty()); + // append entry point args to avoid conflicts in local variable names. + local_variable_names.insert(resource_names.begin(), resource_names.end()); + // If entry point function has variables that require early declaration, // ensure they each have an empty initializer, creating one if needed. // This is done at this late stage because the initialization expression @@ -10567,7 +10661,7 @@ string CompilerMSL::convert_row_major_matrix(string exp_str, const SPIRType &exp { if (!is_matrix(exp_type)) { - return CompilerGLSL::convert_row_major_matrix(move(exp_str), exp_type, physical_type_id, is_packed); + return CompilerGLSL::convert_row_major_matrix(std::move(exp_str), exp_type, physical_type_id, is_packed); } else { @@ -12077,16 +12171,6 @@ void CompilerMSL::fix_up_shader_inputs_outputs() }); } break; - case BuiltInHelperInvocation: - if (msl_options.is_ios() && !msl_options.supports_msl_version(2, 3)) - SPIRV_CROSS_THROW("simd_is_helper_thread() requires version 2.3 on iOS."); - else if (msl_options.is_macos() && !msl_options.supports_msl_version(2, 1)) - SPIRV_CROSS_THROW("simd_is_helper_thread() requires version 2.1 on macOS."); - - entry_func.fixup_hooks_in.push_back([=]() { - statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = simd_is_helper_thread();"); - }); - break; case BuiltInInvocationId: // This is direct-mapped without multi-patch workgroups. if (get_execution_model() != ExecutionModelTessellationControl || !msl_options.multi_patch_workgroup) @@ -12979,8 +13063,8 @@ string CompilerMSL::to_name(uint32_t id, bool allow_alias) const return Compiler::to_name(id, allow_alias); } -// Returns a name that combines the name of the struct with the name of the member, except for Builtins -string CompilerMSL::to_qualified_member_name(const SPIRType &type, uint32_t index) +// Appends the name of the member to the variable qualifier string, except for Builtins. +string CompilerMSL::append_member_name(const string &qualifier, const SPIRType &type, uint32_t index) { // Don't qualify Builtin names because they are unique and are treated as such when building expressions BuiltIn builtin = BuiltInMax; @@ -12991,7 +13075,7 @@ string CompilerMSL::to_qualified_member_name(const SPIRType &type, uint32_t inde string mbr_name = to_member_name(type, index); size_t startPos = mbr_name.find_first_not_of("_"); mbr_name = (startPos != string::npos) ? mbr_name.substr(startPos) : ""; - return join(to_name(type.self), "_", mbr_name); + return join(qualifier, "_", mbr_name); } // Ensures that the specified name is permanently usable by prepending a prefix @@ -13969,7 +14053,7 @@ void CompilerMSL::emit_subgroup_op(const Instruction &i) switch (op) { case OpGroupNonUniformElect: - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) emit_op(result_type, id, "quad_is_first()", false); else emit_op(result_type, id, "simd_is_first()", false); @@ -14042,14 +14126,14 @@ void CompilerMSL::emit_subgroup_op(const Instruction &i) break; case OpGroupNonUniformAll: - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) emit_unary_func_op(result_type, id, ops[3], "quad_all"); else emit_unary_func_op(result_type, id, ops[3], "simd_all"); break; case OpGroupNonUniformAny: - if (msl_options.is_ios() && !msl_options.ios_use_simdgroup_functions) + if (msl_options.use_quadgroup_operation()) emit_unary_func_op(result_type, id, ops[3], "quad_any"); else emit_unary_func_op(result_type, id, ops[3], "simd_any"); @@ -14381,6 +14465,14 @@ string CompilerMSL::builtin_to_glsl(BuiltIn builtin, StorageClass storage) } break; + case BuiltInHelperInvocation: + if (msl_options.is_ios() && !msl_options.supports_msl_version(2, 3)) + SPIRV_CROSS_THROW("simd_is_helper_thread() requires version 2.3 on iOS."); + else if (msl_options.is_macos() && !msl_options.supports_msl_version(2, 1)) + SPIRV_CROSS_THROW("simd_is_helper_thread() requires version 2.1 on macOS."); + // In SPIR-V 1.6 with Volatile HelperInvocation, we cannot emit a fixup early. + return "simd_is_helper_thread()"; + default: break; } @@ -14547,7 +14639,7 @@ string CompilerMSL::builtin_qualifier(BuiltIn builtin) SPIRV_CROSS_THROW("NumSubgroups is handled specially with emulation."); if (!msl_options.supports_msl_version(2)) SPIRV_CROSS_THROW("Subgroup builtins require Metal 2.0."); - return msl_options.is_ios() ? "quadgroups_per_threadgroup" : "simdgroups_per_threadgroup"; + return msl_options.use_quadgroup_operation() ? "quadgroups_per_threadgroup" : "simdgroups_per_threadgroup"; case BuiltInSubgroupId: if (msl_options.emulate_subgroups) @@ -14555,7 +14647,7 @@ string CompilerMSL::builtin_qualifier(BuiltIn builtin) SPIRV_CROSS_THROW("SubgroupId is handled specially with emulation."); if (!msl_options.supports_msl_version(2)) SPIRV_CROSS_THROW("Subgroup builtins require Metal 2.0."); - return msl_options.is_ios() ? "quadgroup_index_in_threadgroup" : "simdgroup_index_in_threadgroup"; + return msl_options.use_quadgroup_operation() ? "quadgroup_index_in_threadgroup" : "simdgroup_index_in_threadgroup"; case BuiltInSubgroupLocalInvocationId: if (msl_options.emulate_subgroups) @@ -14574,7 +14666,7 @@ string CompilerMSL::builtin_qualifier(BuiltIn builtin) // We are generating a Metal kernel function. if (!msl_options.supports_msl_version(2)) SPIRV_CROSS_THROW("Subgroup builtins in kernel functions require Metal 2.0."); - return msl_options.is_ios() ? "thread_index_in_quadgroup" : "thread_index_in_simdgroup"; + return msl_options.use_quadgroup_operation() ? "thread_index_in_quadgroup" : "thread_index_in_simdgroup"; } else SPIRV_CROSS_THROW("Subgroup builtins are not available in this type of function."); diff --git a/3rdparty/spirv-cross/spirv_msl.hpp b/3rdparty/spirv-cross/spirv_msl.hpp index 6591e47c5..b3b706ec7 100644 --- a/3rdparty/spirv-cross/spirv_msl.hpp +++ b/3rdparty/spirv-cross/spirv_msl.hpp @@ -393,7 +393,7 @@ public: // and will be addressed using the current ViewIndex. bool arrayed_subpass_input = false; - // Whether to use SIMD-group or quadgroup functions to implement group nnon-uniform + // Whether to use SIMD-group or quadgroup functions to implement group non-uniform // operations. Some GPUs on iOS do not support the SIMD-group functions, only the // quadgroup functions. bool ios_use_simdgroup_functions = false; @@ -445,6 +445,11 @@ public: return platform == macOS; } + bool use_quadgroup_operation() const + { + return is_ios() && !ios_use_simdgroup_functions; + } + void set_msl_version(uint32_t major, uint32_t minor = 0, uint32_t patch = 0) { msl_version = make_msl_version(major, minor, patch); @@ -826,12 +831,20 @@ protected: bool add_component_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref, SPIRVariable &var, const SPIRType &type, InterfaceBlockMeta &meta); - void add_plain_member_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref, - SPIRType &ib_type, SPIRVariable &var, uint32_t index, - InterfaceBlockMeta &meta); - void add_composite_member_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref, - SPIRType &ib_type, SPIRVariable &var, uint32_t index, - InterfaceBlockMeta &meta); + void add_plain_member_variable_to_interface_block(spv::StorageClass storage, + const std::string &ib_var_ref, SPIRType &ib_type, + SPIRVariable &var, SPIRType &var_type, + uint32_t mbr_idx, InterfaceBlockMeta &meta, + const std::string &mbr_name_qual, + const std::string &var_chain_qual, + uint32_t &location, uint32_t &var_mbr_idx); + void add_composite_member_variable_to_interface_block(spv::StorageClass storage, + const std::string &ib_var_ref, SPIRType &ib_type, + SPIRVariable &var, SPIRType &var_type, + uint32_t mbr_idx, InterfaceBlockMeta &meta, + const std::string &mbr_name_qual, + const std::string &var_chain_qual, + uint32_t &location, uint32_t &var_mbr_idx); void add_tess_level_input_to_interface_block(const std::string &ib_var_ref, SPIRType &ib_type, SPIRVariable &var); void fix_up_interface_member_indices(spv::StorageClass storage, uint32_t ib_type_id); @@ -858,7 +871,7 @@ protected: std::string entry_point_arg_stage_in(); void entry_point_args_builtin(std::string &args); void entry_point_args_discrete_descriptors(std::string &args); - std::string to_qualified_member_name(const SPIRType &type, uint32_t index); + std::string append_member_name(const std::string &qualifier, const SPIRType &type, uint32_t index); std::string ensure_valid_name(std::string name, std::string pfx); std::string to_sampler_expression(uint32_t id); std::string to_swizzle_expression(uint32_t id); diff --git a/3rdparty/spirv-cross/spirv_parser.cpp b/3rdparty/spirv-cross/spirv_parser.cpp index 55700fdbb..1296f841f 100644 --- a/3rdparty/spirv-cross/spirv_parser.cpp +++ b/3rdparty/spirv-cross/spirv_parser.cpp @@ -31,7 +31,7 @@ namespace SPIRV_CROSS_NAMESPACE { Parser::Parser(vector spirv) { - ir.spirv = move(spirv); + ir.spirv = std::move(spirv); } Parser::Parser(const uint32_t *spirv_data, size_t word_count) @@ -259,7 +259,7 @@ void Parser::parse(const Instruction &instruction) case OpExtension: { auto ext = extract_string(ir.spirv, instruction.offset); - ir.declared_extensions.push_back(move(ext)); + ir.declared_extensions.push_back(std::move(ext)); break; } @@ -291,7 +291,15 @@ void Parser::parse(const Instruction &instruction) { // The SPIR-V debug information extended instructions might come at global scope. if (current_block) + { current_block->ops.push_back(instruction); + if (length >= 2) + { + const auto *type = maybe_get(ops[0]); + if (type) + ir.load_type_width.insert({ ops[1], type->width }); + } + } break; } @@ -1059,6 +1067,7 @@ void Parser::parse(const Instruction &instruction) } case OpKill: + case OpTerminateInvocation: { if (!current_block) SPIRV_CROSS_THROW("Trying to end a non-existing block."); @@ -1211,10 +1220,9 @@ void Parser::parse(const Instruction &instruction) { const auto *type = maybe_get(ops[0]); if (type) - { ir.load_type_width.insert({ ops[1], type->width }); - } } + if (!current_block) SPIRV_CROSS_THROW("Currently no block to insert opcode.");