From b61c3cc15ad64e13743b4c6afc5759eef769566c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=D0=91=D1=80=D0=B0=D0=BD=D0=B8=D0=BC=D0=B8=D1=80=20=D0=9A?= =?UTF-8?q?=D0=B0=D1=80=D0=B0=D1=9F=D0=B8=D1=9B?= Date: Sun, 2 Oct 2022 22:57:10 -0700 Subject: [PATCH] Updated spirv-cross. --- 3rdparty/spirv-cross/main.cpp | 43 +- 3rdparty/spirv-cross/spirv.h | 23 + 3rdparty/spirv-cross/spirv.hpp | 23 + 3rdparty/spirv-cross/spirv_common.hpp | 17 +- 3rdparty/spirv-cross/spirv_cross.cpp | 36 +- 3rdparty/spirv-cross/spirv_cross.hpp | 3 + 3rdparty/spirv-cross/spirv_cross_c.cpp | 73 +- 3rdparty/spirv-cross/spirv_cross_c.h | 41 +- .../spirv-cross/spirv_cross_parsed_ir.cpp | 4 + 3rdparty/spirv-cross/spirv_glsl.cpp | 198 ++++- 3rdparty/spirv-cross/spirv_glsl.hpp | 10 +- 3rdparty/spirv-cross/spirv_msl.cpp | 784 ++++++++++++++++-- 3rdparty/spirv-cross/spirv_msl.hpp | 74 +- 3rdparty/spirv-cross/spirv_parser.cpp | 25 + 3rdparty/spirv-cross/spirv_parser.hpp | 2 + 15 files changed, 1186 insertions(+), 170 deletions(-) diff --git a/3rdparty/spirv-cross/main.cpp b/3rdparty/spirv-cross/main.cpp index 81db89ce0..e83284a03 100644 --- a/3rdparty/spirv-cross/main.cpp +++ b/3rdparty/spirv-cross/main.cpp @@ -681,7 +681,8 @@ struct CLIArguments SmallVector msl_device_argument_buffers; SmallVector> msl_dynamic_buffers; SmallVector> msl_inline_uniform_blocks; - SmallVector msl_shader_inputs; + SmallVector msl_shader_inputs; + SmallVector msl_shader_outputs; SmallVector pls_in; SmallVector pls_out; SmallVector remaps; @@ -874,6 +875,10 @@ static void print_help_msl() "\t\t can be 'any32', 'any16', 'u16', 'u8', or 'other', to indicate a 32-bit opaque value, 16-bit opaque value, 16-bit unsigned integer, 8-bit unsigned integer, " "or other-typed variable. is the vector length of the variable, which must be greater than or equal to that declared in the shader.\n" "\t\tUseful if shader stage interfaces don't match up, as pipeline creation might otherwise fail.\n" + "\t[--msl-shader-output ]:\n\t\tSpecify the format of the shader output at .\n" + "\t\t can be 'any32', 'any16', 'u16', 'u8', or 'other', to indicate a 32-bit opaque value, 16-bit opaque value, 16-bit unsigned integer, 8-bit unsigned integer, " + "or other-typed variable. is the vector length of the variable, which must be greater than or equal to that declared in the shader.\n" + "\t\tUseful if shader stage interfaces don't match up, as pipeline creation might otherwise fail.\n" "\t[--msl-multi-patch-workgroup]:\n\t\tUse the new style of tessellation control processing, where multiple patches are processed per workgroup.\n" "\t\tThis should increase throughput by ensuring all the GPU's SIMD lanes are occupied, but it is not compatible with the old style.\n" "\t\tIn addition, this style also passes input variables in buffers directly instead of using vertex attribute processing.\n" @@ -1082,6 +1087,10 @@ static ExecutionModel stage_to_execution_model(const std::string &stage) return ExecutionModelMissKHR; else if (stage == "rcall") return ExecutionModelCallableKHR; + else if (stage == "mesh") + return spv::ExecutionModelMeshEXT; + else if (stage == "task") + return spv::ExecutionModelTaskEXT; else SPIRV_CROSS_THROW("Invalid stage."); } @@ -1178,6 +1187,8 @@ static string compile_iteration(const CLIArguments &args, std::vector msl_comp->add_inline_uniform_block(v.first, v.second); for (auto &v : args.msl_shader_inputs) msl_comp->add_msl_shader_input(v); + for (auto &v : args.msl_shader_outputs) + msl_comp->add_msl_shader_output(v); if (args.msl_combined_sampler_suffix) msl_comp->set_combined_sampler_suffix(args.msl_combined_sampler_suffix); } @@ -1581,23 +1592,41 @@ static int main_inner(int argc, char *argv[]) cbs.add("--msl-no-clip-distance-user-varying", [&args](CLIParser &) { args.msl_enable_clip_distance_user_varying = false; }); cbs.add("--msl-shader-input", [&args](CLIParser &parser) { - MSLShaderInput input; + MSLShaderInterfaceVariable input; // Make sure next_uint() is called in-order. input.location = parser.next_uint(); const char *format = parser.next_value_string("other"); if (strcmp(format, "any32") == 0) - input.format = MSL_SHADER_INPUT_FORMAT_ANY32; + input.format = MSL_SHADER_VARIABLE_FORMAT_ANY32; else if (strcmp(format, "any16") == 0) - input.format = MSL_SHADER_INPUT_FORMAT_ANY16; + input.format = MSL_SHADER_VARIABLE_FORMAT_ANY16; else if (strcmp(format, "u16") == 0) - input.format = MSL_SHADER_INPUT_FORMAT_UINT16; + input.format = MSL_SHADER_VARIABLE_FORMAT_UINT16; else if (strcmp(format, "u8") == 0) - input.format = MSL_SHADER_INPUT_FORMAT_UINT8; + input.format = MSL_SHADER_VARIABLE_FORMAT_UINT8; else - input.format = MSL_SHADER_INPUT_FORMAT_OTHER; + input.format = MSL_SHADER_VARIABLE_FORMAT_OTHER; input.vecsize = parser.next_uint(); args.msl_shader_inputs.push_back(input); }); + cbs.add("--msl-shader-output", [&args](CLIParser &parser) { + MSLShaderInterfaceVariable output; + // Make sure next_uint() is called in-order. + output.location = parser.next_uint(); + const char *format = parser.next_value_string("other"); + if (strcmp(format, "any32") == 0) + output.format = MSL_SHADER_VARIABLE_FORMAT_ANY32; + else if (strcmp(format, "any16") == 0) + output.format = MSL_SHADER_VARIABLE_FORMAT_ANY16; + else if (strcmp(format, "u16") == 0) + output.format = MSL_SHADER_VARIABLE_FORMAT_UINT16; + else if (strcmp(format, "u8") == 0) + output.format = MSL_SHADER_VARIABLE_FORMAT_UINT8; + else + output.format = MSL_SHADER_VARIABLE_FORMAT_OTHER; + output.vecsize = parser.next_uint(); + args.msl_shader_outputs.push_back(output); + }); cbs.add("--msl-multi-patch-workgroup", [&args](CLIParser &) { args.msl_multi_patch_workgroup = true; }); cbs.add("--msl-vertex-for-tessellation", [&args](CLIParser &) { args.msl_vertex_for_tessellation = true; }); cbs.add("--msl-additional-fixed-sample-mask", diff --git a/3rdparty/spirv-cross/spirv.h b/3rdparty/spirv-cross/spirv.h index 38f558748..5b6e8aaf4 100644 --- a/3rdparty/spirv-cross/spirv.h +++ b/3rdparty/spirv-cross/spirv.h @@ -98,6 +98,8 @@ typedef enum SpvExecutionModel_ { SpvExecutionModelMissNV = 5317, SpvExecutionModelCallableKHR = 5318, SpvExecutionModelCallableNV = 5318, + SpvExecutionModelTaskEXT = 5364, + SpvExecutionModelMeshEXT = 5365, SpvExecutionModelMax = 0x7fffffff, } SpvExecutionModel; @@ -165,11 +167,21 @@ typedef enum SpvExecutionMode_ { SpvExecutionModeSignedZeroInfNanPreserve = 4461, SpvExecutionModeRoundingModeRTE = 4462, SpvExecutionModeRoundingModeRTZ = 4463, + SpvExecutionModeEarlyAndLateFragmentTestsAMD = 5017, SpvExecutionModeStencilRefReplacingEXT = 5027, + SpvExecutionModeStencilRefUnchangedFrontAMD = 5079, + SpvExecutionModeStencilRefGreaterFrontAMD = 5080, + SpvExecutionModeStencilRefLessFrontAMD = 5081, + SpvExecutionModeStencilRefUnchangedBackAMD = 5082, + SpvExecutionModeStencilRefGreaterBackAMD = 5083, + SpvExecutionModeStencilRefLessBackAMD = 5084, + SpvExecutionModeOutputLinesEXT = 5269, SpvExecutionModeOutputLinesNV = 5269, + SpvExecutionModeOutputPrimitivesEXT = 5270, SpvExecutionModeOutputPrimitivesNV = 5270, SpvExecutionModeDerivativeGroupQuadsNV = 5289, SpvExecutionModeDerivativeGroupLinearNV = 5290, + SpvExecutionModeOutputTrianglesEXT = 5298, SpvExecutionModeOutputTrianglesNV = 5298, SpvExecutionModePixelInterlockOrderedEXT = 5366, SpvExecutionModePixelInterlockUnorderedEXT = 5367, @@ -219,6 +231,7 @@ typedef enum SpvStorageClass_ { SpvStorageClassShaderRecordBufferNV = 5343, SpvStorageClassPhysicalStorageBuffer = 5349, SpvStorageClassPhysicalStorageBufferEXT = 5349, + SpvStorageClassTaskPayloadWorkgroupEXT = 5402, SpvStorageClassCodeSectionINTEL = 5605, SpvStorageClassDeviceOnlyINTEL = 5936, SpvStorageClassHostOnlyINTEL = 5937, @@ -501,6 +514,7 @@ typedef enum SpvDecoration_ { SpvDecorationPassthroughNV = 5250, SpvDecorationViewportRelativeNV = 5252, SpvDecorationSecondaryViewportRelativeNV = 5256, + SpvDecorationPerPrimitiveEXT = 5271, SpvDecorationPerPrimitiveNV = 5271, SpvDecorationPerViewNV = 5272, SpvDecorationPerTaskNV = 5273, @@ -650,6 +664,10 @@ typedef enum SpvBuiltIn_ { SpvBuiltInFragmentSizeNV = 5292, SpvBuiltInFragInvocationCountEXT = 5293, SpvBuiltInInvocationsPerPixelNV = 5293, + SpvBuiltInPrimitivePointIndicesEXT = 5294, + SpvBuiltInPrimitiveLineIndicesEXT = 5295, + SpvBuiltInPrimitiveTriangleIndicesEXT = 5296, + SpvBuiltInCullPrimitiveEXT = 5299, SpvBuiltInLaunchIdKHR = 5319, SpvBuiltInLaunchIdNV = 5319, SpvBuiltInLaunchSizeKHR = 5320, @@ -990,6 +1008,7 @@ typedef enum SpvCapability_ { SpvCapabilityFragmentFullyCoveredEXT = 5265, SpvCapabilityMeshShadingNV = 5266, SpvCapabilityImageFootprintNV = 5282, + SpvCapabilityMeshShadingEXT = 5283, SpvCapabilityFragmentBarycentricKHR = 5284, SpvCapabilityFragmentBarycentricNV = 5284, SpvCapabilityComputeDerivativeGroupQuadsNV = 5288, @@ -1589,6 +1608,8 @@ typedef enum SpvOp_ { SpvOpFragmentFetchAMD = 5012, SpvOpReadClockKHR = 5056, SpvOpImageSampleFootprintNV = 5283, + SpvOpEmitMeshTasksEXT = 5294, + SpvOpSetMeshOutputsEXT = 5295, SpvOpGroupNonUniformPartitionNV = 5296, SpvOpWritePackedPrimitiveIndices4x8NV = 5299, SpvOpReportIntersectionKHR = 5334, @@ -2262,6 +2283,8 @@ inline void SpvHasResultAndType(SpvOp opcode, bool *hasResult, bool *hasResultTy case SpvOpFragmentFetchAMD: *hasResult = true; *hasResultType = true; break; case SpvOpReadClockKHR: *hasResult = true; *hasResultType = true; break; case SpvOpImageSampleFootprintNV: *hasResult = true; *hasResultType = true; break; + case SpvOpEmitMeshTasksEXT: *hasResult = false; *hasResultType = false; break; + case SpvOpSetMeshOutputsEXT: *hasResult = false; *hasResultType = false; break; case SpvOpGroupNonUniformPartitionNV: *hasResult = true; *hasResultType = true; break; case SpvOpWritePackedPrimitiveIndices4x8NV: *hasResult = false; *hasResultType = false; break; case SpvOpReportIntersectionNV: *hasResult = true; *hasResultType = true; break; diff --git a/3rdparty/spirv-cross/spirv.hpp b/3rdparty/spirv-cross/spirv.hpp index 48d93d64c..e25264af2 100644 --- a/3rdparty/spirv-cross/spirv.hpp +++ b/3rdparty/spirv-cross/spirv.hpp @@ -94,6 +94,8 @@ enum ExecutionModel { ExecutionModelMissNV = 5317, ExecutionModelCallableKHR = 5318, ExecutionModelCallableNV = 5318, + ExecutionModelTaskEXT = 5364, + ExecutionModelMeshEXT = 5365, ExecutionModelMax = 0x7fffffff, }; @@ -161,11 +163,21 @@ enum ExecutionMode { ExecutionModeSignedZeroInfNanPreserve = 4461, ExecutionModeRoundingModeRTE = 4462, ExecutionModeRoundingModeRTZ = 4463, + ExecutionModeEarlyAndLateFragmentTestsAMD = 5017, ExecutionModeStencilRefReplacingEXT = 5027, + ExecutionModeStencilRefUnchangedFrontAMD = 5079, + ExecutionModeStencilRefGreaterFrontAMD = 5080, + ExecutionModeStencilRefLessFrontAMD = 5081, + ExecutionModeStencilRefUnchangedBackAMD = 5082, + ExecutionModeStencilRefGreaterBackAMD = 5083, + ExecutionModeStencilRefLessBackAMD = 5084, + ExecutionModeOutputLinesEXT = 5269, ExecutionModeOutputLinesNV = 5269, + ExecutionModeOutputPrimitivesEXT = 5270, ExecutionModeOutputPrimitivesNV = 5270, ExecutionModeDerivativeGroupQuadsNV = 5289, ExecutionModeDerivativeGroupLinearNV = 5290, + ExecutionModeOutputTrianglesEXT = 5298, ExecutionModeOutputTrianglesNV = 5298, ExecutionModePixelInterlockOrderedEXT = 5366, ExecutionModePixelInterlockUnorderedEXT = 5367, @@ -215,6 +227,7 @@ enum StorageClass { StorageClassShaderRecordBufferNV = 5343, StorageClassPhysicalStorageBuffer = 5349, StorageClassPhysicalStorageBufferEXT = 5349, + StorageClassTaskPayloadWorkgroupEXT = 5402, StorageClassCodeSectionINTEL = 5605, StorageClassDeviceOnlyINTEL = 5936, StorageClassHostOnlyINTEL = 5937, @@ -497,6 +510,7 @@ enum Decoration { DecorationPassthroughNV = 5250, DecorationViewportRelativeNV = 5252, DecorationSecondaryViewportRelativeNV = 5256, + DecorationPerPrimitiveEXT = 5271, DecorationPerPrimitiveNV = 5271, DecorationPerViewNV = 5272, DecorationPerTaskNV = 5273, @@ -646,6 +660,10 @@ enum BuiltIn { BuiltInFragmentSizeNV = 5292, BuiltInFragInvocationCountEXT = 5293, BuiltInInvocationsPerPixelNV = 5293, + BuiltInPrimitivePointIndicesEXT = 5294, + BuiltInPrimitiveLineIndicesEXT = 5295, + BuiltInPrimitiveTriangleIndicesEXT = 5296, + BuiltInCullPrimitiveEXT = 5299, BuiltInLaunchIdKHR = 5319, BuiltInLaunchIdNV = 5319, BuiltInLaunchSizeKHR = 5320, @@ -986,6 +1004,7 @@ enum Capability { CapabilityFragmentFullyCoveredEXT = 5265, CapabilityMeshShadingNV = 5266, CapabilityImageFootprintNV = 5282, + CapabilityMeshShadingEXT = 5283, CapabilityFragmentBarycentricKHR = 5284, CapabilityFragmentBarycentricNV = 5284, CapabilityComputeDerivativeGroupQuadsNV = 5288, @@ -1585,6 +1604,8 @@ enum Op { OpFragmentFetchAMD = 5012, OpReadClockKHR = 5056, OpImageSampleFootprintNV = 5283, + OpEmitMeshTasksEXT = 5294, + OpSetMeshOutputsEXT = 5295, OpGroupNonUniformPartitionNV = 5296, OpWritePackedPrimitiveIndices4x8NV = 5299, OpReportIntersectionKHR = 5334, @@ -2258,6 +2279,8 @@ inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) { case OpFragmentFetchAMD: *hasResult = true; *hasResultType = true; break; case OpReadClockKHR: *hasResult = true; *hasResultType = true; break; case OpImageSampleFootprintNV: *hasResult = true; *hasResultType = true; break; + case OpEmitMeshTasksEXT: *hasResult = false; *hasResultType = false; break; + case OpSetMeshOutputsEXT: *hasResult = false; *hasResultType = false; break; case OpGroupNonUniformPartitionNV: *hasResult = true; *hasResultType = true; break; case OpWritePackedPrimitiveIndices4x8NV: *hasResult = false; *hasResultType = false; break; case OpReportIntersectionNV: *hasResult = true; *hasResultType = true; break; diff --git a/3rdparty/spirv-cross/spirv_common.hpp b/3rdparty/spirv-cross/spirv_common.hpp index 1c8a7253c..5c2ad7476 100644 --- a/3rdparty/spirv-cross/spirv_common.hpp +++ b/3rdparty/spirv-cross/spirv_common.hpp @@ -682,6 +682,7 @@ struct SPIREntryPoint } workgroup_size; uint32_t invocations = 0; uint32_t output_vertices = 0; + uint32_t output_primitives = 0; spv::ExecutionModel model = spv::ExecutionModelMax; bool geometry_passthrough = false; }; @@ -776,7 +777,8 @@ struct SPIRBlock : IVariant Unreachable, // Noop Kill, // Discard IgnoreIntersection, // Ray Tracing - TerminateRay // Ray Tracing + TerminateRay, // Ray Tracing + EmitMeshTasks // Mesh shaders }; enum Merge @@ -838,6 +840,13 @@ struct SPIRBlock : IVariant BlockID false_block = 0; BlockID default_block = 0; + // If terminator is EmitMeshTasksEXT. + struct + { + ID groups[3]; + ID payload; + } mesh = {}; + SmallVector ops; struct Phi @@ -1636,6 +1645,12 @@ enum ExtendedDecorations // results of interpolation can. SPIRVCrossDecorationInterpolantComponentExpr, + // Apply to any struct type that is used in the Workgroup storage class. + // This causes matrices in MSL prior to Metal 3.0 to be emitted using a special + // class that is convertible to the standard matrix type, to work around the + // lack of constructors in the 'threadgroup' address space. + SPIRVCrossDecorationWorkgroupStruct, + SPIRVCrossDecorationCount }; diff --git a/3rdparty/spirv-cross/spirv_cross.cpp b/3rdparty/spirv-cross/spirv_cross.cpp index 050c875e8..3f30ee936 100644 --- a/3rdparty/spirv-cross/spirv_cross.cpp +++ b/3rdparty/spirv-cross/spirv_cross.cpp @@ -98,7 +98,8 @@ bool Compiler::block_is_pure(const SPIRBlock &block) // This is a global side effect of the function. if (block.terminator == SPIRBlock::Kill || block.terminator == SPIRBlock::TerminateRay || - block.terminator == SPIRBlock::IgnoreIntersection) + block.terminator == SPIRBlock::IgnoreIntersection || + block.terminator == SPIRBlock::EmitMeshTasks) return false; for (auto &i : block.ops) @@ -154,6 +155,11 @@ bool Compiler::block_is_pure(const SPIRBlock &block) case OpEmitVertex: return false; + // Mesh shader functions modify global state. + // (EmitMeshTasks is a terminator). + case OpSetMeshOutputsEXT: + return false; + // Barriers disallow any reordering, so we should treat blocks with barrier as writing. case OpControlBarrier: case OpMemoryBarrier: @@ -1069,8 +1075,11 @@ void Compiler::parse_fixup() { auto &var = id.get(); if (var.storage == StorageClassPrivate || var.storage == StorageClassWorkgroup || + var.storage == StorageClassTaskPayloadWorkgroupEXT || var.storage == StorageClassOutput) + { global_variables.push_back(var.self); + } if (variable_storage_is_aliased(var)) aliased_variables.push_back(var.self); } @@ -2177,6 +2186,10 @@ void Compiler::set_execution_mode(ExecutionMode mode, uint32_t arg0, uint32_t ar execution.output_vertices = arg0; break; + case ExecutionModeOutputPrimitivesEXT: + execution.output_primitives = arg0; + break; + default: break; } @@ -2297,6 +2310,9 @@ uint32_t Compiler::get_execution_mode_argument(spv::ExecutionMode mode, uint32_t case ExecutionModeOutputVertices: return execution.output_vertices; + case ExecutionModeOutputPrimitivesEXT: + return execution.output_primitives; + default: return 0; } @@ -2359,6 +2375,19 @@ void Compiler::add_implied_read_expression(SPIRAccessChain &e, uint32_t source) e.implied_read_expressions.push_back(source); } +void Compiler::add_active_interface_variable(uint32_t var_id) +{ + active_interface_variables.insert(var_id); + + // In SPIR-V 1.4 and up we must also track the interface variable in the entry point. + if (ir.get_spirv_version() >= 0x10400) + { + auto &vars = get_entry_point().interface_variables; + if (find(begin(vars), end(vars), VariableID(var_id)) == end(vars)) + vars.push_back(var_id); + } +} + void Compiler::inherit_expression_dependencies(uint32_t dst, uint32_t source_expression) { // Don't inherit any expression dependencies if the expression in dst @@ -4410,6 +4439,7 @@ void Compiler::analyze_image_and_sampler_usage() comparison_ids = std::move(handler.comparison_ids); need_subpass_input = handler.need_subpass_input; + need_subpass_input_ms = handler.need_subpass_input_ms; // Forward information from separate images and samplers into combined image samplers. for (auto &combined : combined_image_samplers) @@ -4576,7 +4606,11 @@ bool Compiler::CombinedImageSamplerUsageHandler::handle(Op opcode, const uint32_ // If we load an image, we're going to use it and there is little harm in declaring an unused gl_FragCoord. auto &type = compiler.get(args[0]); if (type.image.dim == DimSubpassData) + { need_subpass_input = true; + if (type.image.ms) + need_subpass_input_ms = true; + } // If we load a SampledImage and it will be used with Dref, propagate the state up. if (dref_combined_samplers.count(args[1]) != 0) diff --git a/3rdparty/spirv-cross/spirv_cross.hpp b/3rdparty/spirv-cross/spirv_cross.hpp index 1d7e7c480..f1c347dd5 100644 --- a/3rdparty/spirv-cross/spirv_cross.hpp +++ b/3rdparty/spirv-cross/spirv_cross.hpp @@ -755,6 +755,7 @@ protected: void inherit_expression_dependencies(uint32_t dst, uint32_t source); void add_implied_read_expression(SPIRExpression &e, uint32_t source); void add_implied_read_expression(SPIRAccessChain &e, uint32_t source); + void add_active_interface_variable(uint32_t var_id); // For proper multiple entry point support, allow querying if an Input or Output // variable is part of that entry points interface. @@ -930,6 +931,7 @@ protected: // Similar is implemented for images, as well as if subpass inputs are needed. std::unordered_set comparison_ids; bool need_subpass_input = false; + bool need_subpass_input_ms = false; // In certain backends, we will need to use a dummy sampler to be able to emit code. // GLSL does not support texelFetch on texture2D objects, but SPIR-V does, @@ -969,6 +971,7 @@ protected: void add_hierarchy_to_comparison_ids(uint32_t ids); bool need_subpass_input = false; + bool need_subpass_input_ms = false; void add_dependency(uint32_t dst, uint32_t src); }; diff --git a/3rdparty/spirv-cross/spirv_cross_c.cpp b/3rdparty/spirv-cross/spirv_cross_c.cpp index 4a62b635c..2d9401b85 100644 --- a/3rdparty/spirv-cross/spirv_cross_c.cpp +++ b/3rdparty/spirv-cross/spirv_cross_c.cpp @@ -1136,9 +1136,9 @@ spvc_result spvc_compiler_msl_add_vertex_attribute(spvc_compiler compiler, const } auto &msl = *static_cast(compiler->compiler.get()); - MSLShaderInput attr; + MSLShaderInterfaceVariable attr; attr.location = va->location; - attr.format = static_cast(va->format); + attr.format = static_cast(va->format); attr.builtin = static_cast(va->builtin); msl.add_msl_shader_input(attr); return SPVC_SUCCESS; @@ -1149,7 +1149,7 @@ spvc_result spvc_compiler_msl_add_vertex_attribute(spvc_compiler compiler, const #endif } -spvc_result spvc_compiler_msl_add_shader_input(spvc_compiler compiler, const spvc_msl_shader_input *si) +spvc_result spvc_compiler_msl_add_shader_input(spvc_compiler compiler, const spvc_msl_shader_interface_var *si) { #if SPIRV_CROSS_C_API_MSL if (compiler->backend != SPVC_BACKEND_MSL) @@ -1159,9 +1159,9 @@ spvc_result spvc_compiler_msl_add_shader_input(spvc_compiler compiler, const spv } auto &msl = *static_cast(compiler->compiler.get()); - MSLShaderInput input; + MSLShaderInterfaceVariable input; input.location = si->location; - input.format = static_cast(si->format); + input.format = static_cast(si->format); input.builtin = static_cast(si->builtin); input.vecsize = si->vecsize; msl.add_msl_shader_input(input); @@ -1173,6 +1173,30 @@ spvc_result spvc_compiler_msl_add_shader_input(spvc_compiler compiler, const spv #endif } +spvc_result spvc_compiler_msl_add_shader_output(spvc_compiler compiler, const spvc_msl_shader_interface_var *so) +{ +#if SPIRV_CROSS_C_API_MSL + if (compiler->backend != SPVC_BACKEND_MSL) + { + compiler->context->report_error("MSL function used on a non-MSL backend."); + return SPVC_ERROR_INVALID_ARGUMENT; + } + + auto &msl = *static_cast(compiler->compiler.get()); + MSLShaderInterfaceVariable output; + output.location = so->location; + output.format = static_cast(so->format); + output.builtin = static_cast(so->builtin); + output.vecsize = so->vecsize; + msl.add_msl_shader_output(output); + return SPVC_SUCCESS; +#else + (void)so; + compiler->context->report_error("MSL function used on a non-MSL backend."); + return SPVC_ERROR_INVALID_ARGUMENT; +#endif +} + spvc_result spvc_compiler_msl_add_resource_binding(spvc_compiler compiler, const spvc_msl_resource_binding *binding) { @@ -1298,6 +1322,24 @@ spvc_bool spvc_compiler_msl_is_shader_input_used(spvc_compiler compiler, unsigne #endif } +spvc_bool spvc_compiler_msl_is_shader_output_used(spvc_compiler compiler, unsigned location) +{ +#if SPIRV_CROSS_C_API_MSL + if (compiler->backend != SPVC_BACKEND_MSL) + { + compiler->context->report_error("MSL function used on a non-MSL backend."); + return SPVC_FALSE; + } + + auto &msl = *static_cast(compiler->compiler.get()); + return msl.is_msl_shader_output_used(location) ? SPVC_TRUE : SPVC_FALSE; +#else + (void)location; + compiler->context->report_error("MSL function used on a non-MSL backend."); + return SPVC_FALSE; +#endif +} + spvc_bool spvc_compiler_msl_is_vertex_attribute_used(spvc_compiler compiler, unsigned location) { return spvc_compiler_msl_is_shader_input_used(compiler, location); @@ -2511,7 +2553,7 @@ void spvc_msl_vertex_attribute_init(spvc_msl_vertex_attribute *attr) { #if SPIRV_CROSS_C_API_MSL // Crude, but works. - MSLShaderInput attr_default; + MSLShaderInterfaceVariable attr_default; attr->location = attr_default.location; attr->format = static_cast(attr_default.format); attr->builtin = static_cast(attr_default.builtin); @@ -2520,19 +2562,24 @@ void spvc_msl_vertex_attribute_init(spvc_msl_vertex_attribute *attr) #endif } -void spvc_msl_shader_input_init(spvc_msl_shader_input *input) +void spvc_msl_shader_interface_var_init(spvc_msl_shader_interface_var *var) { #if SPIRV_CROSS_C_API_MSL - MSLShaderInput input_default; - input->location = input_default.location; - input->format = static_cast(input_default.format); - input->builtin = static_cast(input_default.builtin); - input->vecsize = input_default.vecsize; + MSLShaderInterfaceVariable var_default; + var->location = var_default.location; + var->format = static_cast(var_default.format); + var->builtin = static_cast(var_default.builtin); + var->vecsize = var_default.vecsize; #else - memset(input, 0, sizeof(*input)); + memset(var, 0, sizeof(*var)); #endif } +void spvc_msl_shader_input_init(spvc_msl_shader_input *input) +{ + spvc_msl_shader_interface_var_init(input); +} + void spvc_msl_resource_binding_init(spvc_msl_resource_binding *binding) { #if SPIRV_CROSS_C_API_MSL diff --git a/3rdparty/spirv-cross/spirv_cross_c.h b/3rdparty/spirv-cross/spirv_cross_c.h index a35a5d651..4e6c63978 100644 --- a/3rdparty/spirv-cross/spirv_cross_c.h +++ b/3rdparty/spirv-cross/spirv_cross_c.h @@ -290,23 +290,29 @@ typedef enum spvc_msl_index_type } spvc_msl_index_type; /* Maps to C++ API. */ -typedef enum spvc_msl_shader_input_format +typedef enum spvc_msl_shader_variable_format { - SPVC_MSL_SHADER_INPUT_FORMAT_OTHER = 0, - SPVC_MSL_SHADER_INPUT_FORMAT_UINT8 = 1, - SPVC_MSL_SHADER_INPUT_FORMAT_UINT16 = 2, - SPVC_MSL_SHADER_INPUT_FORMAT_ANY16 = 3, - SPVC_MSL_SHADER_INPUT_FORMAT_ANY32 = 4, + SPVC_MSL_SHADER_VARIABLE_FORMAT_OTHER = 0, + SPVC_MSL_SHADER_VARIABLE_FORMAT_UINT8 = 1, + SPVC_MSL_SHADER_VARIABLE_FORMAT_UINT16 = 2, + SPVC_MSL_SHADER_VARIABLE_FORMAT_ANY16 = 3, + SPVC_MSL_SHADER_VARIABLE_FORMAT_ANY32 = 4, /* Deprecated names. */ - SPVC_MSL_VERTEX_FORMAT_OTHER = SPVC_MSL_SHADER_INPUT_FORMAT_OTHER, - SPVC_MSL_VERTEX_FORMAT_UINT8 = SPVC_MSL_SHADER_INPUT_FORMAT_UINT8, - SPVC_MSL_VERTEX_FORMAT_UINT16 = SPVC_MSL_SHADER_INPUT_FORMAT_UINT16, + SPVC_MSL_VERTEX_FORMAT_OTHER = SPVC_MSL_SHADER_VARIABLE_FORMAT_OTHER, + SPVC_MSL_VERTEX_FORMAT_UINT8 = SPVC_MSL_SHADER_VARIABLE_FORMAT_UINT8, + SPVC_MSL_VERTEX_FORMAT_UINT16 = SPVC_MSL_SHADER_VARIABLE_FORMAT_UINT16, + SPVC_MSL_SHADER_INPUT_FORMAT_OTHER = SPVC_MSL_SHADER_VARIABLE_FORMAT_OTHER, + SPVC_MSL_SHADER_INPUT_FORMAT_UINT8 = SPVC_MSL_SHADER_VARIABLE_FORMAT_UINT8, + SPVC_MSL_SHADER_INPUT_FORMAT_UINT16 = SPVC_MSL_SHADER_VARIABLE_FORMAT_UINT16, + SPVC_MSL_SHADER_INPUT_FORMAT_ANY16 = SPVC_MSL_SHADER_VARIABLE_FORMAT_ANY16, + SPVC_MSL_SHADER_INPUT_FORMAT_ANY32 = SPVC_MSL_SHADER_VARIABLE_FORMAT_ANY32, + SPVC_MSL_SHADER_INPUT_FORMAT_INT_MAX = 0x7fffffff -} spvc_msl_shader_input_format, spvc_msl_vertex_format; +} spvc_msl_shader_variable_format, spvc_msl_shader_input_format, spvc_msl_vertex_format; -/* Maps to C++ API. Deprecated; use spvc_msl_shader_input. */ +/* Maps to C++ API. Deprecated; use spvc_msl_shader_interface_var. */ typedef struct spvc_msl_vertex_attribute { unsigned location; @@ -330,17 +336,21 @@ typedef struct spvc_msl_vertex_attribute SPVC_PUBLIC_API void spvc_msl_vertex_attribute_init(spvc_msl_vertex_attribute *attr); /* Maps to C++ API. */ -typedef struct spvc_msl_shader_input +typedef struct spvc_msl_shader_interface_var { unsigned location; spvc_msl_vertex_format format; SpvBuiltIn builtin; unsigned vecsize; -} spvc_msl_shader_input; +} spvc_msl_shader_interface_var, spvc_msl_shader_input; /* * Initializes the shader input struct. */ +SPVC_PUBLIC_API void spvc_msl_shader_interface_var_init(spvc_msl_shader_interface_var *var); +/* + * Deprecated. Use spvc_msl_shader_interface_var_init(). + */ SPVC_PUBLIC_API void spvc_msl_shader_input_init(spvc_msl_shader_input *input); /* Maps to C++ API. */ @@ -786,13 +796,16 @@ SPVC_PUBLIC_API spvc_result spvc_compiler_msl_add_vertex_attribute(spvc_compiler SPVC_PUBLIC_API spvc_result spvc_compiler_msl_add_resource_binding(spvc_compiler compiler, const spvc_msl_resource_binding *binding); SPVC_PUBLIC_API spvc_result spvc_compiler_msl_add_shader_input(spvc_compiler compiler, - const spvc_msl_shader_input *input); + const spvc_msl_shader_interface_var *input); +SPVC_PUBLIC_API spvc_result spvc_compiler_msl_add_shader_output(spvc_compiler compiler, + const spvc_msl_shader_interface_var *output); SPVC_PUBLIC_API spvc_result spvc_compiler_msl_add_discrete_descriptor_set(spvc_compiler compiler, unsigned desc_set); SPVC_PUBLIC_API spvc_result spvc_compiler_msl_set_argument_buffer_device_address_space(spvc_compiler compiler, unsigned desc_set, spvc_bool device_address); /* Obsolete, use is_shader_input_used. */ SPVC_PUBLIC_API spvc_bool spvc_compiler_msl_is_vertex_attribute_used(spvc_compiler compiler, unsigned location); SPVC_PUBLIC_API spvc_bool spvc_compiler_msl_is_shader_input_used(spvc_compiler compiler, unsigned location); +SPVC_PUBLIC_API spvc_bool spvc_compiler_msl_is_shader_output_used(spvc_compiler compiler, unsigned location); SPVC_PUBLIC_API spvc_bool spvc_compiler_msl_is_resource_used(spvc_compiler compiler, SpvExecutionModel model, diff --git a/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp b/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp index d399d4159..3b4d6e5ca 100644 --- a/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp +++ b/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp @@ -330,6 +330,10 @@ void ParsedIR::fixup_reserved_names() { for (uint32_t id : meta_needing_name_fixup) { + // Don't rename remapped variables like 'gl_LastFragDepthARM'. + if (ids[id].get_type() == TypeVariable && get(id).remapped_variable) + continue; + auto &m = meta[id]; sanitize_identifier(m.decoration.alias, false, false); for (auto &memb : m.members) diff --git a/3rdparty/spirv-cross/spirv_glsl.cpp b/3rdparty/spirv-cross/spirv_glsl.cpp index f47ac62a7..a81296411 100644 --- a/3rdparty/spirv-cross/spirv_glsl.cpp +++ b/3rdparty/spirv-cross/spirv_glsl.cpp @@ -497,6 +497,15 @@ void CompilerGLSL::find_static_extensions() require_extension_internal("GL_NV_ray_tracing"); break; + case ExecutionModelMeshEXT: + case ExecutionModelTaskEXT: + if (options.es || options.version < 450) + SPIRV_CROSS_THROW("Mesh shaders require GLSL 450 or above."); + if (!options.vulkan_semantics) + SPIRV_CROSS_THROW("Mesh shaders require Vulkan semantics."); + require_extension_internal("GL_EXT_mesh_shader"); + break; + default: break; } @@ -649,7 +658,7 @@ string CompilerGLSL::compile() { // only NV_gpu_shader5 supports divergent indexing on OpenGL, and it does so without extra qualifiers backend.nonuniform_qualifier = ""; - backend.needs_row_major_load_workaround = true; + backend.needs_row_major_load_workaround = options.enable_row_major_load_workaround; } backend.allow_precision_qualifiers = options.vulkan_semantics || options.es; backend.force_gl_in_out_block = true; @@ -1060,6 +1069,8 @@ void CompilerGLSL::emit_header() break; case ExecutionModelGLCompute: + case ExecutionModelTaskEXT: + case ExecutionModelMeshEXT: { if (execution.workgroup_size.constant != 0 || execution.flags.get(ExecutionModeLocalSizeId)) { @@ -1078,6 +1089,18 @@ void CompilerGLSL::emit_header() inputs.push_back(join("local_size_y = ", execution.workgroup_size.y)); inputs.push_back(join("local_size_z = ", execution.workgroup_size.z)); } + + if (execution.model == ExecutionModelMeshEXT) + { + outputs.push_back(join("max_vertices = ", execution.output_vertices)); + outputs.push_back(join("max_primitives = ", execution.output_primitives)); + if (execution.flags.get(ExecutionModeOutputTrianglesEXT)) + outputs.push_back("triangles"); + else if (execution.flags.get(ExecutionModeOutputLinesEXT)) + outputs.push_back("lines"); + else if (execution.flags.get(ExecutionModeOutputPoints)) + outputs.push_back("points"); + } break; } @@ -1235,6 +1258,8 @@ string CompilerGLSL::to_interpolation_qualifiers(const Bitset &flags) res += "sample "; if (flags.get(DecorationInvariant)) res += "invariant "; + if (flags.get(DecorationPerPrimitiveEXT)) + res += "perprimitiveEXT "; if (flags.get(DecorationExplicitInterpAMD)) { @@ -2624,7 +2649,7 @@ void CompilerGLSL::emit_interface_block(const SPIRVariable &var) } // Workaround to make sure we can emit "patch in/out" correctly. - fixup_io_block_patch_qualifiers(var); + fixup_io_block_patch_primitive_qualifiers(var); // Block names should never alias. auto block_name = to_name(type.self, false); @@ -2647,8 +2672,15 @@ void CompilerGLSL::emit_interface_block(const SPIRVariable &var) // Instance names cannot alias block names. resource_names.insert(block_name); - bool is_patch = has_decoration(var.self, DecorationPatch); - statement(layout_for_variable(var), (is_patch ? "patch " : ""), qual, block_name); + const char *block_qualifier; + if (has_decoration(var.self, DecorationPatch)) + block_qualifier = "patch "; + else if (has_decoration(var.self, DecorationPerPrimitiveEXT)) + block_qualifier = "perprimitiveEXT "; + else + block_qualifier = ""; + + statement(layout_for_variable(var), block_qualifier, qual, block_name); begin_scope(); type.member_name_cache.clear(); @@ -3084,7 +3116,8 @@ bool CompilerGLSL::should_force_emit_builtin_block(StorageClass storage) }); // If we're declaring clip/cull planes with control points we need to force block declaration. - if (get_execution_model() == ExecutionModelTessellationControl && + if ((get_execution_model() == ExecutionModelTessellationControl || + get_execution_model() == ExecutionModelMeshEXT) && (clip_distance_count || cull_distance_count)) { should_force = true; @@ -3093,7 +3126,7 @@ bool CompilerGLSL::should_force_emit_builtin_block(StorageClass storage) return should_force; } -void CompilerGLSL::fixup_implicit_builtin_block_names() +void CompilerGLSL::fixup_implicit_builtin_block_names(ExecutionModel model) { ir.for_each_typed_id([&](uint32_t, SPIRVariable &var) { auto &type = this->get(var.basetype); @@ -3101,11 +3134,22 @@ void CompilerGLSL::fixup_implicit_builtin_block_names() if ((var.storage == StorageClassOutput || var.storage == StorageClassInput) && block && is_builtin_variable(var)) { - // Make sure the array has a supported name in the code. - if (var.storage == StorageClassOutput) - set_name(var.self, "gl_out"); - else if (var.storage == StorageClassInput) - set_name(var.self, "gl_in"); + if (model != ExecutionModelMeshEXT) + { + // Make sure the array has a supported name in the code. + if (var.storage == StorageClassOutput) + set_name(var.self, "gl_out"); + else if (var.storage == StorageClassInput) + set_name(var.self, "gl_in"); + } + else + { + auto flags = get_buffer_block_flags(var.self); + if (flags.get(DecorationPerPrimitiveEXT)) + set_name(var.self, "gl_MeshPrimitivesEXT"); + else + set_name(var.self, "gl_MeshVerticesEXT"); + } } }); } @@ -3129,6 +3173,11 @@ void CompilerGLSL::emit_declared_builtin_block(StorageClass storage, ExecutionMo uint32_t xfb_stride = 0, xfb_buffer = 0, geom_stream = 0; std::unordered_map builtin_xfb_offsets; + const auto builtin_is_per_vertex_set = [](BuiltIn builtin) -> bool { + return builtin == BuiltInPosition || builtin == BuiltInPointSize || + builtin == BuiltInClipDistance || builtin == BuiltInCullDistance; + }; + ir.for_each_typed_id([&](uint32_t, SPIRVariable &var) { auto &type = this->get(var.basetype); bool block = has_decoration(type.self, DecorationBlock); @@ -3139,7 +3188,7 @@ void CompilerGLSL::emit_declared_builtin_block(StorageClass storage, ExecutionMo uint32_t index = 0; for (auto &m : ir.meta[type.self].members) { - if (m.builtin) + if (m.builtin && builtin_is_per_vertex_set(m.builtin_type)) { builtins.set(m.builtin_type); if (m.builtin_type == BuiltInCullDistance) @@ -3192,7 +3241,7 @@ void CompilerGLSL::emit_declared_builtin_block(StorageClass storage, ExecutionMo { // While we're at it, collect all declared global builtins (HLSL mostly ...). auto &m = ir.meta[var.self].decoration; - if (m.builtin) + if (m.builtin && builtin_is_per_vertex_set(m.builtin_type)) { global_builtins.set(m.builtin_type); if (m.builtin_type == BuiltInCullDistance) @@ -3281,7 +3330,9 @@ void CompilerGLSL::emit_declared_builtin_block(StorageClass storage, ExecutionMo attr.push_back(join("stream = ", geom_stream)); } - if (!attr.empty()) + if (model == ExecutionModelMeshEXT) + statement("out gl_MeshPerVertexEXT"); + else if (!attr.empty()) statement("layout(", merge(attr), ") out gl_PerVertex"); else statement("out gl_PerVertex"); @@ -3399,7 +3450,8 @@ void CompilerGLSL::emit_resources() case ExecutionModelGeometry: case ExecutionModelTessellationControl: case ExecutionModelTessellationEvaluation: - fixup_implicit_builtin_block_names(); + case ExecutionModelMeshEXT: + fixup_implicit_builtin_block_names(execution.model); break; default: @@ -3419,6 +3471,7 @@ void CompilerGLSL::emit_resources() break; case ExecutionModelVertex: + case ExecutionModelMeshEXT: emit_declared_builtin_block(StorageClassOutput, execution.model); break; @@ -8890,6 +8943,15 @@ string CompilerGLSL::builtin_to_glsl(BuiltIn builtin, StorageClass storage) SPIRV_CROSS_THROW("Need desktop GL to use GL_NV_conservative_raster_underestimation."); return "gl_FragFullyCoveredNV"; + case BuiltInPrimitiveTriangleIndicesEXT: + return "gl_PrimitiveTriangleIndicesEXT"; + case BuiltInPrimitiveLineIndicesEXT: + return "gl_PrimitiveLineIndicesEXT"; + case BuiltInPrimitivePointIndicesEXT: + return "gl_PrimitivePointIndicesEXT"; + case BuiltInCullPrimitiveEXT: + return "gl_CullPrimitiveEXT"; + default: return join("gl_BuiltIn_", convert_to_string(builtin)); } @@ -8913,20 +8975,31 @@ const char *CompilerGLSL::index_to_swizzle(uint32_t index) } void CompilerGLSL::access_chain_internal_append_index(std::string &expr, uint32_t /*base*/, const SPIRType * /*type*/, - AccessChainFlags flags, bool & /*access_chain_is_arrayed*/, + AccessChainFlags flags, bool &access_chain_is_arrayed, uint32_t index) { bool index_is_literal = (flags & ACCESS_CHAIN_INDEX_IS_LITERAL_BIT) != 0; + bool ptr_chain = (flags & ACCESS_CHAIN_PTR_CHAIN_BIT) != 0; bool register_expression_read = (flags & ACCESS_CHAIN_SKIP_REGISTER_EXPRESSION_READ_BIT) == 0; - expr += "["; + string idx_expr = index_is_literal ? convert_to_string(index) : to_unpacked_expression(index, register_expression_read); - if (index_is_literal) - expr += convert_to_string(index); + // For the case where the base of an OpPtrAccessChain already ends in [n], + // we need to use the index as an offset to the existing index, otherwise, + // we can just use the index directly. + if (ptr_chain && access_chain_is_arrayed) + { + size_t split_pos = expr.find_last_of(']'); + string expr_front = expr.substr(0, split_pos); + string expr_back = expr.substr(split_pos); + expr = expr_front + " + " + enclose_expression(idx_expr) + expr_back; + } else - expr += to_unpacked_expression(index, register_expression_read); - - expr += "]"; + { + expr += "["; + expr += idx_expr; + expr += "]"; + } } bool CompilerGLSL::access_chain_needs_stage_io_builtin_translation(uint32_t) @@ -8987,10 +9060,12 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice bool pending_array_enclose = false; bool dimension_flatten = false; - const auto append_index = [&](uint32_t index, bool is_literal) { + const auto append_index = [&](uint32_t index, bool is_literal, bool is_ptr_chain = false) { AccessChainFlags mod_flags = flags; if (!is_literal) mod_flags &= ~ACCESS_CHAIN_INDEX_IS_LITERAL_BIT; + if (!is_ptr_chain) + mod_flags &= ~ACCESS_CHAIN_PTR_CHAIN_BIT; access_chain_internal_append_index(expr, base, type, mod_flags, access_chain_is_arrayed, index); check_physical_type_cast(expr, type, physical_type); }; @@ -9043,7 +9118,7 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice } else { - append_index(index, is_literal); + append_index(index, is_literal, true); } if (type->basetype == SPIRType::ControlPointArray) @@ -9078,14 +9153,19 @@ 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. + // Similar concerns apply for mesh shaders where we have to redirect to gl_MeshVerticesEXT or MeshPrimitivesEXT. auto builtin = ir.meta[base].decoration.builtin_type; + bool mesh_shader = get_execution_model() == ExecutionModelMeshEXT; + switch (builtin) { // case BuiltInCullDistance: // These are already arrays, need to figure out rules for these in tess/geom. // case BuiltInClipDistance: case BuiltInPosition: case BuiltInPointSize: - if (var->storage == StorageClassInput) + if (mesh_shader) + expr = join("gl_MeshVerticesEXT[", to_expression(index, register_expression_read), "].", expr); + else if (var->storage == StorageClassInput) expr = join("gl_in[", to_expression(index, register_expression_read), "].", expr); else if (var->storage == StorageClassOutput) expr = join("gl_out[", to_expression(index, register_expression_read), "].", expr); @@ -9093,6 +9173,17 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice append_index(index, is_literal); break; + case BuiltInPrimitiveId: + case BuiltInLayer: + case BuiltInViewportIndex: + case BuiltInCullPrimitiveEXT: + case BuiltInPrimitiveShadingRateKHR: + if (mesh_shader) + expr = join("gl_MeshPrimitivesEXT[", to_expression(index, register_expression_read), "].", expr); + else + append_index(index, is_literal); + break; + default: append_index(index, is_literal); break; @@ -10741,9 +10832,15 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) if (expr_type.vecsize > type.vecsize) expr = enclose_expression(expr + vector_swizzle(type.vecsize, 0)); + if (forward && ptr_expression) + ptr_expression->need_transpose = old_need_transpose; + // We might need to cast in order to load from a builtin. cast_from_variable_load(ptr, expr, type); + if (forward && ptr_expression) + ptr_expression->need_transpose = false; + // We might be trying to load a gl_Position[N], where we should be // doing float4[](gl_in[i].gl_Position, ...) instead. // Similar workarounds are required for input arrays in tessellation. @@ -11261,8 +11358,13 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) // forcing temporaries is not going to help. // This is similar for Constant and Undef inputs. // The only safe thing to RMW is SPIRExpression. + // If the expression has already been used (i.e. used in a continue block), we have to keep using + // that loop variable, since we won't be able to override the expression after the fact. + // If the composite is hoisted, we might never be able to properly invalidate any usage + // of that composite in a subsequent loop iteration. if (invalid_expressions.count(composite) || block_composite_insert_overwrite.count(composite) || + hoisted_temporaries.count(id) || hoisted_temporaries.count(composite) || maybe_get(composite) == nullptr) { can_modify_in_place = false; @@ -12495,12 +12597,12 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) uint32_t result_type = ops[0]; uint32_t id = ops[1]; uint32_t img = ops[2]; + auto &type = expression_type(img); + auto &imgtype = get(type.self); std::string fname = "textureSize"; if (is_legacy_desktop()) { - auto &type = expression_type(img); - auto &imgtype = get(type.self); fname = legacy_tex_op(fname, imgtype, img); } else if (is_legacy_es()) @@ -12508,6 +12610,11 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) auto expr = join(fname, "(", convert_separate_image_to_expression(img), ", ", bitcast_expression(SPIRType::Int, ops[3]), ")"); + + // ES needs to emulate 1D images as 2D. + if (type.image.dim == Dim1D && options.es) + expr = join(expr, ".x"); + auto &restype = get(ops[0]); expr = bitcast_expression(restype, SPIRType::Int, expr); emit_op(result_type, id, expr, true); @@ -13499,6 +13606,10 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) } break; + case OpSetMeshOutputsEXT: + statement("SetMeshOutputsEXT(", to_unpacked_expression(ops[0]), ", ", to_unpacked_expression(ops[1]), ");"); + break; + default: statement("// unimplemented op ", instruction.op); break; @@ -13807,28 +13918,41 @@ string CompilerGLSL::to_precision_qualifiers_glsl(uint32_t id) return flags_to_qualifiers_glsl(type, ir.meta[id].decoration.decoration_flags); } -void CompilerGLSL::fixup_io_block_patch_qualifiers(const SPIRVariable &var) +void CompilerGLSL::fixup_io_block_patch_primitive_qualifiers(const SPIRVariable &var) { // Works around weird behavior in glslangValidator where // a patch out block is translated to just block members getting the decoration. // To make glslang not complain when we compile again, we have to transform this back to a case where // the variable itself has Patch decoration, and not members. + // Same for perprimitiveEXT. auto &type = get(var.basetype); if (has_decoration(type.self, DecorationBlock)) { uint32_t member_count = uint32_t(type.member_types.size()); + Decoration promoted_decoration = {}; + bool do_promote_decoration = false; for (uint32_t i = 0; i < member_count; i++) { if (has_member_decoration(type.self, i, DecorationPatch)) { - set_decoration(var.self, DecorationPatch); + promoted_decoration = DecorationPatch; + do_promote_decoration = true; + break; + } + else if (has_member_decoration(type.self, i, DecorationPerPrimitiveEXT)) + { + promoted_decoration = DecorationPerPrimitiveEXT; + do_promote_decoration = true; break; } } - if (has_decoration(var.self, DecorationPatch)) + if (do_promote_decoration) + { + set_decoration(var.self, promoted_decoration); for (uint32_t i = 0; i < member_count; i++) - unset_member_decoration(type.self, i, DecorationPatch); + unset_member_decoration(type.self, i, promoted_decoration); + } } } @@ -13841,6 +13965,8 @@ string CompilerGLSL::to_qualifiers_glsl(uint32_t id) if (var && var->storage == StorageClassWorkgroup && !backend.shared_is_implied) res += "shared "; + else if (var && var->storage == StorageClassTaskPayloadWorkgroupEXT) + res += "taskPayloadSharedEXT "; res += to_interpolation_qualifiers(flags); if (var) @@ -15998,6 +16124,13 @@ void CompilerGLSL::emit_block_chain(SPIRBlock &block) statement("terminateRayEXT;"); break; + case SPIRBlock::EmitMeshTasks: + statement("EmitMeshTasksEXT(", + to_unpacked_expression(block.mesh.groups[0]), ", ", + to_unpacked_expression(block.mesh.groups[1]), ", ", + to_unpacked_expression(block.mesh.groups[2]), ");"); + break; + default: SPIRV_CROSS_THROW("Unimplemented block terminator."); } @@ -16326,6 +16459,9 @@ void CompilerGLSL::cast_from_variable_load(uint32_t source_id, std::string &expr case BuiltInIncomingRayFlagsNV: case BuiltInLaunchIdNV: case BuiltInLaunchSizeNV: + case BuiltInPrimitiveTriangleIndicesEXT: + case BuiltInPrimitiveLineIndicesEXT: + case BuiltInPrimitivePointIndicesEXT: expected_type = SPIRType::UInt; break; diff --git a/3rdparty/spirv-cross/spirv_glsl.hpp b/3rdparty/spirv-cross/spirv_glsl.hpp index a798737cf..2d1dad6cf 100644 --- a/3rdparty/spirv-cross/spirv_glsl.hpp +++ b/3rdparty/spirv-cross/spirv_glsl.hpp @@ -145,6 +145,12 @@ public: // compares. bool relax_nan_checks = false; + // Loading row-major matrices from UBOs on older AMD Windows OpenGL drivers is problematic. + // To load these types correctly, we must generate a wrapper. them in a dummy function which only purpose is to + // ensure row_major decoration is actually respected. + // This workaround may cause significant performance degeneration on some Android devices. + bool enable_row_major_load_workaround = true; + // If non-zero, controls layout(num_views = N) in; in GL_OVR_multiview2. uint32_t ovr_multiview_view_count = 0; @@ -622,7 +628,7 @@ protected: void emit_buffer_reference_block(uint32_t type_id, bool forward_declaration); void emit_buffer_block_legacy(const SPIRVariable &var); void emit_buffer_block_flattened(const SPIRVariable &type); - void fixup_implicit_builtin_block_names(); + void fixup_implicit_builtin_block_names(spv::ExecutionModel model); void emit_declared_builtin_block(spv::StorageClass storage, spv::ExecutionModel model); bool should_force_emit_builtin_block(spv::StorageClass storage); void emit_push_constant_block_vulkan(const SPIRVariable &var); @@ -766,7 +772,7 @@ protected: std::string type_to_glsl_constructor(const SPIRType &type); std::string argument_decl(const SPIRFunction::Parameter &arg); virtual std::string to_qualifiers_glsl(uint32_t id); - void fixup_io_block_patch_qualifiers(const SPIRVariable &var); + void fixup_io_block_patch_primitive_qualifiers(const SPIRVariable &var); void emit_output_variable_initializer(const SPIRVariable &var); std::string to_precision_qualifiers_glsl(uint32_t id); virtual const char *to_storage_qualifiers_glsl(const SPIRVariable &var); diff --git a/3rdparty/spirv-cross/spirv_msl.cpp b/3rdparty/spirv-cross/spirv_msl.cpp index efd29879d..d81a81892 100644 --- a/3rdparty/spirv-cross/spirv_msl.cpp +++ b/3rdparty/spirv-cross/spirv_msl.cpp @@ -56,13 +56,20 @@ CompilerMSL::CompilerMSL(ParsedIR &&ir_) { } -void CompilerMSL::add_msl_shader_input(const MSLShaderInput &si) +void CompilerMSL::add_msl_shader_input(const MSLShaderInterfaceVariable &si) { inputs_by_location[{si.location, si.component}] = si; if (si.builtin != BuiltInMax && !inputs_by_builtin.count(si.builtin)) inputs_by_builtin[si.builtin] = si; } +void CompilerMSL::add_msl_shader_output(const MSLShaderInterfaceVariable &so) +{ + outputs_by_location[{so.location, so.component}] = so; + if (so.builtin != BuiltInMax && !outputs_by_builtin.count(so.builtin)) + outputs_by_builtin[so.builtin] = so; +} + void CompilerMSL::add_msl_resource_binding(const MSLResourceBinding &binding) { StageSetBinding tuple = { binding.stage, binding.desc_set, binding.binding }; @@ -150,6 +157,13 @@ bool CompilerMSL::is_msl_shader_input_used(uint32_t location) location_inputs_in_use_fallback.count(location) == 0; } +bool CompilerMSL::is_msl_shader_output_used(uint32_t location) +{ + // Don't report internal location allocations to app. + return location_outputs_in_use.count(location) != 0 && + location_outputs_in_use_fallback.count(location) == 0; +} + uint32_t CompilerMSL::get_automatic_builtin_input_location(spv::BuiltIn builtin) const { auto itr = builtin_to_automatic_input_location.find(builtin); @@ -159,6 +173,15 @@ uint32_t CompilerMSL::get_automatic_builtin_input_location(spv::BuiltIn builtin) return itr->second; } +uint32_t CompilerMSL::get_automatic_builtin_output_location(spv::BuiltIn builtin) const +{ + auto itr = builtin_to_automatic_output_location.find(builtin); + if (itr == builtin_to_automatic_output_location.end()) + return k_unknown_location; + else + return itr->second; +} + bool CompilerMSL::is_msl_resource_binding_used(ExecutionModel model, uint32_t desc_set, uint32_t binding) const { StageSetBinding tuple = { model, desc_set, binding }; @@ -1291,7 +1314,7 @@ void CompilerMSL::emit_entry_point_declarations() } // Emit buffer arrays here. - for (uint32_t array_id : buffer_arrays) + for (uint32_t array_id : buffer_arrays_discrete) { const auto &var = get(array_id); const auto &type = get_variable_data_type(var); @@ -1305,8 +1328,57 @@ void CompilerMSL::emit_entry_point_declarations() end_scope_decl(); statement_no_indent(""); } - // For some reason, without this, we end up emitting the arrays twice. - buffer_arrays.clear(); + // Discrete descriptors are processed in entry point emission every compiler iteration. + buffer_arrays_discrete.clear(); + + // Emit buffer aliases here. + for (auto &var_id : buffer_aliases_discrete) + { + const auto &var = get(var_id); + const auto &type = get_variable_data_type(var); + auto addr_space = get_argument_address_space(var); + auto name = to_name(var_id); + + uint32_t desc_set = get_decoration(var_id, DecorationDescriptorSet); + uint32_t desc_binding = get_decoration(var_id, DecorationBinding); + auto alias_name = join("spvBufferAliasSet", desc_set, "Binding", desc_binding); + + statement(addr_space, " auto& ", to_restrict(var_id), + name, + " = *(", addr_space, " ", type_to_glsl(type), "*)", alias_name, ";"); + } + // Discrete descriptors are processed in entry point emission every compiler iteration. + buffer_aliases_discrete.clear(); + + for (auto &var_pair : buffer_aliases_argument) + { + uint32_t var_id = var_pair.first; + uint32_t alias_id = var_pair.second; + + const auto &var = get(var_id); + const auto &type = get_variable_data_type(var); + auto addr_space = get_argument_address_space(var); + + if (type.array.empty()) + { + statement(addr_space, " auto& ", to_restrict(var_id), to_name(var_id), " = (", addr_space, " ", + type_to_glsl(type), "&)", ir.meta[alias_id].decoration.qualified_alias, ";"); + } + else + { + const char *desc_addr_space = descriptor_address_space(var_id, var.storage, "thread"); + + // Esoteric type cast. Reference to array of pointers. + // Auto here defers to UBO or SSBO. The address space of the reference needs to refer to the + // address space of the argument buffer itself, which is usually constant, but can be const device for + // large argument buffers. + is_using_builtin_array = true; + statement(desc_addr_space, " auto& ", to_restrict(var_id), to_name(var_id), " = (", addr_space, " ", + type_to_glsl(type), "* ", desc_addr_space, " (&)", + type_to_array_glsl(type), ")", ir.meta[alias_id].decoration.qualified_alias, ";"); + is_using_builtin_array = false; + } + } // Emit disabled fragment outputs. std::sort(disabled_frag_outputs.begin(), disabled_frag_outputs.end()); @@ -1391,19 +1463,19 @@ string CompilerMSL::compile() activate_argument_buffer_resources(); if (swizzle_buffer_id) - active_interface_variables.insert(swizzle_buffer_id); + add_active_interface_variable(swizzle_buffer_id); if (buffer_size_buffer_id) - active_interface_variables.insert(buffer_size_buffer_id); + add_active_interface_variable(buffer_size_buffer_id); if (view_mask_buffer_id) - active_interface_variables.insert(view_mask_buffer_id); + add_active_interface_variable(view_mask_buffer_id); if (dynamic_offsets_buffer_id) - active_interface_variables.insert(dynamic_offsets_buffer_id); + add_active_interface_variable(dynamic_offsets_buffer_id); if (builtin_layer_id) - active_interface_variables.insert(builtin_layer_id); + add_active_interface_variable(builtin_layer_id); if (builtin_dispatch_base_id && !msl_options.supports_msl_version(1, 2)) - active_interface_variables.insert(builtin_dispatch_base_id); + add_active_interface_variable(builtin_dispatch_base_id); if (builtin_sample_mask_id) - active_interface_variables.insert(builtin_sample_mask_id); + add_active_interface_variable(builtin_sample_mask_id); // Create structs to hold input, output and uniform variables. // Do output first to ensure out. is declared at top of entry function. @@ -1509,7 +1581,7 @@ void CompilerMSL::preprocess_op_codes() // that function would add gl_FragCoord. if (preproc.needs_sample_id || msl_options.force_sample_rate_shading || (is_sample_rate() && (active_input_builtins.get(BuiltInFragCoord) || - (need_subpass_input && !msl_options.use_framebuffer_fetch_subpasses)))) + (need_subpass_input_ms && !msl_options.use_framebuffer_fetch_subpasses)))) needs_sample_id = true; if (is_intersection_query()) @@ -1966,6 +2038,13 @@ void CompilerMSL::mark_packable_structs() (has_decoration(type.self, DecorationBlock) || has_decoration(type.self, DecorationBufferBlock))) mark_as_packable(type); } + + if (var.storage == StorageClassWorkgroup) + { + auto *type = &this->get(var.basetype); + if (type->basetype == SPIRType::Struct) + mark_as_workgroup_struct(*type); + } }); // Physical storage buffer pointers can appear outside of the context of a variable, if the address @@ -2008,19 +2087,63 @@ void CompilerMSL::mark_as_packable(SPIRType &type) } } +// If the specified type is a struct, it and any nested structs +// are marked as used with workgroup storage using the SPIRVCrossDecorationWorkgroupStruct decoration. +void CompilerMSL::mark_as_workgroup_struct(SPIRType &type) +{ + // If this is not the base type (eg. it's a pointer or array), tunnel down + if (type.parent_type) + { + mark_as_workgroup_struct(get(type.parent_type)); + return; + } + + // Handle possible recursion when a struct contains a pointer to its own type nested somewhere. + if (type.basetype == SPIRType::Struct && !has_extended_decoration(type.self, SPIRVCrossDecorationWorkgroupStruct)) + { + set_extended_decoration(type.self, SPIRVCrossDecorationWorkgroupStruct); + + // Recurse + uint32_t mbr_cnt = uint32_t(type.member_types.size()); + for (uint32_t mbr_idx = 0; mbr_idx < mbr_cnt; mbr_idx++) + { + uint32_t mbr_type_id = type.member_types[mbr_idx]; + auto &mbr_type = get(mbr_type_id); + mark_as_workgroup_struct(mbr_type); + if (mbr_type.type_alias) + { + auto &mbr_type_alias = get(mbr_type.type_alias); + mark_as_workgroup_struct(mbr_type_alias); + } + } + } +} + // If a shader input exists at the location, it is marked as being used by this shader void CompilerMSL::mark_location_as_used_by_shader(uint32_t location, const SPIRType &type, StorageClass storage, bool fallback) { - if (storage != StorageClassInput) - return; - uint32_t count = type_to_location_count(type); - for (uint32_t i = 0; i < count; i++) + switch (storage) { - location_inputs_in_use.insert(location + i); - if (fallback) - location_inputs_in_use_fallback.insert(location + i); + case StorageClassInput: + for (uint32_t i = 0; i < count; i++) + { + location_inputs_in_use.insert(location + i); + if (fallback) + location_inputs_in_use_fallback.insert(location + i); + } + break; + case StorageClassOutput: + for (uint32_t i = 0; i < count; i++) + { + location_outputs_in_use.insert(location + i); + if (fallback) + location_outputs_in_use_fallback.insert(location + i); + } + break; + default: + return; } } @@ -2321,6 +2444,12 @@ void CompilerMSL::add_plain_variable_to_interface_block(StorageClass storage, co set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn); mark_location_as_used_by_shader(locn, type, storage); } + else if (is_builtin && capture_output_to_buffer && storage == StorageClassOutput && outputs_by_builtin.count(builtin)) + { + uint32_t locn = outputs_by_builtin[builtin].location; + set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn); + mark_location_as_used_by_shader(locn, type, storage); + } if (get_decoration_bitset(var.self).get(DecorationComponent)) { @@ -2486,6 +2615,12 @@ void CompilerMSL::add_composite_variable_to_interface_block(StorageClass storage set_member_decoration(ib_type.self, ib_mbr_idx, DecorationLocation, locn); mark_location_as_used_by_shader(locn, *usable_type, storage); } + else if (is_builtin && capture_output_to_buffer && storage == StorageClassOutput && outputs_by_builtin.count(builtin)) + { + uint32_t locn = outputs_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); + } else if (is_builtin && (builtin == BuiltInClipDistance || builtin == BuiltInCullDistance)) { // Declare the Clip/CullDistance as [[user(clip/cullN)]]. @@ -2700,6 +2835,13 @@ void CompilerMSL::add_composite_member_variable_to_interface_block(StorageClass mark_location_as_used_by_shader(location, *usable_type, storage); location++; } + else if (is_builtin && capture_output_to_buffer && storage == StorageClassOutput && outputs_by_builtin.count(builtin)) + { + location = outputs_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)) { // Declare the Clip/CullDistance as [[user(clip/cullN)]]. @@ -2894,6 +3036,13 @@ void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass stor mark_location_as_used_by_shader(location, get(mbr_type_id), storage); location++; } + else if (is_builtin && capture_output_to_buffer && storage == StorageClassOutput && outputs_by_builtin.count(builtin)) + { + location = outputs_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. if (has_member_decoration(var_type.self, mbr_idx, DecorationComponent)) @@ -3786,12 +3935,12 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch) SPIRType type; switch (input.second.format) { - case MSL_SHADER_INPUT_FORMAT_UINT16: - case MSL_SHADER_INPUT_FORMAT_ANY16: + case MSL_SHADER_VARIABLE_FORMAT_UINT16: + case MSL_SHADER_VARIABLE_FORMAT_ANY16: type.basetype = SPIRType::UShort; type.width = 16; break; - case MSL_SHADER_INPUT_FORMAT_ANY32: + case MSL_SHADER_VARIABLE_FORMAT_ANY32: default: type.basetype = SPIRType::UInt; type.width = 32; @@ -3823,6 +3972,67 @@ uint32_t CompilerMSL::add_interface_block(StorageClass storage, bool patch) } } + if (capture_output_to_buffer && storage == StorageClassOutput) + { + // For captured output, add all inputs from the next stage to ensure + // the struct containing them is the correct size and layout. This is + // necessary for certain implicit builtins that may nonetheless be read, + // even when they aren't written. + for (auto &output : outputs_by_location) + { + if (location_outputs_in_use.count(output.first.location) != 0) + continue; + + // Create a fake variable to put at the location. + uint32_t offset = ir.increase_bound_by(4); + uint32_t type_id = offset; + uint32_t array_type_id = offset + 1; + uint32_t ptr_type_id = offset + 2; + uint32_t var_id = offset + 3; + + SPIRType type; + switch (output.second.format) + { + case MSL_SHADER_VARIABLE_FORMAT_UINT16: + case MSL_SHADER_VARIABLE_FORMAT_ANY16: + type.basetype = SPIRType::UShort; + type.width = 16; + break; + case MSL_SHADER_VARIABLE_FORMAT_ANY32: + default: + type.basetype = SPIRType::UInt; + type.width = 32; + break; + } + type.vecsize = output.second.vecsize; + set(type_id, type); + + if (get_execution_model() == ExecutionModelTessellationControl) + { + type.array.push_back(0); + type.array_size_literal.push_back(true); + type.parent_type = type_id; + set(array_type_id, type); + } + + type.pointer = true; + type.pointer_depth++; + type.parent_type = get_execution_model() == ExecutionModelTessellationControl ? array_type_id : type_id; + type.storage = storage; + auto &ptr_type = set(ptr_type_id, type); + ptr_type.self = type.parent_type; + + auto &fake_var = set(var_id, ptr_type_id, storage); + set_decoration(var_id, DecorationLocation, output.first.location); + if (output.first.component) + set_decoration(var_id, DecorationComponent, output.first.component); + + meta.strip_array = true; + meta.allow_local_declaration = false; + add_variable_to_interface_block(storage, ib_var_ref, ib_type, fake_var, meta); + } + } + // When multiple variables need to access same location, // unroll locations one by one and we will flatten output or input as necessary. for (auto &loc : meta.location_meta) @@ -3992,7 +4202,7 @@ uint32_t CompilerMSL::ensure_correct_input_type(uint32_t type_id, uint32_t locat switch (p_va->second.format) { - case MSL_SHADER_INPUT_FORMAT_UINT8: + case MSL_SHADER_VARIABLE_FORMAT_UINT8: { switch (type.basetype) { @@ -4016,7 +4226,7 @@ uint32_t CompilerMSL::ensure_correct_input_type(uint32_t type_id, uint32_t locat } } - case MSL_SHADER_INPUT_FORMAT_UINT16: + case MSL_SHADER_VARIABLE_FORMAT_UINT16: { switch (type.basetype) { @@ -4828,6 +5038,10 @@ void CompilerMSL::add_typedef_line(const string &line) // Template struct like spvUnsafeArray<> need to be declared *before* any resources are declared void CompilerMSL::emit_custom_templates() { + static const char * const address_spaces[] = { + "thread", "constant", "device", "threadgroup", "threadgroup_imageblock", "ray_data", "object_data" + }; + for (const auto &spv_func : spv_function_implementations) { switch (spv_func) @@ -4873,6 +5087,122 @@ void CompilerMSL::emit_custom_templates() statement(""); break; + case SPVFuncImplStorageMatrix: + statement("template"); + statement("struct spvStorageMatrix"); + begin_scope(); + statement("vec columns[Cols];"); + statement(""); + for (size_t method_idx = 0; method_idx < sizeof(address_spaces) / sizeof(address_spaces[0]); ++method_idx) + { + // Some address spaces require particular features. + if (method_idx == 4) // threadgroup_imageblock + statement("#ifdef __HAVE_IMAGEBLOCKS__"); + else if (method_idx == 5) // ray_data + statement("#ifdef __HAVE_RAYTRACING__"); + else if (method_idx == 6) // object_data + statement("#ifdef __HAVE_MESH__"); + const string &method_as = address_spaces[method_idx]; + statement("spvStorageMatrix() ", method_as, " = default;"); + if (method_idx != 1) // constant + { + statement(method_as, " spvStorageMatrix& operator=(initializer_list> cols) ", + method_as); + begin_scope(); + statement("size_t i;"); + statement("thread vec* col;"); + statement("for (i = 0, col = cols.begin(); i < Cols; ++i, ++col)"); + statement(" columns[i] = *col;"); + statement("return *this;"); + end_scope(); + } + statement(""); + for (size_t param_idx = 0; param_idx < sizeof(address_spaces) / sizeof(address_spaces[0]); ++param_idx) + { + if (param_idx != method_idx) + { + if (param_idx == 4) // threadgroup_imageblock + statement("#ifdef __HAVE_IMAGEBLOCKS__"); + else if (param_idx == 5) // ray_data + statement("#ifdef __HAVE_RAYTRACING__"); + else if (param_idx == 6) // object_data + statement("#ifdef __HAVE_MESH__"); + } + const string ¶m_as = address_spaces[param_idx]; + statement("spvStorageMatrix(const ", param_as, " matrix& m) ", method_as); + begin_scope(); + statement("for (size_t i = 0; i < Cols; ++i)"); + statement(" columns[i] = m.columns[i];"); + end_scope(); + statement("spvStorageMatrix(const ", param_as, " spvStorageMatrix& m) ", method_as, " = default;"); + if (method_idx != 1) // constant + { + statement(method_as, " spvStorageMatrix& operator=(const ", param_as, + " matrix& m) ", method_as); + begin_scope(); + statement("for (size_t i = 0; i < Cols; ++i)"); + statement(" columns[i] = m.columns[i];"); + statement("return *this;"); + end_scope(); + statement(method_as, " spvStorageMatrix& operator=(const ", param_as, " spvStorageMatrix& m) ", + method_as, " = default;"); + } + if (param_idx != method_idx && param_idx >= 4) + statement("#endif"); + statement(""); + } + statement("operator matrix() const ", method_as); + begin_scope(); + statement("matrix m;"); + statement("for (int i = 0; i < Cols; ++i)"); + statement(" m.columns[i] = columns[i];"); + statement("return m;"); + end_scope(); + statement(""); + statement("vec operator[](size_t idx) const ", method_as); + begin_scope(); + statement("return columns[idx];"); + end_scope(); + if (method_idx != 1) // constant + { + statement(method_as, " vec& operator[](size_t idx) ", method_as); + begin_scope(); + statement("return columns[idx];"); + end_scope(); + } + if (method_idx >= 4) + statement("#endif"); + statement(""); + } + end_scope_decl(); + statement(""); + statement("template"); + statement("matrix transpose(spvStorageMatrix m)"); + begin_scope(); + statement("return transpose(matrix(m));"); + end_scope(); + statement(""); + statement("typedef spvStorageMatrix spvStorage_half2x2;"); + statement("typedef spvStorageMatrix spvStorage_half2x3;"); + statement("typedef spvStorageMatrix spvStorage_half2x4;"); + statement("typedef spvStorageMatrix spvStorage_half3x2;"); + statement("typedef spvStorageMatrix spvStorage_half3x3;"); + statement("typedef spvStorageMatrix spvStorage_half3x4;"); + statement("typedef spvStorageMatrix spvStorage_half4x2;"); + statement("typedef spvStorageMatrix spvStorage_half4x3;"); + statement("typedef spvStorageMatrix spvStorage_half4x4;"); + statement("typedef spvStorageMatrix spvStorage_float2x2;"); + statement("typedef spvStorageMatrix spvStorage_float2x3;"); + statement("typedef spvStorageMatrix spvStorage_float2x4;"); + statement("typedef spvStorageMatrix spvStorage_float3x2;"); + statement("typedef spvStorageMatrix spvStorage_float3x3;"); + statement("typedef spvStorageMatrix spvStorage_float3x4;"); + statement("typedef spvStorageMatrix spvStorage_float4x2;"); + statement("typedef spvStorageMatrix spvStorage_float4x3;"); + statement("typedef spvStorageMatrix spvStorage_float4x4;"); + statement(""); + break; + default: break; } @@ -6958,6 +7288,23 @@ void CompilerMSL::emit_specialization_constants_and_structs() statement(""); } +void CompilerMSL::emit_binary_ptr_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op) +{ + bool forward = should_forward(op0) && should_forward(op1); + emit_op(result_type, result_id, join(to_ptr_expression(op0), " ", op, " ", to_ptr_expression(op1)), forward); + inherit_expression_dependencies(result_id, op0); + inherit_expression_dependencies(result_id, op1); +} + +string CompilerMSL::to_ptr_expression(uint32_t id, bool register_expression_read) +{ + auto *e = maybe_get(id); + auto expr = enclose_expression(e && e->need_transpose ? e->expression : to_expression(id, register_expression_read)); + if (!should_dereference(id)) + expr = address_of_expression(expr); + return expr; +} + void CompilerMSL::emit_binary_unord_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op) { @@ -7746,6 +8093,7 @@ void CompilerMSL::check_physical_type_cast(std::string &expr, const SPIRType *ty void CompilerMSL::emit_instruction(const Instruction &instruction) { #define MSL_BOP(op) emit_binary_op(ops[0], ops[1], ops[2], ops[3], #op) +#define MSL_PTR_BOP(op) emit_binary_ptr_op(ops[0], ops[1], ops[2], ops[3], #op) #define MSL_BOP_CAST(op, type) \ emit_binary_op_cast(ops[0], ops[1], ops[2], ops[3], #op, type, opcode_is_sign_invariant(opcode)) #define MSL_UOP(op) emit_unary_op(ops[0], ops[1], ops[2], #op) @@ -7885,6 +8233,19 @@ void CompilerMSL::emit_instruction(const Instruction &instruction) MSL_UNORD_BOP(<=); break; + // Pointer math + case OpPtrEqual: + MSL_PTR_BOP(==); + break; + + case OpPtrNotEqual: + MSL_PTR_BOP(!=); + break; + + case OpPtrDiff: + MSL_PTR_BOP(-); + break; + // Derivatives case OpDPdx: case OpDPdxFine: @@ -8692,6 +9053,34 @@ void CompilerMSL::emit_instruction(const Instruction &instruction) #undef MSL_RAY_QUERY_IS_OP2 #undef MSL_RAY_QUERY_GET_OP2 #undef MSL_RAY_QUERY_OP_INNER2 + + case OpConvertPtrToU: + case OpConvertUToPtr: + case OpBitcast: + { + auto &type = get(ops[0]); + auto &input_type = expression_type(ops[2]); + + if (opcode != OpBitcast || type.pointer || input_type.pointer) + { + string op; + + if (type.vecsize == 1 && input_type.vecsize == 1) + op = join("reinterpret_cast<", type_to_glsl(type), ">(", to_unpacked_expression(ops[2]), ")"); + else if (input_type.vecsize == 2) + op = join("reinterpret_cast<", type_to_glsl(type), ">(as_type(", to_unpacked_expression(ops[2]), "))"); + else + op = join("as_type<", type_to_glsl(type), ">(reinterpret_cast(", to_unpacked_expression(ops[2]), "))"); + + emit_op(ops[0], ops[1], op, should_forward(ops[2])); + inherit_expression_dependencies(ops[1], ops[2]); + } + else + CompilerGLSL::emit_instruction(instruction); + + break; + } + default: CompilerGLSL::emit_instruction(instruction); break; @@ -9024,6 +9413,9 @@ void CompilerMSL::emit_atomic_func_op(uint32_t result_type, uint32_t result_id, else if (opcode == OpAtomicSMax || opcode == OpAtomicSMin) expected_type = to_signed_basetype(type.width); + if (type.width == 64) + SPIRV_CROSS_THROW("MSL currently does not support 64-bit atomics."); + auto remapped_type = type; remapped_type.basetype = expected_type; @@ -10562,6 +10954,10 @@ string CompilerMSL::to_func_call_arg(const SPIRFunction::Parameter &arg, uint32_ constants.push_back(id); } } + // Dereference pointer variables where needed. + // FIXME: This dereference is actually backwards. We should really just support passing pointer variables between functions. + else if (should_dereference(id)) + arg_str += dereference_expression(type, CompilerGLSL::to_func_call_arg(arg, id)); else arg_str += CompilerGLSL::to_func_call_arg(arg, id); @@ -10872,12 +11268,23 @@ string CompilerMSL::to_struct_member(const SPIRType &type, uint32_t member_type_ else if (!is_scalar(physical_type)) // scalar type is already packed. pack_pfx = "packed_"; } - else if (row_major) + else if (is_matrix(physical_type)) { - // Need to declare type with flipped vecsize/columns. - row_major_physical_type = physical_type; - swap(row_major_physical_type.vecsize, row_major_physical_type.columns); - declared_type = &row_major_physical_type; + if (!msl_options.supports_msl_version(3, 0) && + has_extended_decoration(type.self, SPIRVCrossDecorationWorkgroupStruct)) + { + pack_pfx = "spvStorage_"; + add_spv_func_and_recompile(SPVFuncImplStorageMatrix); + // The pack prefix causes problems with array wrappers. + is_using_builtin_array = true; + } + if (row_major) + { + // Need to declare type with flipped vecsize/columns. + row_major_physical_type = physical_type; + swap(row_major_physical_type.vecsize, row_major_physical_type.columns); + declared_type = &row_major_physical_type; + } } // Very specifically, image load-store in argument buffers are disallowed on MSL on iOS. @@ -10907,8 +11314,8 @@ string CompilerMSL::to_struct_member(const SPIRType &type, uint32_t member_type_ array_type = type_to_array_glsl(physical_type); } - auto result = join(pack_pfx, type_to_glsl(*declared_type, orig_id), " ", qualifier, to_member_name(type, index), - member_attribute_qualifier(type, index), array_type, ";"); + auto result = join(pack_pfx, type_to_glsl(*declared_type, orig_id, true), " ", qualifier, + to_member_name(type, index), member_attribute_qualifier(type, index), array_type, ";"); is_using_builtin_array = false; return result; @@ -11043,6 +11450,16 @@ string CompilerMSL::member_attribute_qualifier(const SPIRType &type, uint32_t in return join(" [[", loc_qual, "]]"); } + if (execution.model == ExecutionModelVertex && msl_options.vertex_for_tessellation && type.storage == StorageClassOutput) + { + // For this type of shader, we always arrange for it to capture its + // output to a buffer. For this reason, qualifiers are irrelevant here. + if (is_builtin) + // We still have to assign a location so the output struct will sort correctly. + get_or_allocate_builtin_output_member_location(builtin, type.self, index); + return ""; + } + // Tessellation control function inputs if (execution.model == ExecutionModelTessellationControl && type.storage == StorageClassInput) { @@ -11085,6 +11502,9 @@ string CompilerMSL::member_attribute_qualifier(const SPIRType &type, uint32_t in { // For this type of shader, we always arrange for it to capture its // output to a buffer. For this reason, qualifiers are irrelevant here. + if (is_builtin) + // We still have to assign a location so the output struct will sort correctly. + get_or_allocate_builtin_output_member_location(builtin, type.self, index); return ""; } @@ -11375,6 +11795,50 @@ uint32_t CompilerMSL::get_or_allocate_builtin_input_member_location(spv::BuiltIn return loc; } +uint32_t CompilerMSL::get_or_allocate_builtin_output_member_location(spv::BuiltIn builtin, + uint32_t type_id, uint32_t index, + uint32_t *comp) +{ + uint32_t loc = get_member_location(type_id, index, comp); + if (loc != k_unknown_location) + return loc; + loc = 0; + + if (comp) + *comp = k_unknown_component; + + // Late allocation. Find a location which is unused by the application. + // This can happen for built-in outputs in tessellation which are mixed and matched with user inputs. + auto &mbr_type = get(get(type_id).member_types[index]); + uint32_t count = type_to_location_count(mbr_type); + + const auto location_range_in_use = [this](uint32_t location, uint32_t location_count) -> bool { + for (uint32_t i = 0; i < location_count; i++) + if (location_outputs_in_use.count(location + i) != 0) + return true; + return false; + }; + + while (location_range_in_use(loc, count)) + loc++; + + set_member_decoration(type_id, index, DecorationLocation, loc); + + // Triangle tess level inputs are shared in one packed float4; + // mark both builtins as sharing one location. + if (get_execution_mode_bitset().get(ExecutionModeTriangles) && + (builtin == BuiltInTessLevelInner || builtin == BuiltInTessLevelOuter)) + { + builtin_to_automatic_output_location[BuiltInTessLevelInner] = loc; + builtin_to_automatic_output_location[BuiltInTessLevelOuter] = loc; + } + else + builtin_to_automatic_output_location[builtin] = loc; + + mark_location_as_used_by_shader(loc, mbr_type, StorageClassOutput, true); + return loc; +} + // Returns the type declaration for a function, including the // entry type if the current function is the entry point function string CompilerMSL::func_type_decl(SPIRType &type) @@ -11659,7 +12123,7 @@ bool CompilerMSL::is_sample_rate() const return get_execution_model() == ExecutionModelFragment && (msl_options.force_sample_rate_shading || std::find(caps.begin(), caps.end(), CapabilitySampleRateShading) != caps.end() || - (msl_options.use_framebuffer_fetch_subpasses && need_subpass_input)); + (msl_options.use_framebuffer_fetch_subpasses && need_subpass_input_ms)); } bool CompilerMSL::is_intersection_query() const @@ -11980,6 +12444,7 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) struct Resource { SPIRVariable *var; + SPIRVariable *descriptor_alias; string name; SPIRType::BaseType basetype; uint32_t index; @@ -12003,6 +12468,37 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) return; } + // Handle descriptor aliasing. We can handle aliasing of buffers by casting pointers, + // but not for typed resources. + SPIRVariable *descriptor_alias = nullptr; + if (var.storage == StorageClassUniform || var.storage == StorageClassStorageBuffer) + { + for (auto &resource : resources) + { + if (get_decoration(resource.var->self, DecorationDescriptorSet) == + get_decoration(var_id, DecorationDescriptorSet) && + get_decoration(resource.var->self, DecorationBinding) == + get_decoration(var_id, DecorationBinding) && + resource.basetype == SPIRType::Struct && type.basetype == SPIRType::Struct && + (resource.var->storage == StorageClassUniform || + resource.var->storage == StorageClassStorageBuffer)) + { + // Possible, but horrible to implement, ignore for now. + if (!type.array.empty()) + SPIRV_CROSS_THROW("Aliasing arrayed discrete descriptors is currently not supported."); + + descriptor_alias = resource.var; + // Self-reference marks that we should declare the resource, + // and it's being used as an alias (so we can emit void* instead). + resource.descriptor_alias = resource.var; + // Need to promote interlocked usage so that the primary declaration is correct. + if (interlocked_resources.count(var_id)) + interlocked_resources.insert(resource.var->self); + break; + } + } + } + const MSLConstexprSampler *constexpr_sampler = nullptr; if (type.basetype == SPIRType::SampledImage || type.basetype == SPIRType::Sampler) { @@ -12030,12 +12526,12 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) plane_count = constexpr_sampler->planes; for (uint32_t i = 0; i < plane_count; i++) - resources.push_back({ &var, to_name(var_id), SPIRType::Image, + resources.push_back({ &var, descriptor_alias, to_name(var_id), SPIRType::Image, get_metal_resource_index(var, SPIRType::Image, i), i, secondary_index }); if (type.image.dim != DimBuffer && !constexpr_sampler) { - resources.push_back({ &var, to_sampler_expression(var_id), SPIRType::Sampler, + resources.push_back({ &var, descriptor_alias, to_sampler_expression(var_id), SPIRType::Sampler, get_metal_resource_index(var, SPIRType::Sampler), 0, 0 }); } } @@ -12043,13 +12539,19 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) { // constexpr samplers are not declared as resources. add_resource_name(var_id); - resources.push_back({ &var, to_name(var_id), type.basetype, - get_metal_resource_index(var, type.basetype), 0, secondary_index }); + + // Don't allocate resource indices for aliases. + uint32_t resource_index = ~0u; + if (!descriptor_alias) + resource_index = get_metal_resource_index(var, type.basetype); + + resources.push_back({ &var, descriptor_alias, to_name(var_id), type.basetype, + resource_index, 0, secondary_index }); } } }); - sort(resources.begin(), resources.end(), [](const Resource &lhs, const Resource &rhs) { + stable_sort(resources.begin(), resources.end(), [](const Resource &lhs, const Resource &rhs) { return tie(lhs.basetype, lhs.index) < tie(rhs.basetype, rhs.index); }); @@ -12067,7 +12569,29 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) auto &m = ir.meta[type.self]; if (m.members.size() == 0) break; - if (!type.array.empty()) + + if (r.descriptor_alias) + { + if (r.var == r.descriptor_alias) + { + auto primary_name = join("spvBufferAliasSet", + get_decoration(var_id, DecorationDescriptorSet), + "Binding", + get_decoration(var_id, DecorationBinding)); + + // Declare the primary alias as void* + if (!ep_args.empty()) + ep_args += ", "; + ep_args += get_argument_address_space(var) + " void* " + primary_name; + ep_args += " [[buffer(" + convert_to_string(r.index) + ")"; + if (interlocked_resources.count(var_id)) + ep_args += ", raster_order_group(0)"; + ep_args += "]]"; + } + + buffer_aliases_discrete.push_back(r.var->self); + } + else if (!type.array.empty()) { if (type.array.size() > 1) SPIRV_CROSS_THROW("Arrays of arrays of buffers are not supported."); @@ -12082,7 +12606,7 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) // Allow Metal to use the array template to make arrays a value type is_using_builtin_array = true; - buffer_arrays.push_back(var_id); + buffer_arrays_discrete.push_back(var_id); for (uint32_t i = 0; i < array_size; ++i) { if (!ep_args.empty()) @@ -13213,6 +13737,8 @@ const std::unordered_set &CompilerMSL::get_reserved_keyword_set() "vertex", "fragment", "compute", + "constant", + "device", "bias", "level", "gradient2d", @@ -13339,6 +13865,8 @@ const std::unordered_set &CompilerMSL::get_reserved_keyword_set() "M_SQRT2", "M_SQRT1_2", "quad_broadcast", + "thread", + "threadgroup", }; return keywords; @@ -13581,7 +14109,7 @@ string CompilerMSL::to_qualifiers_glsl(uint32_t id) // The optional id parameter indicates the object whose type we are trying // to find the description for. It is optional. Most type descriptions do not // depend on a specific object's use of that type. -string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id) +string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id, bool member) { string type_name; @@ -13671,9 +14199,7 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id) // Need to special-case threadgroup booleans. They are supposed to be logical // storage, but MSL compilers will sometimes crash if you use threadgroup bool. // Workaround this by using 16-bit types instead and fixup on load-store to this data. - // FIXME: We have no sane way of working around this problem if a struct member is boolean - // and that struct is used as a threadgroup variable, but ... sigh. - if ((var && var->storage == StorageClassWorkgroup) || type.storage == StorageClassWorkgroup) + if ((var && var->storage == StorageClassWorkgroup) || type.storage == StorageClassWorkgroup || member) type_name = "short"; else type_name = "bool"; @@ -13735,7 +14261,24 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id) // Matrix? if (type.columns > 1) + { + auto *var = maybe_get_backing_variable(id); + if (var && var->basevariable) + var = &get(var->basevariable); + + // Need to special-case threadgroup matrices. Due to an oversight, Metal's + // matrix struct prior to Metal 3 lacks constructors in the threadgroup AS, + // preventing us from default-constructing or initializing matrices in threadgroup storage. + // Work around this by using our own type as storage. + if (((var && var->storage == StorageClassWorkgroup) || type.storage == StorageClassWorkgroup) && + !msl_options.supports_msl_version(3, 0)) + { + add_spv_func_and_recompile(SPVFuncImplStorageMatrix); + type_name = "spvStorage_" + type_name; + } + type_name += to_string(type.columns) + "x"; + } // Vector or Matrix? if (type.vecsize > 1) @@ -13765,6 +14308,11 @@ string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id) } } +string CompilerMSL::type_to_glsl(const SPIRType &type, uint32_t id) +{ + return type_to_glsl(type, id, false); +} + string CompilerMSL::type_to_array_glsl(const SPIRType &type) { // Allow Metal to use the array template to make arrays a value type @@ -13773,17 +14321,14 @@ string CompilerMSL::type_to_array_glsl(const SPIRType &type) case SPIRType::AtomicCounter: case SPIRType::ControlPointArray: case SPIRType::RayQuery: - { return CompilerGLSL::type_to_array_glsl(type); - } + default: - { - if (using_builtin_array()) + if (type_is_array_of_pointers(type) || using_builtin_array()) return CompilerGLSL::type_to_array_glsl(type); else return ""; } - } } string CompilerMSL::constant_op_expression(const SPIRConstantOp &cop) @@ -14400,26 +14945,10 @@ string CompilerMSL::bitcast_glsl_op(const SPIRType &out_type, const SPIRType &in return type_to_glsl(out_type); } -bool CompilerMSL::emit_complex_bitcast(uint32_t result_type, uint32_t id, uint32_t op0) +bool CompilerMSL::emit_complex_bitcast(uint32_t, uint32_t, uint32_t) { - auto &out_type = get(result_type); - auto &in_type = expression_type(op0); - bool uvec2_to_ptr = (in_type.basetype == SPIRType::UInt && in_type.vecsize == 2 && - out_type.pointer && out_type.storage == StorageClassPhysicalStorageBuffer); - bool ptr_to_uvec2 = (in_type.pointer && in_type.storage == StorageClassPhysicalStorageBuffer && - out_type.basetype == SPIRType::UInt && out_type.vecsize == 2); - string expr; - - // Casting between uvec2 and buffer storage pointer per GL_EXT_buffer_reference_uvec2 - if (uvec2_to_ptr) - expr = join("((", type_to_glsl(out_type), ")as_type(", to_unpacked_expression(op0), "))"); - else if (ptr_to_uvec2) - expr = join("as_type<", type_to_glsl(out_type), ">((uint64_t)", to_unpacked_expression(op0), ")"); - else - return false; - - emit_op(result_type, id, expr, should_forward(op0)); - return true; + // This is handled from the outside where we deal with PtrToU/UToPtr and friends. + return false; } // Returns an MSL string identifying the name of a SPIR-V builtin. @@ -15786,13 +16315,40 @@ void CompilerMSL::remap_constexpr_sampler_by_binding(uint32_t desc_set, uint32_t void CompilerMSL::cast_from_variable_load(uint32_t source_id, std::string &expr, const SPIRType &expr_type) { + bool is_packed = has_extended_decoration(source_id, SPIRVCrossDecorationPhysicalTypePacked); + auto *source_expr = maybe_get(source_id); auto *var = maybe_get_backing_variable(source_id); + const SPIRType *var_type, *phys_type; + if (uint32_t phys_id = get_extended_decoration(source_id, SPIRVCrossDecorationPhysicalTypeID)) + phys_type = &get(phys_id); + else + phys_type = &expr_type; if (var) + { source_id = var->self; + var_type = &get_variable_data_type(*var); + } // Type fixups for workgroup variables if they are booleans. - if (var && var->storage == StorageClassWorkgroup && expr_type.basetype == SPIRType::Boolean) + if (var && (var->storage == StorageClassWorkgroup || var_type->basetype == SPIRType::Struct) && + expr_type.basetype == SPIRType::Boolean) expr = join(type_to_glsl(expr_type), "(", expr, ")"); + // Type fixups for workgroup variables if they are matrices. + // Don't do fixup for packed types; those are handled specially. + // FIXME: Maybe use a type like spvStorageMatrix for packed matrices? + if (!msl_options.supports_msl_version(3, 0) && var && + (var->storage == StorageClassWorkgroup || + (var_type->basetype == SPIRType::Struct && + has_extended_decoration(var_type->self, SPIRVCrossDecorationWorkgroupStruct) && !is_packed)) && + expr_type.columns > 1) + { + SPIRType matrix_type = *phys_type; + if (source_expr && source_expr->need_transpose) + swap(matrix_type.vecsize, matrix_type.columns); + matrix_type.array.clear(); + matrix_type.array_size_literal.clear(); + expr = join(type_to_glsl(matrix_type), "(", expr, ")"); + } // Only interested in standalone builtin variables in the switch below. if (!has_decoration(source_id, DecorationBuiltIn)) @@ -15885,17 +16441,42 @@ void CompilerMSL::cast_from_variable_load(uint32_t source_id, std::string &expr, void CompilerMSL::cast_to_variable_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type) { + bool is_packed = has_extended_decoration(target_id, SPIRVCrossDecorationPhysicalTypePacked); + auto *target_expr = maybe_get(target_id); auto *var = maybe_get_backing_variable(target_id); + const SPIRType *var_type, *phys_type; + if (uint32_t phys_id = get_extended_decoration(target_id, SPIRVCrossDecorationPhysicalTypeID)) + phys_type = &get(phys_id); + else + phys_type = &expr_type; if (var) + { target_id = var->self; + var_type = &get_variable_data_type(*var); + } // Type fixups for workgroup variables if they are booleans. - if (var && var->storage == StorageClassWorkgroup && expr_type.basetype == SPIRType::Boolean) + if (var && (var->storage == StorageClassWorkgroup || var_type->basetype == SPIRType::Struct) && + expr_type.basetype == SPIRType::Boolean) { auto short_type = expr_type; short_type.basetype = SPIRType::Short; expr = join(type_to_glsl(short_type), "(", expr, ")"); } + // Type fixups for workgroup variables if they are matrices. + // Don't do fixup for packed types; those are handled specially. + // FIXME: Maybe use a type like spvStorageMatrix for packed matrices? + if (!msl_options.supports_msl_version(3, 0) && var && + (var->storage == StorageClassWorkgroup || + (var_type->basetype == SPIRType::Struct && + has_extended_decoration(var_type->self, SPIRVCrossDecorationWorkgroupStruct) && !is_packed)) && + expr_type.columns > 1) + { + SPIRType matrix_type = *phys_type; + if (target_expr && target_expr->need_transpose) + swap(matrix_type.vecsize, matrix_type.columns); + expr = join("spvStorage_", type_to_glsl(matrix_type), "(", expr, ")"); + } // Only interested in standalone builtin variables. if (!has_decoration(target_id, DecorationBuiltIn)) @@ -16006,6 +16587,7 @@ void CompilerMSL::analyze_argument_buffers() struct Resource { SPIRVariable *var; + SPIRVariable *descriptor_alias; string name; SPIRType::BaseType basetype; uint32_t index; @@ -16045,6 +16627,32 @@ void CompilerMSL::analyze_argument_buffers() } } + // Handle descriptor aliasing as well as we can. + // We can handle aliasing of buffers by casting pointers, but not for typed resources. + // Inline UBOs cannot be handled since it's not a pointer, but inline data. + SPIRVariable *descriptor_alias = nullptr; + if (var.storage == StorageClassUniform || var.storage == StorageClassStorageBuffer) + { + for (auto &resource : resources_in_set[desc_set]) + { + if (get_decoration(resource.var->self, DecorationBinding) == + get_decoration(var_id, DecorationBinding) && + resource.basetype == SPIRType::Struct && type.basetype == SPIRType::Struct && + (resource.var->storage == StorageClassUniform || + resource.var->storage == StorageClassStorageBuffer)) + { + descriptor_alias = resource.var; + // Self-reference marks that we should declare the resource, + // and it's being used as an alias (so we can emit void* instead). + resource.descriptor_alias = resource.var; + // Need to promote interlocked usage so that the primary declaration is correct. + if (interlocked_resources.count(var_id)) + interlocked_resources.insert(resource.var->self); + break; + } + } + } + uint32_t binding = get_decoration(var_id, DecorationBinding); if (type.basetype == SPIRType::SampledImage) { @@ -16058,14 +16666,14 @@ void CompilerMSL::analyze_argument_buffers() { uint32_t image_resource_index = get_metal_resource_index(var, SPIRType::Image, i); resources_in_set[desc_set].push_back( - { &var, to_name(var_id), SPIRType::Image, image_resource_index, i }); + { &var, descriptor_alias, to_name(var_id), SPIRType::Image, image_resource_index, i }); } if (type.image.dim != DimBuffer && !constexpr_sampler) { uint32_t sampler_resource_index = get_metal_resource_index(var, SPIRType::Sampler); resources_in_set[desc_set].push_back( - { &var, to_sampler_expression(var_id), SPIRType::Sampler, sampler_resource_index, 0 }); + { &var, descriptor_alias, to_sampler_expression(var_id), SPIRType::Sampler, sampler_resource_index, 0 }); } } else if (inline_uniform_blocks.count(SetBindingPair{ desc_set, binding })) @@ -16077,15 +16685,20 @@ void CompilerMSL::analyze_argument_buffers() // constexpr samplers are not declared as resources. // Inline uniform blocks are always emitted at the end. add_resource_name(var_id); + + uint32_t resource_index = ~0u; + if (!descriptor_alias) + resource_index = get_metal_resource_index(var, type.basetype); + resources_in_set[desc_set].push_back( - { &var, to_name(var_id), type.basetype, get_metal_resource_index(var, type.basetype), 0 }); + { &var, descriptor_alias, to_name(var_id), type.basetype, resource_index, 0 }); // Emulate texture2D atomic operations if (atomic_image_vars.count(var.self)) { uint32_t buffer_resource_index = get_metal_resource_index(var, SPIRType::AtomicCounter, 0); resources_in_set[desc_set].push_back( - { &var, to_name(var_id) + "_atomic", SPIRType::Struct, buffer_resource_index, 0 }); + { &var, descriptor_alias, to_name(var_id) + "_atomic", SPIRType::Struct, buffer_resource_index, 0 }); } } @@ -16132,7 +16745,7 @@ void CompilerMSL::analyze_argument_buffers() set_decoration(var_id, DecorationDescriptorSet, desc_set); set_decoration(var_id, DecorationBinding, kSwizzleBufferBinding); resources_in_set[desc_set].push_back( - { &var, to_name(var_id), SPIRType::UInt, get_metal_resource_index(var, SPIRType::UInt), 0 }); + { &var, nullptr, to_name(var_id), SPIRType::UInt, get_metal_resource_index(var, SPIRType::UInt), 0 }); } if (set_needs_buffer_sizes[desc_set]) @@ -16143,7 +16756,7 @@ void CompilerMSL::analyze_argument_buffers() set_decoration(var_id, DecorationDescriptorSet, desc_set); set_decoration(var_id, DecorationBinding, kBufferSizeBufferBinding); resources_in_set[desc_set].push_back( - { &var, to_name(var_id), SPIRType::UInt, get_metal_resource_index(var, SPIRType::UInt), 0 }); + { &var, nullptr, to_name(var_id), SPIRType::UInt, get_metal_resource_index(var, SPIRType::UInt), 0 }); } } } @@ -16155,7 +16768,7 @@ void CompilerMSL::analyze_argument_buffers() uint32_t desc_set = get_decoration(var_id, DecorationDescriptorSet); add_resource_name(var_id); resources_in_set[desc_set].push_back( - { &var, to_name(var_id), SPIRType::Struct, get_metal_resource_index(var, SPIRType::Struct), 0 }); + { &var, nullptr, to_name(var_id), SPIRType::Struct, get_metal_resource_index(var, SPIRType::Struct), 0 }); } for (uint32_t desc_set = 0; desc_set < kMaxArgumentBuffers; desc_set++) @@ -16199,7 +16812,7 @@ void CompilerMSL::analyze_argument_buffers() set_name(buffer_variable_id, join("spvDescriptorSet", desc_set)); // Ids must be emitted in ID order. - sort(begin(resources), end(resources), [&](const Resource &lhs, const Resource &rhs) -> bool { + stable_sort(begin(resources), end(resources), [&](const Resource &lhs, const Resource &rhs) -> bool { return tie(lhs.index, lhs.basetype) < tie(rhs.index, rhs.basetype); }); @@ -16304,12 +16917,18 @@ void CompilerMSL::analyze_argument_buffers() } else if (buffers_requiring_dynamic_offset.count(pair)) { + if (resource.descriptor_alias) + SPIRV_CROSS_THROW("Descriptor aliasing is currently not supported with dynamic offsets."); + // Don't set the qualified name here; we'll define a variable holding the corrected buffer address later. buffer_type.member_types.push_back(var.basetype); buffers_requiring_dynamic_offset[pair].second = var.self; } else if (inline_uniform_blocks.count(pair)) { + if (resource.descriptor_alias) + SPIRV_CROSS_THROW("Descriptor aliasing is currently not supported with inline UBOs."); + // Put the buffer block itself into the argument buffer. buffer_type.member_types.push_back(get_variable_data_type_id(var)); set_qualified_name(var.self, join(to_name(buffer_variable_id), ".", mbr_name)); @@ -16341,9 +16960,12 @@ void CompilerMSL::analyze_argument_buffers() } else { - // Resources will be declared as pointers not references, so automatically dereference as appropriate. - buffer_type.member_types.push_back(var.basetype); - if (type.array.empty()) + if (!resource.descriptor_alias || resource.descriptor_alias == resource.var) + buffer_type.member_types.push_back(var.basetype); + + if (resource.descriptor_alias && resource.descriptor_alias != resource.var) + buffer_aliases_argument.push_back({ var.self, resource.descriptor_alias->self }); + else if (type.array.empty()) set_qualified_name(var.self, join("(*", to_name(buffer_variable_id), ".", mbr_name, ")")); else set_qualified_name(var.self, join(to_name(buffer_variable_id), ".", mbr_name)); @@ -16487,7 +17109,7 @@ void CompilerMSL::activate_argument_buffer_resources() uint32_t desc_set = get_decoration(self, DecorationDescriptorSet); if (descriptor_set_is_argument_buffer(desc_set)) - active_interface_variables.insert(self); + add_active_interface_variable(self); }); } diff --git a/3rdparty/spirv-cross/spirv_msl.hpp b/3rdparty/spirv-cross/spirv_msl.hpp index c0317c7a0..1a7ee5c03 100644 --- a/3rdparty/spirv-cross/spirv_msl.hpp +++ b/3rdparty/spirv-cross/spirv_msl.hpp @@ -34,34 +34,39 @@ namespace SPIRV_CROSS_NAMESPACE { -// Indicates the format of a shader input. Currently limited to specifying +// Indicates the format of a shader interface variable. Currently limited to specifying // if the input is an 8-bit unsigned integer, 16-bit unsigned integer, or // some other format. -enum MSLShaderInputFormat +enum MSLShaderVariableFormat { - MSL_SHADER_INPUT_FORMAT_OTHER = 0, - MSL_SHADER_INPUT_FORMAT_UINT8 = 1, - MSL_SHADER_INPUT_FORMAT_UINT16 = 2, - MSL_SHADER_INPUT_FORMAT_ANY16 = 3, - MSL_SHADER_INPUT_FORMAT_ANY32 = 4, + MSL_SHADER_VARIABLE_FORMAT_OTHER = 0, + MSL_SHADER_VARIABLE_FORMAT_UINT8 = 1, + MSL_SHADER_VARIABLE_FORMAT_UINT16 = 2, + MSL_SHADER_VARIABLE_FORMAT_ANY16 = 3, + MSL_SHADER_VARIABLE_FORMAT_ANY32 = 4, // Deprecated aliases. - MSL_VERTEX_FORMAT_OTHER = MSL_SHADER_INPUT_FORMAT_OTHER, - MSL_VERTEX_FORMAT_UINT8 = MSL_SHADER_INPUT_FORMAT_UINT8, - MSL_VERTEX_FORMAT_UINT16 = MSL_SHADER_INPUT_FORMAT_UINT16, + MSL_VERTEX_FORMAT_OTHER = MSL_SHADER_VARIABLE_FORMAT_OTHER, + MSL_VERTEX_FORMAT_UINT8 = MSL_SHADER_VARIABLE_FORMAT_UINT8, + MSL_VERTEX_FORMAT_UINT16 = MSL_SHADER_VARIABLE_FORMAT_UINT16, + MSL_SHADER_INPUT_FORMAT_OTHER = MSL_SHADER_VARIABLE_FORMAT_OTHER, + MSL_SHADER_INPUT_FORMAT_UINT8 = MSL_SHADER_VARIABLE_FORMAT_UINT8, + MSL_SHADER_INPUT_FORMAT_UINT16 = MSL_SHADER_VARIABLE_FORMAT_UINT16, + MSL_SHADER_INPUT_FORMAT_ANY16 = MSL_SHADER_VARIABLE_FORMAT_ANY16, + MSL_SHADER_INPUT_FORMAT_ANY32 = MSL_SHADER_VARIABLE_FORMAT_ANY32, - MSL_SHADER_INPUT_FORMAT_INT_MAX = 0x7fffffff + MSL_SHADER_VARIABLE_FORMAT_INT_MAX = 0x7fffffff }; -// Defines MSL characteristics of an input variable at a particular location. +// Defines MSL characteristics of a shader interface variable at a particular location. // After compilation, it is possible to query whether or not this location was used. // If vecsize is nonzero, it must be greater than or equal to the vecsize declared in the shader, // or behavior is undefined. -struct MSLShaderInput +struct MSLShaderInterfaceVariable { uint32_t location = 0; uint32_t component = 0; - MSLShaderInputFormat format = MSL_SHADER_INPUT_FORMAT_OTHER; + MSLShaderVariableFormat format = MSL_SHADER_VARIABLE_FORMAT_OTHER; spv::BuiltIn builtin = spv::BuiltInMax; uint32_t vecsize = 0; }; @@ -539,10 +544,15 @@ public: explicit CompilerMSL(const ParsedIR &ir); explicit CompilerMSL(ParsedIR &&ir); - // input is a shader input description used to fix up shader input variables. + // input is a shader interface variable description used to fix up shader input variables. // If shader inputs are provided, is_msl_shader_input_used() will return true after - // calling ::compile() if the location was used by the MSL code. - void add_msl_shader_input(const MSLShaderInput &input); + // calling ::compile() if the location were used by the MSL code. + void add_msl_shader_input(const MSLShaderInterfaceVariable &input); + + // output is a shader interface variable description used to fix up shader output variables. + // If shader outputs are provided, is_msl_shader_output_used() will return true after + // calling ::compile() if the location were used by the MSL code. + void add_msl_shader_output(const MSLShaderInterfaceVariable &output); // resource is a resource binding to indicate the MSL buffer, // texture or sampler index to use for a particular SPIR-V description set @@ -577,6 +587,9 @@ public: // Query after compilation is done. This allows you to check if an input location was used by the shader. bool is_msl_shader_input_used(uint32_t location); + // Query after compilation is done. This allows you to check if an output location were used by the shader. + bool is_msl_shader_output_used(uint32_t location); + // If not using add_msl_shader_input, it's possible // that certain builtin attributes need to be automatically assigned locations. // This is typical for tessellation builtin inputs such as tess levels, gl_Position, etc. @@ -584,6 +597,13 @@ public: // add_msl_shader_input or the builtin is not used, otherwise returns N in [[attribute(N)]]. uint32_t get_automatic_builtin_input_location(spv::BuiltIn builtin) const; + // If not using add_msl_shader_output, it's possible + // that certain builtin attributes need to be automatically assigned locations. + // This is typical for tessellation builtin outputs such as tess levels, gl_Position, etc. + // This returns k_unknown_location if the location were explicitly assigned with + // add_msl_shader_output or the builtin were not used, otherwise returns N in [[attribute(N)]]. + uint32_t get_automatic_builtin_output_location(spv::BuiltIn builtin) const; + // NOTE: Only resources which are remapped using add_msl_resource_binding will be reported here. // Constexpr samplers are always assumed to be emitted. // No specific MSLResourceBinding remapping is required for constexpr samplers as long as they are remapped @@ -665,6 +685,7 @@ protected: SPVFuncImplQuantizeToF16, SPVFuncImplCubemapTo2DArrayFace, SPVFuncImplUnsafeArray, // Allow Metal to use the array template to make arrays a value type + SPVFuncImplStorageMatrix, // Allow threadgroup construction of matrices SPVFuncImplInverse4x4, SPVFuncImplInverse3x3, SPVFuncImplInverse2x2, @@ -718,6 +739,8 @@ protected: // If the underlying resource has been used for comparison then duplicate loads of that resource must be too // Use Metal's native frame-buffer fetch API for subpass inputs. void emit_texture_op(const Instruction &i, bool sparse) override; + void emit_binary_ptr_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op); + std::string to_ptr_expression(uint32_t id, bool register_expression_read = true); void emit_binary_unord_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op); void emit_instruction(const Instruction &instr) override; void emit_glsl_op(uint32_t result_type, uint32_t result_id, uint32_t op, const uint32_t *args, @@ -736,6 +759,7 @@ protected: void emit_struct_member(const SPIRType &type, uint32_t member_type_id, uint32_t index, const std::string &qualifier = "", uint32_t base_offset = 0) override; void emit_struct_padding_target(const SPIRType &type) override; + std::string type_to_glsl(const SPIRType &type, uint32_t id, bool member); std::string type_to_glsl(const SPIRType &type, uint32_t id = 0) override; void emit_block_hints(const SPIRBlock &block) override; @@ -796,6 +820,7 @@ protected: void extract_global_variables_from_functions(); void mark_packable_structs(); void mark_as_packable(SPIRType &type); + void mark_as_workgroup_struct(SPIRType &type); std::unordered_map> function_global_vars; void extract_global_variables_from_function(uint32_t func_id, std::set &added_arg_ids, @@ -891,6 +916,8 @@ protected: uint32_t get_member_location(uint32_t type_id, uint32_t index, uint32_t *comp = nullptr) const; uint32_t get_or_allocate_builtin_input_member_location(spv::BuiltIn builtin, uint32_t type_id, uint32_t index, uint32_t *comp = nullptr); + uint32_t get_or_allocate_builtin_output_member_location(spv::BuiltIn builtin, + uint32_t type_id, uint32_t index, uint32_t *comp = nullptr); uint32_t get_physical_tess_level_array_size(spv::BuiltIn builtin) const; @@ -1001,12 +1028,17 @@ protected: Options msl_options; std::set spv_function_implementations; // Must be ordered to ensure declarations are in a specific order. - std::map inputs_by_location; - std::unordered_map inputs_by_builtin; + std::map inputs_by_location; + std::unordered_map inputs_by_builtin; + std::map outputs_by_location; + std::unordered_map outputs_by_builtin; std::unordered_set location_inputs_in_use; std::unordered_set location_inputs_in_use_fallback; + std::unordered_set location_outputs_in_use; + std::unordered_set location_outputs_in_use_fallback; std::unordered_map fragment_output_components; std::unordered_map builtin_to_automatic_input_location; + std::unordered_map builtin_to_automatic_output_location; std::set pragma_lines; std::set typedef_lines; SmallVector vars_needing_early_declaration; @@ -1075,7 +1107,9 @@ protected: const MSLConstexprSampler *find_constexpr_sampler(uint32_t id) const; std::unordered_set buffers_requiring_array_length; - SmallVector buffer_arrays; + SmallVector buffer_arrays_discrete; + SmallVector> buffer_aliases_argument; + SmallVector buffer_aliases_discrete; std::unordered_set atomic_image_vars; // Emulate texture2D atomic operations std::unordered_set pull_model_inputs; diff --git a/3rdparty/spirv-cross/spirv_parser.cpp b/3rdparty/spirv-cross/spirv_parser.cpp index c290a5ebb..49eb1933c 100644 --- a/3rdparty/spirv-cross/spirv_parser.cpp +++ b/3rdparty/spirv-cross/spirv_parser.cpp @@ -183,6 +183,15 @@ void Parser::parse(const Instruction &instruction) auto op = static_cast(instruction.op); uint32_t length = instruction.length; + // HACK for glslang that might emit OpEmitMeshTasksEXT followed by return / branch. + // Instead of failing hard, just ignore it. + if (ignore_trailing_block_opcodes) + { + ignore_trailing_block_opcodes = false; + if (op == OpReturn || op == OpBranch || op == OpUnreachable) + return; + } + switch (op) { case OpSourceContinued: @@ -349,6 +358,10 @@ void Parser::parse(const Instruction &instruction) execution.output_vertices = ops[2]; break; + case ExecutionModeOutputPrimitivesEXT: + execution.output_primitives = ops[2]; + break; + default: break; } @@ -1103,6 +1116,18 @@ void Parser::parse(const Instruction &instruction) current_block = nullptr; break; + case OpEmitMeshTasksEXT: + if (!current_block) + SPIRV_CROSS_THROW("Trying to end a non-existing block."); + current_block->terminator = SPIRBlock::EmitMeshTasks; + for (uint32_t i = 0; i < 3; i++) + current_block->mesh.groups[i] = ops[i]; + current_block->mesh.payload = length >= 4 ? ops[3] : 0; + current_block = nullptr; + // Currently glslang is bugged and does not treat EmitMeshTasksEXT as a terminator. + ignore_trailing_block_opcodes = true; + break; + case OpReturn: { if (!current_block) diff --git a/3rdparty/spirv-cross/spirv_parser.hpp b/3rdparty/spirv-cross/spirv_parser.hpp index d72fc71d8..dabc0e224 100644 --- a/3rdparty/spirv-cross/spirv_parser.hpp +++ b/3rdparty/spirv-cross/spirv_parser.hpp @@ -46,6 +46,8 @@ private: ParsedIR ir; SPIRFunction *current_function = nullptr; SPIRBlock *current_block = nullptr; + // For workarounds. + bool ignore_trailing_block_opcodes = false; void parse(const Instruction &instr); const uint32_t *stream(const Instruction &instr) const;