From 9f5ebeefa40def6f193b8a90da1a7d3e171f757a 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, 10 Jan 2021 11:25:26 -0800 Subject: [PATCH] Updated spirv-cross. --- 3rdparty/spirv-cross/main.cpp | 27 +- 3rdparty/spirv-cross/spirv.h | 108 ++++- 3rdparty/spirv-cross/spirv.hpp | 109 ++++- 3rdparty/spirv-cross/spirv_common.hpp | 4 +- 3rdparty/spirv-cross/spirv_cross.cpp | 109 +++-- 3rdparty/spirv-cross/spirv_cross.hpp | 3 + .../spirv-cross/spirv_cross_parsed_ir.cpp | 5 + .../spirv-cross/spirv_cross_parsed_ir.hpp | 2 + 3rdparty/spirv-cross/spirv_glsl.cpp | 385 ++++++++++++++---- 3rdparty/spirv-cross/spirv_glsl.hpp | 4 + 3rdparty/spirv-cross/spirv_hlsl.cpp | 58 ++- 3rdparty/spirv-cross/spirv_hlsl.hpp | 1 + 3rdparty/spirv-cross/spirv_msl.cpp | 251 +++++++++--- 3rdparty/spirv-cross/spirv_msl.hpp | 7 + 3rdparty/spirv-cross/spirv_parser.cpp | 18 +- 15 files changed, 922 insertions(+), 169 deletions(-) diff --git a/3rdparty/spirv-cross/main.cpp b/3rdparty/spirv-cross/main.cpp index 6f511358b..2d50e1b72 100644 --- a/3rdparty/spirv-cross/main.cpp +++ b/3rdparty/spirv-cross/main.cpp @@ -38,6 +38,11 @@ #include #include +#ifdef _WIN32 +#include +#include +#endif + #ifdef HAVE_SPIRV_CROSS_GIT_VERSION #include "gitversion.h" #endif @@ -220,8 +225,27 @@ struct CLIParser #pragma warning(disable : 4996) #endif +static vector read_spirv_file_stdin() +{ +#ifdef _WIN32 + setmode(fileno(stdin), O_BINARY); +#endif + + vector buffer; + uint32_t tmp[256]; + size_t ret; + + while ((ret = fread(tmp, sizeof(uint32_t), 256, stdin))) + buffer.insert(buffer.end(), tmp, tmp + ret); + + return buffer; +} + static vector read_spirv_file(const char *path) { + if (path[0] == '-' && path[1] == '\0') + return read_spirv_file_stdin(); + FILE *file = fopen(path, "rb"); if (!file) { @@ -849,7 +873,7 @@ static void print_help() // clang-format off fprintf(stderr, "Usage: spirv-cross <...>\n" "\nBasic:\n" - "\t[SPIR-V file]\n" + "\t[SPIR-V file] (- is stdin)\n" "\t[--output ]: If not provided, prints output to stdout.\n" "\t[--dump-resources]:\n\t\tPrints a basic reflection of the SPIR-V module along with other output.\n" "\t[--help]:\n\t\tPrints this help message.\n" @@ -1548,6 +1572,7 @@ static int main_inner(int argc, char *argv[]) cbs.add("--emit-line-directives", [&args](CLIParser &) { args.emit_line_directives = true; }); cbs.default_handler = [&args](const char *value) { args.input = value; }; + cbs.add("-", [&args](CLIParser &) { args.input = "-"; }); cbs.error_handler = [] { print_help(); }; CLIParser parser{ move(cbs), argc - 1, argv + 1 }; diff --git a/3rdparty/spirv-cross/spirv.h b/3rdparty/spirv-cross/spirv.h index dd9850dbb..949f1980e 100644 --- a/3rdparty/spirv-cross/spirv.h +++ b/3rdparty/spirv-cross/spirv.h @@ -54,11 +54,11 @@ typedef unsigned int SpvId; #define SPV_VERSION 0x10500 -#define SPV_REVISION 3 +#define SPV_REVISION 4 static const unsigned int SpvMagicNumber = 0x07230203; static const unsigned int SpvVersion = 0x00010500; -static const unsigned int SpvRevision = 3; +static const unsigned int SpvRevision = 4; static const unsigned int SpvOpCodeMask = 0xffff; static const unsigned int SpvWordCountShift = 16; @@ -172,6 +172,10 @@ typedef enum SpvExecutionMode_ { SpvExecutionModeSampleInterlockUnorderedEXT = 5369, SpvExecutionModeShadingRateInterlockOrderedEXT = 5370, SpvExecutionModeShadingRateInterlockUnorderedEXT = 5371, + SpvExecutionModeMaxWorkgroupSizeINTEL = 5893, + SpvExecutionModeMaxWorkDimINTEL = 5894, + SpvExecutionModeNoGlobalOffsetINTEL = 5895, + SpvExecutionModeNumSIMDWorkitemsINTEL = 5896, SpvExecutionModeMax = 0x7fffffff, } SpvExecutionMode; @@ -203,6 +207,7 @@ typedef enum SpvStorageClass_ { SpvStorageClassShaderRecordBufferNV = 5343, SpvStorageClassPhysicalStorageBuffer = 5349, SpvStorageClassPhysicalStorageBufferEXT = 5349, + SpvStorageClassCodeSectionINTEL = 5605, SpvStorageClassMax = 0x7fffffff, } SpvStorageClass; @@ -273,6 +278,8 @@ typedef enum SpvImageFormat_ { SpvImageFormatRg8ui = 37, SpvImageFormatR16ui = 38, SpvImageFormatR8ui = 39, + SpvImageFormatR64ui = 40, + SpvImageFormatR64i = 41, SpvImageFormatMax = 0x7fffffff, } SpvImageFormat; @@ -481,11 +488,24 @@ typedef enum SpvDecoration_ { SpvDecorationRestrictPointerEXT = 5355, SpvDecorationAliasedPointer = 5356, SpvDecorationAliasedPointerEXT = 5356, + SpvDecorationReferencedIndirectlyINTEL = 5602, SpvDecorationCounterBuffer = 5634, SpvDecorationHlslCounterBufferGOOGLE = 5634, SpvDecorationHlslSemanticGOOGLE = 5635, SpvDecorationUserSemantic = 5635, SpvDecorationUserTypeGOOGLE = 5636, + SpvDecorationRegisterINTEL = 5825, + SpvDecorationMemoryINTEL = 5826, + SpvDecorationNumbanksINTEL = 5827, + SpvDecorationBankwidthINTEL = 5828, + SpvDecorationMaxPrivateCopiesINTEL = 5829, + SpvDecorationSinglepumpINTEL = 5830, + SpvDecorationDoublepumpINTEL = 5831, + SpvDecorationMaxReplicatesINTEL = 5832, + SpvDecorationSimpleDualPortINTEL = 5833, + SpvDecorationMergeINTEL = 5834, + SpvDecorationBankBitsINTEL = 5835, + SpvDecorationForcePow2DepthINTEL = 5836, SpvDecorationMax = 0x7fffffff, } SpvDecoration; @@ -544,8 +564,10 @@ typedef enum SpvBuiltIn_ { SpvBuiltInBaseVertex = 4424, SpvBuiltInBaseInstance = 4425, SpvBuiltInDrawIndex = 4426, + SpvBuiltInPrimitiveShadingRateKHR = 4432, SpvBuiltInDeviceIndex = 4438, SpvBuiltInViewIndex = 4440, + SpvBuiltInShadingRateKHR = 4444, SpvBuiltInBaryCoordNoPerspAMD = 4992, SpvBuiltInBaryCoordNoPerspCentroidAMD = 4993, SpvBuiltInBaryCoordNoPerspSampleAMD = 4994, @@ -596,7 +618,6 @@ typedef enum SpvBuiltIn_ { SpvBuiltInObjectToWorldNV = 5330, SpvBuiltInWorldToObjectKHR = 5331, SpvBuiltInWorldToObjectNV = 5331, - SpvBuiltInHitTKHR = 5332, SpvBuiltInHitTNV = 5332, SpvBuiltInHitKindKHR = 5333, SpvBuiltInHitKindNV = 5333, @@ -632,6 +653,13 @@ typedef enum SpvLoopControlShift_ { SpvLoopControlIterationMultipleShift = 6, SpvLoopControlPeelCountShift = 7, SpvLoopControlPartialCountShift = 8, + SpvLoopControlInitiationIntervalINTELShift = 16, + SpvLoopControlMaxConcurrencyINTELShift = 17, + SpvLoopControlDependencyArrayINTELShift = 18, + SpvLoopControlPipelineEnableINTELShift = 19, + SpvLoopControlLoopCoalesceINTELShift = 20, + SpvLoopControlMaxInterleavingINTELShift = 21, + SpvLoopControlSpeculatedIterationsINTELShift = 22, SpvLoopControlMax = 0x7fffffff, } SpvLoopControlShift; @@ -646,6 +674,13 @@ typedef enum SpvLoopControlMask_ { SpvLoopControlIterationMultipleMask = 0x00000040, SpvLoopControlPeelCountMask = 0x00000080, SpvLoopControlPartialCountMask = 0x00000100, + SpvLoopControlInitiationIntervalINTELMask = 0x00010000, + SpvLoopControlMaxConcurrencyINTELMask = 0x00020000, + SpvLoopControlDependencyArrayINTELMask = 0x00040000, + SpvLoopControlPipelineEnableINTELMask = 0x00080000, + SpvLoopControlLoopCoalesceINTELMask = 0x00100000, + SpvLoopControlMaxInterleavingINTELMask = 0x00200000, + SpvLoopControlSpeculatedIterationsINTELMask = 0x00400000, } SpvLoopControlMask; typedef enum SpvFunctionControlShift_ { @@ -842,6 +877,7 @@ typedef enum SpvCapability_ { SpvCapabilityGroupNonUniformQuad = 68, SpvCapabilityShaderLayer = 69, SpvCapabilityShaderViewportIndex = 70, + SpvCapabilityFragmentShadingRateKHR = 4422, SpvCapabilitySubgroupBallotKHR = 4423, SpvCapabilityDrawParameters = 4427, SpvCapabilitySubgroupVoteKHR = 4431, @@ -866,12 +902,15 @@ typedef enum SpvCapability_ { SpvCapabilityRoundingModeRTE = 4467, SpvCapabilityRoundingModeRTZ = 4468, SpvCapabilityRayQueryProvisionalKHR = 4471, - SpvCapabilityRayTraversalPrimitiveCullingProvisionalKHR = 4478, + SpvCapabilityRayQueryKHR = 4472, + SpvCapabilityRayTraversalPrimitiveCullingKHR = 4478, + SpvCapabilityRayTracingKHR = 4479, SpvCapabilityFloat16ImageAMD = 5008, SpvCapabilityImageGatherBiasLodAMD = 5009, SpvCapabilityFragmentMaskAMD = 5010, SpvCapabilityStencilExportEXT = 5013, SpvCapabilityImageReadWriteLodAMD = 5015, + SpvCapabilityInt64ImageEXT = 5016, SpvCapabilityShaderClockKHR = 5055, SpvCapabilitySampleMaskOverrideCoverageNV = 5249, SpvCapabilityGeometryShaderPassthroughNV = 5251, @@ -932,9 +971,20 @@ typedef enum SpvCapability_ { SpvCapabilitySubgroupImageBlockIOINTEL = 5570, SpvCapabilitySubgroupImageMediaBlockIOINTEL = 5579, SpvCapabilityIntegerFunctions2INTEL = 5584, + SpvCapabilityFunctionPointersINTEL = 5603, + SpvCapabilityIndirectReferencesINTEL = 5604, SpvCapabilitySubgroupAvcMotionEstimationINTEL = 5696, SpvCapabilitySubgroupAvcMotionEstimationIntraINTEL = 5697, SpvCapabilitySubgroupAvcMotionEstimationChromaINTEL = 5698, + SpvCapabilityFPGAMemoryAttributesINTEL = 5824, + SpvCapabilityUnstructuredLoopControlsINTEL = 5886, + SpvCapabilityFPGALoopControlsINTEL = 5888, + SpvCapabilityKernelAttributesINTEL = 5892, + SpvCapabilityFPGAKernelAttributesINTEL = 5897, + SpvCapabilityBlockingPipesINTEL = 5945, + SpvCapabilityFPGARegINTEL = 5948, + SpvCapabilityAtomicFloat32AddEXT = 6033, + SpvCapabilityAtomicFloat64AddEXT = 6034, SpvCapabilityMax = 0x7fffffff, } SpvCapability; @@ -985,6 +1035,22 @@ typedef enum SpvRayQueryCandidateIntersectionType_ { SpvRayQueryCandidateIntersectionTypeMax = 0x7fffffff, } SpvRayQueryCandidateIntersectionType; +typedef enum SpvFragmentShadingRateShift_ { + SpvFragmentShadingRateVertical2PixelsShift = 0, + SpvFragmentShadingRateVertical4PixelsShift = 1, + SpvFragmentShadingRateHorizontal2PixelsShift = 2, + SpvFragmentShadingRateHorizontal4PixelsShift = 3, + SpvFragmentShadingRateMax = 0x7fffffff, +} SpvFragmentShadingRateShift; + +typedef enum SpvFragmentShadingRateMask_ { + SpvFragmentShadingRateMaskNone = 0, + SpvFragmentShadingRateVertical2PixelsMask = 0x00000001, + SpvFragmentShadingRateVertical4PixelsMask = 0x00000002, + SpvFragmentShadingRateHorizontal2PixelsMask = 0x00000004, + SpvFragmentShadingRateHorizontal4PixelsMask = 0x00000008, +} SpvFragmentShadingRateMask; + typedef enum SpvOp_ { SpvOpNop = 0, SpvOpUndef = 1, @@ -1330,13 +1396,19 @@ typedef enum SpvOp_ { SpvOpPtrEqual = 401, SpvOpPtrNotEqual = 402, SpvOpPtrDiff = 403, + SpvOpTerminateInvocation = 4416, SpvOpSubgroupBallotKHR = 4421, SpvOpSubgroupFirstInvocationKHR = 4422, SpvOpSubgroupAllKHR = 4428, SpvOpSubgroupAnyKHR = 4429, SpvOpSubgroupAllEqualKHR = 4430, SpvOpSubgroupReadInvocationKHR = 4432, - SpvOpTypeRayQueryProvisionalKHR = 4472, + SpvOpTraceRayKHR = 4445, + SpvOpExecuteCallableKHR = 4446, + SpvOpConvertUToAccelerationStructureKHR = 4447, + SpvOpIgnoreIntersectionKHR = 4448, + SpvOpTerminateRayKHR = 4449, + SpvOpTypeRayQueryKHR = 4472, SpvOpRayQueryInitializeKHR = 4473, SpvOpRayQueryTerminateKHR = 4474, SpvOpRayQueryGenerateIntersectionKHR = 4475, @@ -1359,15 +1431,11 @@ typedef enum SpvOp_ { SpvOpWritePackedPrimitiveIndices4x8NV = 5299, SpvOpReportIntersectionKHR = 5334, SpvOpReportIntersectionNV = 5334, - SpvOpIgnoreIntersectionKHR = 5335, SpvOpIgnoreIntersectionNV = 5335, - SpvOpTerminateRayKHR = 5336, SpvOpTerminateRayNV = 5336, SpvOpTraceNV = 5337, - SpvOpTraceRayKHR = 5337, SpvOpTypeAccelerationStructureKHR = 5341, SpvOpTypeAccelerationStructureNV = 5341, - SpvOpExecuteCallableKHR = 5344, SpvOpExecuteCallableNV = 5344, SpvOpTypeCooperativeMatrixNV = 5358, SpvOpCooperativeMatrixLoadNV = 5359, @@ -1402,6 +1470,8 @@ typedef enum SpvOp_ { SpvOpUSubSatINTEL = 5596, SpvOpIMul32x16INTEL = 5597, SpvOpUMul32x16INTEL = 5598, + SpvOpFunctionPointerINTEL = 5600, + SpvOpFunctionPointerCallINTEL = 5601, SpvOpDecorateString = 5632, SpvOpDecorateStringGOOGLE = 5632, SpvOpMemberDecorateString = 5633, @@ -1524,6 +1594,10 @@ typedef enum SpvOp_ { SpvOpSubgroupAvcSicGetPackedSkcLumaCountThresholdINTEL = 5814, SpvOpSubgroupAvcSicGetPackedSkcLumaSumThresholdINTEL = 5815, SpvOpSubgroupAvcSicGetInterRawSadsINTEL = 5816, + SpvOpLoopControlINTEL = 5887, + SpvOpReadPipeBlockingINTEL = 5946, + SpvOpWritePipeBlockingINTEL = 5947, + SpvOpFPGARegINTEL = 5949, SpvOpRayQueryGetRayTMinKHR = 6016, SpvOpRayQueryGetRayFlagsKHR = 6017, SpvOpRayQueryGetIntersectionTKHR = 6018, @@ -1541,6 +1615,7 @@ typedef enum SpvOp_ { SpvOpRayQueryGetWorldRayOriginKHR = 6030, SpvOpRayQueryGetIntersectionObjectToWorldKHR = 6031, SpvOpRayQueryGetIntersectionWorldToObjectKHR = 6032, + SpvOpAtomicFAddEXT = 6035, SpvOpMax = 0x7fffffff, } SpvOp; @@ -1893,13 +1968,19 @@ inline void SpvHasResultAndType(SpvOp opcode, bool *hasResult, bool *hasResultTy case SpvOpPtrEqual: *hasResult = true; *hasResultType = true; break; case SpvOpPtrNotEqual: *hasResult = true; *hasResultType = true; break; case SpvOpPtrDiff: *hasResult = true; *hasResultType = true; break; + case SpvOpTerminateInvocation: *hasResult = false; *hasResultType = false; break; case SpvOpSubgroupBallotKHR: *hasResult = true; *hasResultType = true; break; case SpvOpSubgroupFirstInvocationKHR: *hasResult = true; *hasResultType = true; break; case SpvOpSubgroupAllKHR: *hasResult = true; *hasResultType = true; break; case SpvOpSubgroupAnyKHR: *hasResult = true; *hasResultType = true; break; case SpvOpSubgroupAllEqualKHR: *hasResult = true; *hasResultType = true; break; case SpvOpSubgroupReadInvocationKHR: *hasResult = true; *hasResultType = true; break; - case SpvOpTypeRayQueryProvisionalKHR: *hasResult = true; *hasResultType = false; break; + case SpvOpTraceRayKHR: *hasResult = false; *hasResultType = false; break; + case SpvOpExecuteCallableKHR: *hasResult = false; *hasResultType = false; break; + case SpvOpConvertUToAccelerationStructureKHR: *hasResult = true; *hasResultType = true; break; + case SpvOpIgnoreIntersectionKHR: *hasResult = false; *hasResultType = false; break; + case SpvOpTerminateRayKHR: *hasResult = false; *hasResultType = false; break; + case SpvOpTypeRayQueryKHR: *hasResult = true; *hasResultType = false; break; case SpvOpRayQueryInitializeKHR: *hasResult = false; *hasResultType = false; break; case SpvOpRayQueryTerminateKHR: *hasResult = false; *hasResultType = false; break; case SpvOpRayQueryGenerateIntersectionKHR: *hasResult = false; *hasResultType = false; break; @@ -1959,6 +2040,8 @@ inline void SpvHasResultAndType(SpvOp opcode, bool *hasResult, bool *hasResultTy case SpvOpUSubSatINTEL: *hasResult = true; *hasResultType = true; break; case SpvOpIMul32x16INTEL: *hasResult = true; *hasResultType = true; break; case SpvOpUMul32x16INTEL: *hasResult = true; *hasResultType = true; break; + case SpvOpFunctionPointerINTEL: *hasResult = true; *hasResultType = true; break; + case SpvOpFunctionPointerCallINTEL: *hasResult = true; *hasResultType = true; break; case SpvOpDecorateString: *hasResult = false; *hasResultType = false; break; case SpvOpMemberDecorateString: *hasResult = false; *hasResultType = false; break; case SpvOpVmeImageINTEL: *hasResult = true; *hasResultType = true; break; @@ -2079,6 +2162,10 @@ inline void SpvHasResultAndType(SpvOp opcode, bool *hasResult, bool *hasResultTy case SpvOpSubgroupAvcSicGetPackedSkcLumaCountThresholdINTEL: *hasResult = true; *hasResultType = true; break; case SpvOpSubgroupAvcSicGetPackedSkcLumaSumThresholdINTEL: *hasResult = true; *hasResultType = true; break; case SpvOpSubgroupAvcSicGetInterRawSadsINTEL: *hasResult = true; *hasResultType = true; break; + case SpvOpLoopControlINTEL: *hasResult = false; *hasResultType = false; break; + case SpvOpReadPipeBlockingINTEL: *hasResult = true; *hasResultType = true; break; + case SpvOpWritePipeBlockingINTEL: *hasResult = true; *hasResultType = true; break; + case SpvOpFPGARegINTEL: *hasResult = true; *hasResultType = true; break; case SpvOpRayQueryGetRayTMinKHR: *hasResult = true; *hasResultType = true; break; case SpvOpRayQueryGetRayFlagsKHR: *hasResult = true; *hasResultType = true; break; case SpvOpRayQueryGetIntersectionTKHR: *hasResult = true; *hasResultType = true; break; @@ -2096,6 +2183,7 @@ inline void SpvHasResultAndType(SpvOp opcode, bool *hasResult, bool *hasResultTy case SpvOpRayQueryGetWorldRayOriginKHR: *hasResult = true; *hasResultType = true; break; case SpvOpRayQueryGetIntersectionObjectToWorldKHR: *hasResult = true; *hasResultType = true; break; case SpvOpRayQueryGetIntersectionWorldToObjectKHR: *hasResult = true; *hasResultType = true; break; + case SpvOpAtomicFAddEXT: *hasResult = true; *hasResultType = true; break; } } #endif /* SPV_ENABLE_UTILITY_CODE */ diff --git a/3rdparty/spirv-cross/spirv.hpp b/3rdparty/spirv-cross/spirv.hpp index dae36cf20..43dd2aaee 100644 --- a/3rdparty/spirv-cross/spirv.hpp +++ b/3rdparty/spirv-cross/spirv.hpp @@ -50,11 +50,11 @@ namespace spv { typedef unsigned int Id; #define SPV_VERSION 0x10500 -#define SPV_REVISION 3 +#define SPV_REVISION 4 static const unsigned int MagicNumber = 0x07230203; static const unsigned int Version = 0x00010500; -static const unsigned int Revision = 3; +static const unsigned int Revision = 4; static const unsigned int OpCodeMask = 0xffff; static const unsigned int WordCountShift = 16; @@ -168,6 +168,10 @@ enum ExecutionMode { ExecutionModeSampleInterlockUnorderedEXT = 5369, ExecutionModeShadingRateInterlockOrderedEXT = 5370, ExecutionModeShadingRateInterlockUnorderedEXT = 5371, + ExecutionModeMaxWorkgroupSizeINTEL = 5893, + ExecutionModeMaxWorkDimINTEL = 5894, + ExecutionModeNoGlobalOffsetINTEL = 5895, + ExecutionModeNumSIMDWorkitemsINTEL = 5896, ExecutionModeMax = 0x7fffffff, }; @@ -199,6 +203,7 @@ enum StorageClass { StorageClassShaderRecordBufferNV = 5343, StorageClassPhysicalStorageBuffer = 5349, StorageClassPhysicalStorageBufferEXT = 5349, + StorageClassCodeSectionINTEL = 5605, StorageClassMax = 0x7fffffff, }; @@ -269,6 +274,8 @@ enum ImageFormat { ImageFormatRg8ui = 37, ImageFormatR16ui = 38, ImageFormatR8ui = 39, + ImageFormatR64ui = 40, + ImageFormatR64i = 41, ImageFormatMax = 0x7fffffff, }; @@ -477,11 +484,24 @@ enum Decoration { DecorationRestrictPointerEXT = 5355, DecorationAliasedPointer = 5356, DecorationAliasedPointerEXT = 5356, + DecorationReferencedIndirectlyINTEL = 5602, DecorationCounterBuffer = 5634, DecorationHlslCounterBufferGOOGLE = 5634, DecorationHlslSemanticGOOGLE = 5635, DecorationUserSemantic = 5635, DecorationUserTypeGOOGLE = 5636, + DecorationRegisterINTEL = 5825, + DecorationMemoryINTEL = 5826, + DecorationNumbanksINTEL = 5827, + DecorationBankwidthINTEL = 5828, + DecorationMaxPrivateCopiesINTEL = 5829, + DecorationSinglepumpINTEL = 5830, + DecorationDoublepumpINTEL = 5831, + DecorationMaxReplicatesINTEL = 5832, + DecorationSimpleDualPortINTEL = 5833, + DecorationMergeINTEL = 5834, + DecorationBankBitsINTEL = 5835, + DecorationForcePow2DepthINTEL = 5836, DecorationMax = 0x7fffffff, }; @@ -540,8 +560,10 @@ enum BuiltIn { BuiltInBaseVertex = 4424, BuiltInBaseInstance = 4425, BuiltInDrawIndex = 4426, + BuiltInPrimitiveShadingRateKHR = 4432, BuiltInDeviceIndex = 4438, BuiltInViewIndex = 4440, + BuiltInShadingRateKHR = 4444, BuiltInBaryCoordNoPerspAMD = 4992, BuiltInBaryCoordNoPerspCentroidAMD = 4993, BuiltInBaryCoordNoPerspSampleAMD = 4994, @@ -592,7 +614,6 @@ enum BuiltIn { BuiltInObjectToWorldNV = 5330, BuiltInWorldToObjectKHR = 5331, BuiltInWorldToObjectNV = 5331, - BuiltInHitTKHR = 5332, BuiltInHitTNV = 5332, BuiltInHitKindKHR = 5333, BuiltInHitKindNV = 5333, @@ -628,6 +649,13 @@ enum LoopControlShift { LoopControlIterationMultipleShift = 6, LoopControlPeelCountShift = 7, LoopControlPartialCountShift = 8, + LoopControlInitiationIntervalINTELShift = 16, + LoopControlMaxConcurrencyINTELShift = 17, + LoopControlDependencyArrayINTELShift = 18, + LoopControlPipelineEnableINTELShift = 19, + LoopControlLoopCoalesceINTELShift = 20, + LoopControlMaxInterleavingINTELShift = 21, + LoopControlSpeculatedIterationsINTELShift = 22, LoopControlMax = 0x7fffffff, }; @@ -642,6 +670,13 @@ enum LoopControlMask { LoopControlIterationMultipleMask = 0x00000040, LoopControlPeelCountMask = 0x00000080, LoopControlPartialCountMask = 0x00000100, + LoopControlInitiationIntervalINTELMask = 0x00010000, + LoopControlMaxConcurrencyINTELMask = 0x00020000, + LoopControlDependencyArrayINTELMask = 0x00040000, + LoopControlPipelineEnableINTELMask = 0x00080000, + LoopControlLoopCoalesceINTELMask = 0x00100000, + LoopControlMaxInterleavingINTELMask = 0x00200000, + LoopControlSpeculatedIterationsINTELMask = 0x00400000, }; enum FunctionControlShift { @@ -838,6 +873,7 @@ enum Capability { CapabilityGroupNonUniformQuad = 68, CapabilityShaderLayer = 69, CapabilityShaderViewportIndex = 70, + CapabilityFragmentShadingRateKHR = 4422, CapabilitySubgroupBallotKHR = 4423, CapabilityDrawParameters = 4427, CapabilitySubgroupVoteKHR = 4431, @@ -862,12 +898,15 @@ enum Capability { CapabilityRoundingModeRTE = 4467, CapabilityRoundingModeRTZ = 4468, CapabilityRayQueryProvisionalKHR = 4471, - CapabilityRayTraversalPrimitiveCullingProvisionalKHR = 4478, + CapabilityRayQueryKHR = 4472, + CapabilityRayTraversalPrimitiveCullingKHR = 4478, + CapabilityRayTracingKHR = 4479, CapabilityFloat16ImageAMD = 5008, CapabilityImageGatherBiasLodAMD = 5009, CapabilityFragmentMaskAMD = 5010, CapabilityStencilExportEXT = 5013, CapabilityImageReadWriteLodAMD = 5015, + CapabilityInt64ImageEXT = 5016, CapabilityShaderClockKHR = 5055, CapabilitySampleMaskOverrideCoverageNV = 5249, CapabilityGeometryShaderPassthroughNV = 5251, @@ -928,9 +967,20 @@ enum Capability { CapabilitySubgroupImageBlockIOINTEL = 5570, CapabilitySubgroupImageMediaBlockIOINTEL = 5579, CapabilityIntegerFunctions2INTEL = 5584, + CapabilityFunctionPointersINTEL = 5603, + CapabilityIndirectReferencesINTEL = 5604, CapabilitySubgroupAvcMotionEstimationINTEL = 5696, CapabilitySubgroupAvcMotionEstimationIntraINTEL = 5697, CapabilitySubgroupAvcMotionEstimationChromaINTEL = 5698, + CapabilityFPGAMemoryAttributesINTEL = 5824, + CapabilityUnstructuredLoopControlsINTEL = 5886, + CapabilityFPGALoopControlsINTEL = 5888, + CapabilityKernelAttributesINTEL = 5892, + CapabilityFPGAKernelAttributesINTEL = 5897, + CapabilityBlockingPipesINTEL = 5945, + CapabilityFPGARegINTEL = 5948, + CapabilityAtomicFloat32AddEXT = 6033, + CapabilityAtomicFloat64AddEXT = 6034, CapabilityMax = 0x7fffffff, }; @@ -981,6 +1031,22 @@ enum RayQueryCandidateIntersectionType { RayQueryCandidateIntersectionTypeMax = 0x7fffffff, }; +enum FragmentShadingRateShift { + FragmentShadingRateVertical2PixelsShift = 0, + FragmentShadingRateVertical4PixelsShift = 1, + FragmentShadingRateHorizontal2PixelsShift = 2, + FragmentShadingRateHorizontal4PixelsShift = 3, + FragmentShadingRateMax = 0x7fffffff, +}; + +enum FragmentShadingRateMask { + FragmentShadingRateMaskNone = 0, + FragmentShadingRateVertical2PixelsMask = 0x00000001, + FragmentShadingRateVertical4PixelsMask = 0x00000002, + FragmentShadingRateHorizontal2PixelsMask = 0x00000004, + FragmentShadingRateHorizontal4PixelsMask = 0x00000008, +}; + enum Op { OpNop = 0, OpUndef = 1, @@ -1326,13 +1392,19 @@ enum Op { OpPtrEqual = 401, OpPtrNotEqual = 402, OpPtrDiff = 403, + OpTerminateInvocation = 4416, OpSubgroupBallotKHR = 4421, OpSubgroupFirstInvocationKHR = 4422, OpSubgroupAllKHR = 4428, OpSubgroupAnyKHR = 4429, OpSubgroupAllEqualKHR = 4430, OpSubgroupReadInvocationKHR = 4432, - OpTypeRayQueryProvisionalKHR = 4472, + OpTraceRayKHR = 4445, + OpExecuteCallableKHR = 4446, + OpConvertUToAccelerationStructureKHR = 4447, + OpIgnoreIntersectionKHR = 4448, + OpTerminateRayKHR = 4449, + OpTypeRayQueryKHR = 4472, OpRayQueryInitializeKHR = 4473, OpRayQueryTerminateKHR = 4474, OpRayQueryGenerateIntersectionKHR = 4475, @@ -1355,15 +1427,11 @@ enum Op { OpWritePackedPrimitiveIndices4x8NV = 5299, OpReportIntersectionKHR = 5334, OpReportIntersectionNV = 5334, - OpIgnoreIntersectionKHR = 5335, OpIgnoreIntersectionNV = 5335, - OpTerminateRayKHR = 5336, OpTerminateRayNV = 5336, OpTraceNV = 5337, - OpTraceRayKHR = 5337, OpTypeAccelerationStructureKHR = 5341, OpTypeAccelerationStructureNV = 5341, - OpExecuteCallableKHR = 5344, OpExecuteCallableNV = 5344, OpTypeCooperativeMatrixNV = 5358, OpCooperativeMatrixLoadNV = 5359, @@ -1398,6 +1466,8 @@ enum Op { OpUSubSatINTEL = 5596, OpIMul32x16INTEL = 5597, OpUMul32x16INTEL = 5598, + OpFunctionPointerINTEL = 5600, + OpFunctionPointerCallINTEL = 5601, OpDecorateString = 5632, OpDecorateStringGOOGLE = 5632, OpMemberDecorateString = 5633, @@ -1520,6 +1590,10 @@ enum Op { OpSubgroupAvcSicGetPackedSkcLumaCountThresholdINTEL = 5814, OpSubgroupAvcSicGetPackedSkcLumaSumThresholdINTEL = 5815, OpSubgroupAvcSicGetInterRawSadsINTEL = 5816, + OpLoopControlINTEL = 5887, + OpReadPipeBlockingINTEL = 5946, + OpWritePipeBlockingINTEL = 5947, + OpFPGARegINTEL = 5949, OpRayQueryGetRayTMinKHR = 6016, OpRayQueryGetRayFlagsKHR = 6017, OpRayQueryGetIntersectionTKHR = 6018, @@ -1537,6 +1611,7 @@ enum Op { OpRayQueryGetWorldRayOriginKHR = 6030, OpRayQueryGetIntersectionObjectToWorldKHR = 6031, OpRayQueryGetIntersectionWorldToObjectKHR = 6032, + OpAtomicFAddEXT = 6035, OpMax = 0x7fffffff, }; @@ -1889,13 +1964,19 @@ inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) { case OpPtrEqual: *hasResult = true; *hasResultType = true; break; case OpPtrNotEqual: *hasResult = true; *hasResultType = true; break; case OpPtrDiff: *hasResult = true; *hasResultType = true; break; + case OpTerminateInvocation: *hasResult = false; *hasResultType = false; break; case OpSubgroupBallotKHR: *hasResult = true; *hasResultType = true; break; case OpSubgroupFirstInvocationKHR: *hasResult = true; *hasResultType = true; break; case OpSubgroupAllKHR: *hasResult = true; *hasResultType = true; break; case OpSubgroupAnyKHR: *hasResult = true; *hasResultType = true; break; case OpSubgroupAllEqualKHR: *hasResult = true; *hasResultType = true; break; case OpSubgroupReadInvocationKHR: *hasResult = true; *hasResultType = true; break; - case OpTypeRayQueryProvisionalKHR: *hasResult = true; *hasResultType = false; break; + case OpTraceRayKHR: *hasResult = false; *hasResultType = false; break; + case OpExecuteCallableKHR: *hasResult = false; *hasResultType = false; break; + case OpConvertUToAccelerationStructureKHR: *hasResult = true; *hasResultType = true; break; + case OpIgnoreIntersectionKHR: *hasResult = false; *hasResultType = false; break; + case OpTerminateRayKHR: *hasResult = false; *hasResultType = false; break; + case OpTypeRayQueryKHR: *hasResult = true; *hasResultType = false; break; case OpRayQueryInitializeKHR: *hasResult = false; *hasResultType = false; break; case OpRayQueryTerminateKHR: *hasResult = false; *hasResultType = false; break; case OpRayQueryGenerateIntersectionKHR: *hasResult = false; *hasResultType = false; break; @@ -1955,6 +2036,8 @@ inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) { case OpUSubSatINTEL: *hasResult = true; *hasResultType = true; break; case OpIMul32x16INTEL: *hasResult = true; *hasResultType = true; break; case OpUMul32x16INTEL: *hasResult = true; *hasResultType = true; break; + case OpFunctionPointerINTEL: *hasResult = true; *hasResultType = true; break; + case OpFunctionPointerCallINTEL: *hasResult = true; *hasResultType = true; break; case OpDecorateString: *hasResult = false; *hasResultType = false; break; case OpMemberDecorateString: *hasResult = false; *hasResultType = false; break; case OpVmeImageINTEL: *hasResult = true; *hasResultType = true; break; @@ -2075,6 +2158,10 @@ inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) { case OpSubgroupAvcSicGetPackedSkcLumaCountThresholdINTEL: *hasResult = true; *hasResultType = true; break; case OpSubgroupAvcSicGetPackedSkcLumaSumThresholdINTEL: *hasResult = true; *hasResultType = true; break; case OpSubgroupAvcSicGetInterRawSadsINTEL: *hasResult = true; *hasResultType = true; break; + case OpLoopControlINTEL: *hasResult = false; *hasResultType = false; break; + case OpReadPipeBlockingINTEL: *hasResult = true; *hasResultType = true; break; + case OpWritePipeBlockingINTEL: *hasResult = true; *hasResultType = true; break; + case OpFPGARegINTEL: *hasResult = true; *hasResultType = true; break; case OpRayQueryGetRayTMinKHR: *hasResult = true; *hasResultType = true; break; case OpRayQueryGetRayFlagsKHR: *hasResult = true; *hasResultType = true; break; case OpRayQueryGetIntersectionTKHR: *hasResult = true; *hasResultType = true; break; @@ -2092,6 +2179,7 @@ inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) { case OpRayQueryGetWorldRayOriginKHR: *hasResult = true; *hasResultType = true; break; case OpRayQueryGetIntersectionObjectToWorldKHR: *hasResult = true; *hasResultType = true; break; case OpRayQueryGetIntersectionWorldToObjectKHR: *hasResult = true; *hasResultType = true; break; + case OpAtomicFAddEXT: *hasResult = true; *hasResultType = true; break; } } #endif /* SPV_ENABLE_UTILITY_CODE */ @@ -2107,6 +2195,7 @@ inline MemorySemanticsMask operator|(MemorySemanticsMask a, MemorySemanticsMask inline MemoryAccessMask operator|(MemoryAccessMask a, MemoryAccessMask b) { return MemoryAccessMask(unsigned(a) | unsigned(b)); } inline KernelProfilingInfoMask operator|(KernelProfilingInfoMask a, KernelProfilingInfoMask b) { return KernelProfilingInfoMask(unsigned(a) | unsigned(b)); } inline RayFlagsMask operator|(RayFlagsMask a, RayFlagsMask b) { return RayFlagsMask(unsigned(a) | unsigned(b)); } +inline FragmentShadingRateMask operator|(FragmentShadingRateMask a, FragmentShadingRateMask b) { return FragmentShadingRateMask(unsigned(a) | unsigned(b)); } } // end namespace spv diff --git a/3rdparty/spirv-cross/spirv_common.hpp b/3rdparty/spirv-cross/spirv_common.hpp index ecb840772..8e3b13739 100644 --- a/3rdparty/spirv-cross/spirv_common.hpp +++ b/3rdparty/spirv-cross/spirv_common.hpp @@ -729,7 +729,9 @@ struct SPIRBlock : IVariant Return, // Block ends with return. Unreachable, // Noop - Kill // Discard + Kill, // Discard + IgnoreIntersection, // Ray Tracing + TerminateRay // Ray Tracing }; enum Merge diff --git a/3rdparty/spirv-cross/spirv_cross.cpp b/3rdparty/spirv-cross/spirv_cross.cpp index 5c1b3ebc5..d814cd64d 100644 --- a/3rdparty/spirv-cross/spirv_cross.cpp +++ b/3rdparty/spirv-cross/spirv_cross.cpp @@ -96,7 +96,9 @@ bool Compiler::variable_storage_is_aliased(const SPIRVariable &v) bool Compiler::block_is_pure(const SPIRBlock &block) { // This is a global side effect of the function. - if (block.terminator == SPIRBlock::Kill) + if (block.terminator == SPIRBlock::Kill || + block.terminator == SPIRBlock::TerminateRay || + block.terminator == SPIRBlock::IgnoreIntersection) return false; for (auto &i : block.ops) @@ -158,11 +160,13 @@ bool Compiler::block_is_pure(const SPIRBlock &block) return false; // Ray tracing builtins are impure. - case OpReportIntersectionNV: + case OpReportIntersectionKHR: case OpIgnoreIntersectionNV: case OpTerminateRayNV: case OpTraceNV: + case OpTraceRayKHR: case OpExecuteCallableNV: + case OpExecuteCallableKHR: return false; // OpExtInst is potentially impure depending on extension, but GLSL builtins are at least pure. @@ -805,9 +809,17 @@ unordered_set Compiler::get_active_interface_variables() const InterfaceVariableAccessHandler handler(*this, variables); traverse_all_reachable_opcodes(get(ir.default_entry_point), handler); - // Make sure we preserve output variables which are only initialized, but never accessed by any code. ir.for_each_typed_id([&](uint32_t, const SPIRVariable &var) { - if (var.storage == StorageClassOutput && var.initializer != ID(0)) + if (var.storage != StorageClassOutput) + return; + if (!interface_variable_exists_in_entry_point(var.self)) + return; + + // An output variable which is just declared (but uninitialized) might be read by subsequent stages + // so we should force-enable these outputs, + // since compilation will fail if a subsequent stage attempts to read from the variable in question. + // Also, make sure we preserve output variables which are only initialized, but never accessed by any code. + if (var.initializer != ID(0) || get_execution_model() != ExecutionModelFragment) variables.insert(var.self); }); @@ -2279,16 +2291,22 @@ SPIREntryPoint &Compiler::get_entry_point() bool Compiler::interface_variable_exists_in_entry_point(uint32_t id) const { auto &var = get(id); - if (var.storage != StorageClassInput && var.storage != StorageClassOutput && - var.storage != StorageClassUniformConstant) - SPIRV_CROSS_THROW("Only Input, Output variables and Uniform constants are part of a shader linking interface."); - // This is to avoid potential problems with very old glslang versions which did - // not emit input/output interfaces properly. - // We can assume they only had a single entry point, and single entry point - // shaders could easily be assumed to use every interface variable anyways. - if (ir.entry_points.size() <= 1) - return true; + if (ir.get_spirv_version() < 0x10400) + { + if (var.storage != StorageClassInput && var.storage != StorageClassOutput && + var.storage != StorageClassUniformConstant) + SPIRV_CROSS_THROW("Only Input, Output variables and Uniform constants are part of a shader linking interface."); + + // This is to avoid potential problems with very old glslang versions which did + // not emit input/output interfaces properly. + // We can assume they only had a single entry point, and single entry point + // shaders could easily be assumed to use every interface variable anyways. + if (ir.entry_points.size() <= 1) + return true; + } + + // In SPIR-V 1.4 and later, all global resource variables must be present. auto &execution = get_entry_point(); return find(begin(execution.interface_variables), end(execution.interface_variables), VariableID(id)) != @@ -3839,23 +3857,55 @@ void Compiler::ActiveBuiltinHandler::handle_builtin(const SPIRType &type, BuiltI } } -bool Compiler::ActiveBuiltinHandler::handle(spv::Op opcode, const uint32_t *args, uint32_t length) +void Compiler::ActiveBuiltinHandler::add_if_builtin(uint32_t id, bool allow_blocks) { - const auto add_if_builtin = [&](uint32_t id) { - // Only handles variables here. - // Builtins which are part of a block are handled in AccessChain. - auto *var = compiler.maybe_get(id); - auto &decorations = compiler.ir.meta[id].decoration; - if (var && decorations.builtin) + // Only handle plain variables here. + // Builtins which are part of a block are handled in AccessChain. + // If allow_blocks is used however, this is to handle initializers of blocks, + // which implies that all members are written to. + + auto *var = compiler.maybe_get(id); + auto *m = compiler.ir.find_meta(id); + if (var && m) + { + auto &type = compiler.get(var->basetype); + auto &decorations = m->decoration; + auto &flags = type.storage == StorageClassInput ? + compiler.active_input_builtins : compiler.active_output_builtins; + if (decorations.builtin) { - auto &type = compiler.get(var->basetype); - auto &flags = - type.storage == StorageClassInput ? compiler.active_input_builtins : compiler.active_output_builtins; flags.set(decorations.builtin_type); handle_builtin(type, decorations.builtin_type, decorations.decoration_flags); } - }; + else if (allow_blocks && compiler.has_decoration(type.self, DecorationBlock)) + { + uint32_t member_count = uint32_t(type.member_types.size()); + for (uint32_t i = 0; i < member_count; i++) + { + if (compiler.has_member_decoration(type.self, i, DecorationBuiltIn)) + { + auto &member_type = compiler.get(type.member_types[i]); + BuiltIn builtin = BuiltIn(compiler.get_member_decoration(type.self, i, DecorationBuiltIn)); + flags.set(builtin); + handle_builtin(member_type, builtin, compiler.get_member_decoration_bitset(type.self, i)); + } + } + } + } +} +void Compiler::ActiveBuiltinHandler::add_if_builtin(uint32_t id) +{ + add_if_builtin(id, false); +} + +void Compiler::ActiveBuiltinHandler::add_if_builtin_or_block(uint32_t id) +{ + add_if_builtin(id, true); +} + +bool Compiler::ActiveBuiltinHandler::handle(spv::Op opcode, const uint32_t *args, uint32_t length) +{ switch (opcode) { case OpStore: @@ -3993,6 +4043,17 @@ void Compiler::update_active_builtins() clip_distance_count = 0; ActiveBuiltinHandler handler(*this); traverse_all_reachable_opcodes(get(ir.default_entry_point), handler); + + ir.for_each_typed_id([&](uint32_t, const SPIRVariable &var) { + if (var.storage != StorageClassOutput) + return; + if (!interface_variable_exists_in_entry_point(var.self)) + return; + + // Also, make sure we preserve output variables which are only initialized, but never accessed by any code. + if (var.initializer != ID(0)) + handler.add_if_builtin_or_block(var.self); + }); } // Returns whether this shader uses a builtin of the storage class diff --git a/3rdparty/spirv-cross/spirv_cross.hpp b/3rdparty/spirv-cross/spirv_cross.hpp index 47f1d7949..016d13e53 100644 --- a/3rdparty/spirv-cross/spirv_cross.hpp +++ b/3rdparty/spirv-cross/spirv_cross.hpp @@ -833,6 +833,9 @@ protected: Compiler &compiler; void handle_builtin(const SPIRType &type, spv::BuiltIn builtin, const Bitset &decoration_flags); + void add_if_builtin(uint32_t id); + void add_if_builtin_or_block(uint32_t id); + void add_if_builtin(uint32_t id, bool allow_blocks); }; bool traverse_all_reachable_opcodes(const SPIRBlock &block, OpcodeHandler &handler) const; diff --git a/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp b/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp index ff1b63fbc..a9ade42b1 100644 --- a/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp +++ b/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp @@ -235,6 +235,11 @@ bool ParsedIR::is_globally_reserved_identifier(std::string &str, bool allow_rese return is_reserved_identifier(str, false, allow_reserved_prefixes); } +uint32_t ParsedIR::get_spirv_version() const +{ + return spirv[1]; +} + static string make_unreserved_identifier(const string &name) { if (is_reserved_prefix(name)) diff --git a/3rdparty/spirv-cross/spirv_cross_parsed_ir.hpp b/3rdparty/spirv-cross/spirv_cross_parsed_ir.hpp index ee202d2fa..ca37a9bcd 100644 --- a/3rdparty/spirv-cross/spirv_cross_parsed_ir.hpp +++ b/3rdparty/spirv-cross/spirv_cross_parsed_ir.hpp @@ -222,6 +222,8 @@ public: static void sanitize_identifier(std::string &str, bool member, bool allow_reserved_prefixes); static bool is_globally_reserved_identifier(std::string &str, bool allow_reserved_prefixes); + uint32_t get_spirv_version() const; + private: template T &get(uint32_t id) diff --git a/3rdparty/spirv-cross/spirv_glsl.cpp b/3rdparty/spirv-cross/spirv_glsl.cpp index ca68e87f5..7a740711c 100644 --- a/3rdparty/spirv-cross/spirv_glsl.cpp +++ b/3rdparty/spirv-cross/spirv_glsl.cpp @@ -427,15 +427,37 @@ void CompilerGLSL::find_static_extensions() require_extension_internal("GL_ARB_tessellation_shader"); break; - case ExecutionModelRayGenerationNV: - case ExecutionModelIntersectionNV: - case ExecutionModelAnyHitNV: - case ExecutionModelClosestHitNV: - case ExecutionModelMissNV: - case ExecutionModelCallableNV: + case ExecutionModelRayGenerationKHR: + case ExecutionModelIntersectionKHR: + case ExecutionModelAnyHitKHR: + case ExecutionModelClosestHitKHR: + case ExecutionModelMissKHR: + case ExecutionModelCallableKHR: + // NV enums are aliases. if (options.es || options.version < 460) SPIRV_CROSS_THROW("Ray tracing shaders require non-es profile with version 460 or above."); - require_extension_internal("GL_NV_ray_tracing"); + if (!options.vulkan_semantics) + SPIRV_CROSS_THROW("Ray tracing requires Vulkan semantics."); + + // Need to figure out if we should target KHR or NV extension based on capabilities. + for (auto &cap : ir.declared_capabilities) + { + if (cap == CapabilityRayTracingKHR || cap == CapabilityRayQueryKHR) + { + ray_tracing_is_khr = true; + break; + } + } + + if (ray_tracing_is_khr) + { + // In KHR ray tracing we pass payloads by pointer instead of location, + // so make sure we assign locations properly. + ray_tracing_khr_fixup_locations(); + require_extension_internal("GL_EXT_ray_tracing"); + } + else + require_extension_internal("GL_NV_ray_tracing"); break; default: @@ -512,6 +534,18 @@ void CompilerGLSL::find_static_extensions() } } +void CompilerGLSL::ray_tracing_khr_fixup_locations() +{ + uint32_t location = 0; + ir.for_each_typed_id([&](uint32_t, SPIRVariable &var) { + if (var.storage != StorageClassRayPayloadKHR && var.storage != StorageClassCallableDataKHR) + return; + if (!interface_variable_exists_in_entry_point(var.self)) + return; + set_decoration(var.self, DecorationLocation, location++); + }); +} + string CompilerGLSL::compile() { ir.fixup_reserved_names(); @@ -1620,8 +1654,8 @@ string CompilerGLSL::layout_for_variable(const SPIRVariable &var) if (options.vulkan_semantics && var.storage == StorageClassPushConstant) attr.push_back("push_constant"); - else if (var.storage == StorageClassShaderRecordBufferNV) - attr.push_back("shaderRecordNV"); + else if (var.storage == StorageClassShaderRecordBufferKHR) + attr.push_back(ray_tracing_is_khr ? "shaderRecordEXT" : "shaderRecordNV"); if (flags.get(DecorationRowMajor)) attr.push_back("row_major"); @@ -1777,14 +1811,14 @@ string CompilerGLSL::layout_for_variable(const SPIRVariable &var) // Do not emit set = decoration in regular GLSL output, but // we need to preserve it in Vulkan GLSL mode. - if (var.storage != StorageClassPushConstant && var.storage != StorageClassShaderRecordBufferNV) + if (var.storage != StorageClassPushConstant && var.storage != StorageClassShaderRecordBufferKHR) { if (flags.get(DecorationDescriptorSet) && options.vulkan_semantics) attr.push_back(join("set = ", get_decoration(var.self, DecorationDescriptorSet))); } bool push_constant_block = options.vulkan_semantics && var.storage == StorageClassPushConstant; - bool ssbo_block = var.storage == StorageClassStorageBuffer || var.storage == StorageClassShaderRecordBufferNV || + bool ssbo_block = var.storage == StorageClassStorageBuffer || var.storage == StorageClassShaderRecordBufferKHR || (var.storage == StorageClassUniform && typeflags.get(DecorationBufferBlock)); bool emulated_ubo = var.storage == StorageClassPushConstant && options.emit_push_constant_as_uniform_buffer; bool ubo_block = var.storage == StorageClassUniform && typeflags.get(DecorationBlock); @@ -1806,7 +1840,7 @@ string CompilerGLSL::layout_for_variable(const SPIRVariable &var) if (!can_use_buffer_blocks && var.storage == StorageClassUniform) can_use_binding = false; - if (var.storage == StorageClassShaderRecordBufferNV) + if (var.storage == StorageClassShaderRecordBufferKHR) can_use_binding = false; if (can_use_binding && flags.get(DecorationBinding)) @@ -2086,7 +2120,7 @@ void CompilerGLSL::emit_buffer_block_native(const SPIRVariable &var) auto &type = get(var.basetype); Bitset flags = ir.get_buffer_block_flags(var); - bool ssbo = var.storage == StorageClassStorageBuffer || var.storage == StorageClassShaderRecordBufferNV || + bool ssbo = var.storage == StorageClassStorageBuffer || var.storage == StorageClassShaderRecordBufferKHR || ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); bool is_restrict = ssbo && flags.get(DecorationRestrict); bool is_writeonly = ssbo && flags.get(DecorationNonReadable); @@ -2201,25 +2235,25 @@ const char *CompilerGLSL::to_storage_qualifiers_glsl(const SPIRVariable &var) { return "uniform "; } - else if (var.storage == StorageClassRayPayloadNV) + else if (var.storage == StorageClassRayPayloadKHR) { - return "rayPayloadNV "; + return ray_tracing_is_khr ? "rayPayloadEXT " : "rayPayloadNV "; } - else if (var.storage == StorageClassIncomingRayPayloadNV) + else if (var.storage == StorageClassIncomingRayPayloadKHR) { - return "rayPayloadInNV "; + return ray_tracing_is_khr ? "rayPayloadInEXT " : "rayPayloadInNV "; } - else if (var.storage == StorageClassHitAttributeNV) + else if (var.storage == StorageClassHitAttributeKHR) { - return "hitAttributeNV "; + return ray_tracing_is_khr ? "hitAttributeEXT " : "hitAttributeNV "; } - else if (var.storage == StorageClassCallableDataNV) + else if (var.storage == StorageClassCallableDataKHR) { - return "callableDataNV "; + return ray_tracing_is_khr ? "callableDataEXT " : "callableDataNV "; } - else if (var.storage == StorageClassIncomingCallableDataNV) + else if (var.storage == StorageClassIncomingCallableDataKHR) { - return "callableDataInNV "; + return ray_tracing_is_khr ? "callableDataInEXT " : "callableDataInNV "; } return ""; @@ -2363,6 +2397,9 @@ void CompilerGLSL::emit_interface_block(const SPIRVariable &var) require_extension_internal("GL_EXT_shader_io_blocks"); } + // Workaround to make sure we can emit "patch in/out" correctly. + fixup_io_block_patch_qualifiers(var); + // Block names should never alias. auto block_name = to_name(type.self, false); @@ -2384,7 +2421,8 @@ void CompilerGLSL::emit_interface_block(const SPIRVariable &var) // Instance names cannot alias block names. resource_names.insert(block_name); - statement(layout_for_variable(var), qual, block_name); + bool is_patch = has_decoration(var.self, DecorationPatch); + statement(layout_for_variable(var), (is_patch ? "patch " : ""), qual, block_name); begin_scope(); type.member_name_cache.clear(); @@ -2441,14 +2479,6 @@ void CompilerGLSL::emit_interface_block(const SPIRVariable &var) swap(type.array.back(), old_array_size); swap(type.array_size_literal.back(), old_array_size_literal); } - - // If a StorageClassOutput variable has an initializer, we need to initialize it in main(). - if (var.storage == StorageClassOutput && var.initializer) - { - auto &entry_func = this->get(ir.default_entry_point); - entry_func.fixup_hooks_in.push_back( - [&]() { statement(to_name(var.self), " = ", to_expression(var.initializer), ";"); }); - } } } } @@ -2551,7 +2581,17 @@ void CompilerGLSL::replace_illegal_names(const unordered_set &keywords) return; auto &m = meta->decoration; - if (m.alias.compare(0, 3, "gl_") == 0 || keywords.find(m.alias) != end(keywords)) + if (keywords.find(m.alias) != end(keywords)) + m.alias = join("_", m.alias); + }); + + ir.for_each_typed_id([&](uint32_t, const SPIRFunction &func) { + auto *meta = ir.find_meta(func.self); + if (!meta) + return; + + auto &m = meta->decoration; + if (keywords.find(m.alias) != end(keywords)) m.alias = join("_", m.alias); }); @@ -2561,11 +2601,11 @@ void CompilerGLSL::replace_illegal_names(const unordered_set &keywords) return; auto &m = meta->decoration; - if (m.alias.compare(0, 3, "gl_") == 0 || keywords.find(m.alias) != end(keywords)) + if (keywords.find(m.alias) != end(keywords)) m.alias = join("_", m.alias); for (auto &memb : meta->members) - if (memb.alias.compare(0, 3, "gl_") == 0 || keywords.find(memb.alias) != end(keywords)) + if (keywords.find(memb.alias) != end(keywords)) memb.alias = join("_", memb.alias); }); } @@ -2795,6 +2835,13 @@ 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 && + (clip_distance_count || cull_distance_count)) + { + should_force = true; + } + return should_force; } @@ -3205,8 +3252,8 @@ void CompilerGLSL::emit_resources() // Special case, ray payload and hit attribute blocks are not really blocks, just regular structs. if (type->basetype == SPIRType::Struct && type->pointer && has_decoration(type->self, DecorationBlock) && - (type->storage == StorageClassRayPayloadNV || type->storage == StorageClassIncomingRayPayloadNV || - type->storage == StorageClassHitAttributeNV)) + (type->storage == StorageClassRayPayloadKHR || type->storage == StorageClassIncomingRayPayloadKHR || + type->storage == StorageClassHitAttributeKHR)) { type = &get(type->parent_type); is_natural_struct = true; @@ -3282,7 +3329,7 @@ void CompilerGLSL::emit_resources() auto &type = this->get(var.basetype); bool is_block_storage = type.storage == StorageClassStorageBuffer || type.storage == StorageClassUniform || - type.storage == StorageClassShaderRecordBufferNV; + type.storage == StorageClassShaderRecordBufferKHR; bool has_block_flags = ir.meta[type.self].decoration.decoration_flags.get(DecorationBlock) || ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); @@ -3322,9 +3369,9 @@ void CompilerGLSL::emit_resources() if (var.storage != StorageClassFunction && type.pointer && (type.storage == StorageClassUniformConstant || type.storage == StorageClassAtomicCounter || - type.storage == StorageClassRayPayloadNV || type.storage == StorageClassIncomingRayPayloadNV || - type.storage == StorageClassCallableDataNV || type.storage == StorageClassIncomingCallableDataNV || - type.storage == StorageClassHitAttributeNV) && + type.storage == StorageClassRayPayloadKHR || type.storage == StorageClassIncomingRayPayloadKHR || + type.storage == StorageClassCallableDataKHR || type.storage == StorageClassIncomingCallableDataKHR || + type.storage == StorageClassHitAttributeKHR) && !is_hidden_variable(var)) { emit_uniform(var); @@ -3421,6 +3468,10 @@ void CompilerGLSL::emit_resources() emitted = true; } } + else if (var.initializer && maybe_get(var.initializer) != nullptr) + { + emit_output_variable_initializer(var); + } } if (emitted) @@ -3429,6 +3480,140 @@ void CompilerGLSL::emit_resources() declare_undefined_values(); } +void CompilerGLSL::emit_output_variable_initializer(const SPIRVariable &var) +{ + // If a StorageClassOutput variable has an initializer, we need to initialize it in main(). + auto &entry_func = this->get(ir.default_entry_point); + auto &type = get(var.basetype); + bool is_patch = has_decoration(var.self, DecorationPatch); + bool is_block = has_decoration(type.self, DecorationBlock); + bool is_control_point = get_execution_model() == ExecutionModelTessellationControl && !is_patch; + + if (is_block) + { + uint32_t member_count = uint32_t(type.member_types.size()); + bool type_is_array = type.array.size() == 1; + uint32_t array_size = 1; + if (type_is_array) + array_size = to_array_size_literal(type); + uint32_t iteration_count = is_control_point ? 1 : array_size; + + // If the initializer is a block, we must initialize each block member one at a time. + for (uint32_t i = 0; i < member_count; i++) + { + // These outputs might not have been properly declared, so don't initialize them in that case. + if (has_member_decoration(type.self, i, DecorationBuiltIn)) + { + if (get_member_decoration(type.self, i, DecorationBuiltIn) == BuiltInCullDistance && + !cull_distance_count) + continue; + + if (get_member_decoration(type.self, i, DecorationBuiltIn) == BuiltInClipDistance && + !clip_distance_count) + continue; + } + + // We need to build a per-member array first, essentially transposing from AoS to SoA. + // This code path hits when we have an array of blocks. + string lut_name; + if (type_is_array) + { + lut_name = join("_", var.self, "_", i, "_init"); + uint32_t member_type_id = get(var.basetype).member_types[i]; + auto &member_type = get(member_type_id); + auto array_type = member_type; + array_type.parent_type = member_type_id; + array_type.array.push_back(array_size); + array_type.array_size_literal.push_back(true); + + SmallVector exprs; + exprs.reserve(array_size); + auto &c = get(var.initializer); + for (uint32_t j = 0; j < array_size; j++) + exprs.push_back(to_expression(get(c.subconstants[j]).subconstants[i])); + statement("const ", type_to_glsl(array_type), " ", lut_name, type_to_array_glsl(array_type), " = ", + type_to_glsl_constructor(array_type), "(", merge(exprs, ", "), ");"); + } + + for (uint32_t j = 0; j < iteration_count; j++) + { + entry_func.fixup_hooks_in.push_back([=, &var]() { + AccessChainMeta meta; + auto &c = this->get(var.initializer); + + uint32_t invocation_id = 0; + uint32_t member_index_id = 0; + if (is_control_point) + { + uint32_t ids = ir.increase_bound_by(3); + SPIRType uint_type; + uint_type.basetype = SPIRType::UInt; + uint_type.width = 32; + set(ids, uint_type); + set(ids + 1, builtin_to_glsl(BuiltInInvocationId, StorageClassInput), ids, true); + set(ids + 2, ids, i, false); + invocation_id = ids + 1; + member_index_id = ids + 2; + } + + if (is_patch) + { + statement("if (gl_InvocationID == 0)"); + begin_scope(); + } + + if (type_is_array && !is_control_point) + { + uint32_t indices[2] = { j, i }; + auto chain = access_chain_internal(var.self, indices, 2, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &meta); + statement(chain, " = ", lut_name, "[", j, "];"); + } + else if (is_control_point) + { + uint32_t indices[2] = { invocation_id, member_index_id }; + auto chain = access_chain_internal(var.self, indices, 2, 0, &meta); + statement(chain, " = ", lut_name, "[", builtin_to_glsl(BuiltInInvocationId, StorageClassInput), "];"); + } + else + { + auto chain = + access_chain_internal(var.self, &i, 1, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &meta); + statement(chain, " = ", to_expression(c.subconstants[i]), ";"); + } + + if (is_patch) + end_scope(); + }); + } + } + } + else if (is_control_point) + { + auto lut_name = join("_", var.self, "_init"); + statement("const ", type_to_glsl(type), " ", lut_name, type_to_array_glsl(type), + " = ", to_expression(var.initializer), ";"); + entry_func.fixup_hooks_in.push_back([&, lut_name]() { + statement(to_expression(var.self), "[gl_InvocationID] = ", lut_name, "[gl_InvocationID];"); + }); + } + else + { + auto lut_name = join("_", var.self, "_init"); + statement("const ", type_to_glsl(type), " ", lut_name, + type_to_array_glsl(type), " = ", to_expression(var.initializer), ";"); + entry_func.fixup_hooks_in.push_back([&, lut_name, is_patch]() { + if (is_patch) + { + statement("if (gl_InvocationID == 0)"); + begin_scope(); + } + statement(to_expression(var.self), " = ", lut_name, ";"); + if (is_patch) + end_scope(); + }); + } +} + void CompilerGLSL::emit_extension_workarounds(spv::ExecutionModel model) { static const char *workaround_types[] = { "int", "ivec2", "ivec3", "ivec4", "uint", "uvec2", "uvec3", "uvec4", @@ -7987,34 +8172,35 @@ string CompilerGLSL::builtin_to_glsl(BuiltIn builtin, StorageClass storage) request_subgroup_feature(ShaderSubgroupSupportHelper::SubgroupMask); return "gl_SubgroupLtMask"; - case BuiltInLaunchIdNV: - return "gl_LaunchIDNV"; - case BuiltInLaunchSizeNV: - return "gl_LaunchSizeNV"; - case BuiltInWorldRayOriginNV: - return "gl_WorldRayOriginNV"; - case BuiltInWorldRayDirectionNV: - return "gl_WorldRayDirectionNV"; - case BuiltInObjectRayOriginNV: - return "gl_ObjectRayOriginNV"; - case BuiltInObjectRayDirectionNV: - return "gl_ObjectRayDirectionNV"; - case BuiltInRayTminNV: - return "gl_RayTminNV"; - case BuiltInRayTmaxNV: - return "gl_RayTmaxNV"; - case BuiltInInstanceCustomIndexNV: - return "gl_InstanceCustomIndexNV"; - case BuiltInObjectToWorldNV: - return "gl_ObjectToWorldNV"; - case BuiltInWorldToObjectNV: - return "gl_WorldToObjectNV"; + case BuiltInLaunchIdKHR: + return ray_tracing_is_khr ? "gl_LaunchIDEXT" : "gl_LaunchIDNV"; + case BuiltInLaunchSizeKHR: + return ray_tracing_is_khr ? "gl_LaunchSizeEXT" : "gl_LaunchSizeNV"; + case BuiltInWorldRayOriginKHR: + return ray_tracing_is_khr ? "gl_WorldRayOriginEXT" : "gl_WorldRayOriginNV"; + case BuiltInWorldRayDirectionKHR: + return ray_tracing_is_khr ? "gl_WorldRayDirectionEXT" : "gl_WorldRayDirectionNV"; + case BuiltInObjectRayOriginKHR: + return ray_tracing_is_khr ? "gl_ObjectRayOriginEXT" : "gl_ObjectRayOriginNV"; + case BuiltInObjectRayDirectionKHR: + return ray_tracing_is_khr ? "gl_ObjectRayDirectionEXT" : "gl_ObjectRayDirectionNV"; + case BuiltInRayTminKHR: + return ray_tracing_is_khr ? "gl_RayTminEXT" : "gl_RayTminNV"; + case BuiltInRayTmaxKHR: + return ray_tracing_is_khr ? "gl_RayTmaxEXT" : "gl_RayTmaxNV"; + case BuiltInInstanceCustomIndexKHR: + return ray_tracing_is_khr ? "gl_InstanceCustomIndexEXT" : "gl_InstanceCustomIndexNV"; + case BuiltInObjectToWorldKHR: + return ray_tracing_is_khr ? "gl_ObjectToWorldEXT" : "gl_ObjectToWorldNV"; + case BuiltInWorldToObjectKHR: + return ray_tracing_is_khr ? "gl_WorldToObjectEXT" : "gl_WorldToObjectNV"; case BuiltInHitTNV: + // gl_HitTEXT is an alias of RayTMax in KHR. return "gl_HitTNV"; - case BuiltInHitKindNV: - return "gl_HitKindNV"; - case BuiltInIncomingRayFlagsNV: - return "gl_IncomingRayFlagsNV"; + case BuiltInHitKindKHR: + return ray_tracing_is_khr ? "gl_HitKindEXT" : "gl_HitKindNV"; + case BuiltInIncomingRayFlagsKHR: + return ray_tracing_is_khr ? "gl_IncomingRayFlagsEXT" : "gl_IncomingRayFlagsNV"; case BuiltInBaryCoordNV: { @@ -12010,15 +12196,22 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) break; } - case OpReportIntersectionNV: - statement("reportIntersectionNV(", to_expression(ops[0]), ", ", to_expression(ops[1]), ");"); + case OpReportIntersectionKHR: + // NV is same opcode. + forced_temporaries.insert(ops[1]); + if (ray_tracing_is_khr) + GLSL_BFOP(reportIntersectionEXT); + else + GLSL_BFOP(reportIntersectionNV); flush_control_dependent_expressions(current_emitting_block->self); break; case OpIgnoreIntersectionNV: + // KHR variant is a terminator. statement("ignoreIntersectionNV();"); flush_control_dependent_expressions(current_emitting_block->self); break; case OpTerminateRayNV: + // KHR variant is a terminator. statement("terminateRayNV();"); flush_control_dependent_expressions(current_emitting_block->self); break; @@ -12029,10 +12222,29 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) to_expression(ops[9]), ", ", to_expression(ops[10]), ");"); flush_control_dependent_expressions(current_emitting_block->self); break; + case OpTraceRayKHR: + if (!has_decoration(ops[10], DecorationLocation)) + SPIRV_CROSS_THROW("A memory declaration object must be used in TraceRayKHR."); + statement("traceRayEXT(", to_expression(ops[0]), ", ", to_expression(ops[1]), ", ", to_expression(ops[2]), ", ", + to_expression(ops[3]), ", ", to_expression(ops[4]), ", ", to_expression(ops[5]), ", ", + to_expression(ops[6]), ", ", to_expression(ops[7]), ", ", to_expression(ops[8]), ", ", + to_expression(ops[9]), ", ", get_decoration(ops[10], DecorationLocation), ");"); + flush_control_dependent_expressions(current_emitting_block->self); + break; case OpExecuteCallableNV: statement("executeCallableNV(", to_expression(ops[0]), ", ", to_expression(ops[1]), ");"); flush_control_dependent_expressions(current_emitting_block->self); break; + case OpExecuteCallableKHR: + if (!has_decoration(ops[1], DecorationLocation)) + SPIRV_CROSS_THROW("A memory declaration object must be used in ExecuteCallableKHR."); + statement("executeCallableEXT(", to_expression(ops[0]), ", ", get_decoration(ops[1], DecorationLocation), ");"); + flush_control_dependent_expressions(current_emitting_block->self); + break; + + case OpConvertUToAccelerationStructureKHR: + GLSL_UFOP(accelerationStructureEXT); + break; case OpConvertUToPtr: { @@ -12409,6 +12621,31 @@ const char *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) +{ + // 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. + auto &type = get(var.basetype); + if (has_decoration(type.self, DecorationBlock)) + { + uint32_t member_count = uint32_t(type.member_types.size()); + for (uint32_t i = 0; i < member_count; i++) + { + if (has_member_decoration(type.self, i, DecorationPatch)) + { + set_decoration(var.self, DecorationPatch); + break; + } + } + + if (has_decoration(var.self, DecorationPatch)) + for (uint32_t i = 0; i < member_count; i++) + unset_member_decoration(type.self, i, DecorationPatch); + } +} + string CompilerGLSL::to_qualifiers_glsl(uint32_t id) { auto &flags = ir.meta[id].decoration.decoration_flags; @@ -12802,7 +13039,7 @@ string CompilerGLSL::type_to_glsl(const SPIRType &type, uint32_t id) return comparison_ids.count(id) ? "samplerShadow" : "sampler"; case SPIRType::AccelerationStructure: - return "accelerationStructureNV"; + return ray_tracing_is_khr ? "accelerationStructureEXT" : "accelerationStructureNV"; case SPIRType::Void: return "void"; @@ -14441,6 +14678,14 @@ void CompilerGLSL::emit_block_chain(SPIRBlock &block) emit_next_block = false; break; + case SPIRBlock::IgnoreIntersection: + statement("ignoreIntersectionEXT;"); + break; + + case SPIRBlock::TerminateRay: + statement("terminateRayEXT;"); + break; + default: SPIRV_CROSS_THROW("Unimplemented block terminator."); } diff --git a/3rdparty/spirv-cross/spirv_glsl.hpp b/3rdparty/spirv-cross/spirv_glsl.hpp index c382db678..a32edc0a5 100644 --- a/3rdparty/spirv-cross/spirv_glsl.hpp +++ b/3rdparty/spirv-cross/spirv_glsl.hpp @@ -711,6 +711,8 @@ 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 emit_output_variable_initializer(const SPIRVariable &var); const char *to_precision_qualifiers_glsl(uint32_t id); virtual const char *to_storage_qualifiers_glsl(const SPIRVariable &var); const char *flags_to_qualifiers_glsl(const SPIRType &type, const Bitset &flags); @@ -815,6 +817,8 @@ protected: bool requires_transpose_2x2 = false; bool requires_transpose_3x3 = false; bool requires_transpose_4x4 = false; + bool ray_tracing_is_khr = false; + void ray_tracing_khr_fixup_locations(); bool args_will_forward(uint32_t id, const uint32_t *args, uint32_t num_args, bool pure); void register_call_out_argument(uint32_t id); diff --git a/3rdparty/spirv-cross/spirv_hlsl.cpp b/3rdparty/spirv-cross/spirv_hlsl.cpp index f79e2be7a..24eae9985 100644 --- a/3rdparty/spirv-cross/spirv_hlsl.cpp +++ b/3rdparty/spirv-cross/spirv_hlsl.cpp @@ -839,6 +839,21 @@ std::string CompilerHLSL::to_semantic(uint32_t location, ExecutionModel em, Stor return join("TEXCOORD", location); } +std::string CompilerHLSL::to_initializer_expression(const SPIRVariable &var) +{ + // We cannot emit static const initializer for block constants for practical reasons, + // so just inline the initializer. + // FIXME: There is a theoretical problem here if someone tries to composite extract + // into this initializer since we don't declare it properly, but that is somewhat non-sensical. + auto &type = get(var.basetype); + bool is_block = has_decoration(type.self, DecorationBlock); + auto *c = maybe_get(var.initializer); + if (is_block && c) + return constant_expression(*c); + else + return CompilerGLSL::to_initializer_expression(var); +} + void CompilerHLSL::emit_io_block(const SPIRVariable &var) { auto &execution = get_entry_point(); @@ -1008,12 +1023,43 @@ void CompilerHLSL::emit_builtin_variables() bool need_base_vertex_info = false; + std::unordered_map builtin_to_initializer; + ir.for_each_typed_id([&](uint32_t, SPIRVariable &var) { + if (!is_builtin_variable(var) || var.storage != StorageClassOutput || !var.initializer) + return; + + auto *c = this->maybe_get(var.initializer); + if (!c) + return; + + auto &type = this->get(var.basetype); + if (type.basetype == SPIRType::Struct) + { + uint32_t member_count = uint32_t(type.member_types.size()); + for (uint32_t i = 0; i < member_count; i++) + { + if (has_member_decoration(type.self, i, DecorationBuiltIn)) + { + builtin_to_initializer[get_member_decoration(type.self, i, DecorationBuiltIn)] = + c->subconstants[i]; + } + } + } + else if (has_decoration(var.self, DecorationBuiltIn)) + builtin_to_initializer[get_decoration(var.self, DecorationBuiltIn)] = var.initializer; + }); + // Emit global variables for the interface variables which are statically used by the shader. builtins.for_each_bit([&](uint32_t i) { const char *type = nullptr; auto builtin = static_cast(i); uint32_t array_size = 0; + string init_expr; + auto init_itr = builtin_to_initializer.find(builtin); + if (init_itr != builtin_to_initializer.end()) + init_expr = join(" = ", to_expression(init_itr->second)); + switch (builtin) { case BuiltInFragCoord: @@ -1106,16 +1152,16 @@ void CompilerHLSL::emit_builtin_variables() if (type) { if (array_size) - statement("static ", type, " ", builtin_to_glsl(builtin, storage), "[", array_size, "];"); + statement("static ", type, " ", builtin_to_glsl(builtin, storage), "[", array_size, "]", init_expr, ";"); else - statement("static ", type, " ", builtin_to_glsl(builtin, storage), ";"); + statement("static ", type, " ", builtin_to_glsl(builtin, storage), init_expr, ";"); } // SampleMask can be both in and out with sample builtin, in this case we have already // declared the input variable and we need to add the output one now. if (builtin == BuiltInSampleMask && storage == StorageClassInput && this->active_output_builtins.get(i)) { - statement("static ", type, " ", this->builtin_to_glsl(builtin, StorageClassOutput), ";"); + statement("static ", type, " ", this->builtin_to_glsl(builtin, StorageClassOutput), init_expr, ";"); } }); @@ -1141,7 +1187,11 @@ void CompilerHLSL::emit_composite_constants() return; auto &type = this->get(c.constant_type); - if (type.basetype == SPIRType::Struct || !type.array.empty()) + + // Cannot declare block type constants here. + // We do not have the struct type yet. + bool is_block = has_decoration(type.self, DecorationBlock); + if (!is_block && (type.basetype == SPIRType::Struct || !type.array.empty())) { auto name = to_name(c.self); statement("static const ", variable_decl(type, name), " = ", constant_expression(c), ";"); diff --git a/3rdparty/spirv-cross/spirv_hlsl.hpp b/3rdparty/spirv-cross/spirv_hlsl.hpp index 84e75e913..4b481ddda 100644 --- a/3rdparty/spirv-cross/spirv_hlsl.hpp +++ b/3rdparty/spirv-cross/spirv_hlsl.hpp @@ -244,6 +244,7 @@ private: std::string to_resource_binding(const SPIRVariable &var); std::string to_resource_binding_sampler(const SPIRVariable &var); std::string to_resource_register(HLSLBindingFlagBits flag, char space, uint32_t binding, uint32_t set); + std::string to_initializer_expression(const SPIRVariable &var) override; void emit_sampled_image_op(uint32_t result_type, uint32_t result_id, uint32_t image_id, uint32_t samp_id) override; void emit_access_chain(const Instruction &instruction); void emit_load(const Instruction &instruction); diff --git a/3rdparty/spirv-cross/spirv_msl.cpp b/3rdparty/spirv-cross/spirv_msl.cpp index b0709d590..a66945700 100644 --- a/3rdparty/spirv-cross/spirv_msl.cpp +++ b/3rdparty/spirv-cross/spirv_msl.cpp @@ -198,9 +198,12 @@ void CompilerMSL::build_implicit_builtins() bool has_workgroup_size = false; uint32_t workgroup_id_type = 0; - // FIXME: Investigate the fact that there are no checks for the entry point interface variables. ir.for_each_typed_id([&](uint32_t, SPIRVariable &var) { - if (!ir.meta[var.self].decoration.builtin) + if (var.storage != StorageClassInput && var.storage != StorageClassOutput) + return; + if (!interface_variable_exists_in_entry_point(var.self)) + return; + if (!has_decoration(var.self, DecorationBuiltIn)) return; BuiltIn builtin = ir.meta[var.self].decoration.builtin_type; @@ -1180,6 +1183,7 @@ void CompilerMSL::emit_entry_point_declarations() string CompilerMSL::compile() { + replace_illegal_entry_point_names(); ir.fixup_reserved_names(); // Do not deal with GLES-isms like precision, older extensions and such. @@ -1229,6 +1233,7 @@ string CompilerMSL::compile() fixup_type_alias(); replace_illegal_names(); + sync_entry_point_aliases_and_names(); build_function_control_flow_graphs_and_analyze(); update_active_builtins(); @@ -1992,8 +1997,22 @@ void CompilerMSL::add_plain_variable_to_interface_block(StorageClass storage, co } else { - entry_func.fixup_hooks_in.push_back( - [=, &var]() { statement(qual_var_name, " = ", to_expression(var.initializer), ";"); }); + if (meta.strip_array) + { + entry_func.fixup_hooks_in.push_back([=, &var]() { + uint32_t index = get_extended_decoration(var.self, SPIRVCrossDecorationInterfaceMemberIndex); + statement(to_expression(stage_out_ptr_var_id), "[", + builtin_to_glsl(BuiltInInvocationId, StorageClassInput), "].", + to_member_name(ib_type, index), " = ", to_expression(var.initializer), "[", + builtin_to_glsl(BuiltInInvocationId, StorageClassInput), "];"); + }); + } + else + { + entry_func.fixup_hooks_in.push_back([=, &var]() { + statement(qual_var_name, " = ", to_expression(var.initializer), ";"); + }); + } } } @@ -2510,6 +2529,8 @@ void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass stor qual_var_name += ".interpolate_at_center()"; } + bool flatten_stage_out = false; + if (is_builtin && !meta.strip_array) { // For the builtin gl_PerVertex, we cannot treat it as a block anyways, @@ -2528,6 +2549,7 @@ void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass stor break; case StorageClassOutput: + flatten_stage_out = true; entry_func.fixup_hooks_out.push_back([=, &var, &var_type]() { statement(qual_var_name, " = ", to_name(var.self), ".", to_member_name(var_type, mbr_idx), ";"); }); @@ -2596,6 +2618,35 @@ void CompilerMSL::add_plain_member_variable_to_interface_block(StorageClass stor qual_pos_var_name = qual_var_name; } + const SPIRConstant *c = nullptr; + if (!flatten_stage_out && var.storage == StorageClassOutput && + var.initializer != ID(0) && (c = maybe_get(var.initializer))) + { + if (meta.strip_array) + { + entry_func.fixup_hooks_in.push_back([=, &var]() { + auto &type = this->get(var.basetype); + uint32_t index = get_extended_decoration(var.self, SPIRVCrossDecorationInterfaceMemberIndex); + index += mbr_idx; + + AccessChainMeta chain_meta; + auto constant_chain = access_chain_internal(var.initializer, &builtin_invocation_id_id, 1, 0, &chain_meta); + + statement(to_expression(stage_out_ptr_var_id), "[", + builtin_to_glsl(BuiltInInvocationId, StorageClassInput), "].", + to_member_name(ib_type, index), " = ", + constant_chain, ".", to_member_name(type, mbr_idx), ";"); + }); + } + else + { + entry_func.fixup_hooks_in.push_back([=]() { + statement(qual_var_name, " = ", constant_expression( + this->get(c->subconstants[mbr_idx])), ";"); + }); + } + } + if (storage != StorageClassInput || !pull_model_inputs.count(var.self)) { // Copy interpolation decorations if needed @@ -2806,7 +2857,7 @@ void CompilerMSL::add_variable_to_interface_block(StorageClass storage, const st add_tess_level_input_to_interface_block(ib_var_ref, ib_type, var); } else if (var_type.basetype == SPIRType::Boolean || var_type.basetype == SPIRType::Char || - type_is_integral(var_type) || type_is_floating_point(var_type) || var_type.basetype == SPIRType::Boolean) + type_is_integral(var_type) || type_is_floating_point(var_type)) { if (!is_builtin || has_active_builtin(builtin, storage)) { @@ -6102,6 +6153,15 @@ void CompilerMSL::emit_specialization_constants_and_structs() mark_scalar_layout_structs(type); }); + bool builtin_block_type_is_required = false; + // Very special case. If gl_PerVertex is initialized as an array (tessellation) + // we have to potentially emit the gl_PerVertex struct type so that we can emit a constant LUT. + ir.for_each_typed_id([&](uint32_t, SPIRConstant &c) { + auto &type = this->get(c.constant_type); + if (is_array(type) && has_decoration(type.self, DecorationBlock) && is_builtin_type(type)) + builtin_block_type_is_required = true; + }); + // Very particular use of the soft loop lock. // align_struct may need to create custom types on the fly, but we don't care about // these types for purpose of iterating over them in ir.ids_for_type and friends. @@ -6186,7 +6246,7 @@ void CompilerMSL::emit_specialization_constants_and_structs() has_decoration(type.self, DecorationBlock) || has_decoration(type.self, DecorationBufferBlock); bool is_builtin_block = is_block && is_builtin_type(type); - bool is_declarable_struct = is_struct && !is_builtin_block; + bool is_declarable_struct = is_struct && (!is_builtin_block || builtin_block_type_is_required); // We'll declare this later. if (stage_out_var_id && get_stage_out_struct_type().self == type_id) @@ -9849,7 +9909,14 @@ string CompilerMSL::to_struct_member(const SPIRType &type, uint32_t member_type_ physical_type.basetype != SPIRType::SampledImage) { BuiltIn builtin = BuiltInMax; - if (is_member_builtin(type, index, &builtin)) + + // Special handling. In [[stage_out]] or [[stage_in]] blocks, + // we need flat arrays, but if we're somehow declaring gl_PerVertex for constant array reasons, we want + // template array types to be declared. + bool is_ib_in_out = + ((stage_out_var_id && get_stage_out_struct_type().self == type.self) || + (stage_in_var_id && get_stage_in_struct_type().self == type.self)); + if (is_ib_in_out && is_member_builtin(type, index, &builtin)) is_using_builtin_array = true; array_type = type_to_array_glsl(physical_type); } @@ -10522,11 +10589,14 @@ void CompilerMSL::entry_point_args_builtin(string &ep_args) // Builtin variables SmallVector, 8> active_builtins; ir.for_each_typed_id([&](uint32_t var_id, SPIRVariable &var) { + if (var.storage != StorageClassInput) + return; + auto bi_type = BuiltIn(get_decoration(var_id, DecorationBuiltIn)); // Don't emit SamplePosition as a separate parameter. In the entry // point, we get that by calling get_sample_position() on the sample ID. - if (var.storage == StorageClassInput && is_builtin_variable(var) && + if (is_builtin_variable(var) && get_variable_data_type(var).basetype != SPIRType::Struct && get_variable_data_type(var).basetype != SPIRType::ControlPointArray) { @@ -10560,8 +10630,7 @@ void CompilerMSL::entry_point_args_builtin(string &ep_args) } } - if (var.storage == StorageClassInput && - has_extended_decoration(var_id, SPIRVCrossDecorationBuiltInDispatchBase)) + if (has_extended_decoration(var_id, SPIRVCrossDecorationBuiltInDispatchBase)) { // This is a special implicit builtin, not corresponding to any SPIR-V builtin, // which holds the base that was passed to vkCmdDispatchBase() or vkCmdDrawIndexed(). If it's present, @@ -10573,8 +10642,7 @@ void CompilerMSL::entry_point_args_builtin(string &ep_args) ep_args += type_to_glsl(get_variable_data_type(var)) + " " + to_expression(var_id) + " [[grid_origin]]"; } - if (var.storage == StorageClassInput && - has_extended_decoration(var_id, SPIRVCrossDecorationBuiltInStageInputSize)) + if (has_extended_decoration(var_id, SPIRVCrossDecorationBuiltInStageInputSize)) { // This is another special implicit builtin, not corresponding to any SPIR-V builtin, // which holds the number of vertices and instances to draw. If it's present, @@ -10668,6 +10736,58 @@ void CompilerMSL::entry_point_args_builtin(string &ep_args) ep_args += ", "; ep_args += join("device ", get_tess_factor_struct_name(), "* ", tess_factor_buffer_var_name, " [[buffer(", convert_to_string(msl_options.shader_tess_factor_buffer_index), ")]]"); + + // Initializer for tess factors must be handled specially since it's never declared as a normal variable. + uint32_t outer_factor_initializer_id = 0; + uint32_t inner_factor_initializer_id = 0; + ir.for_each_typed_id([&](uint32_t, SPIRVariable &var) { + if (!has_decoration(var.self, DecorationBuiltIn) || var.storage != StorageClassOutput || !var.initializer) + return; + + BuiltIn builtin = BuiltIn(get_decoration(var.self, DecorationBuiltIn)); + if (builtin == BuiltInTessLevelInner) + inner_factor_initializer_id = var.initializer; + else if (builtin == BuiltInTessLevelOuter) + outer_factor_initializer_id = var.initializer; + }); + + const SPIRConstant *c = nullptr; + + if (outer_factor_initializer_id && (c = maybe_get(outer_factor_initializer_id))) + { + auto &entry_func = get(ir.default_entry_point); + entry_func.fixup_hooks_in.push_back([=]() { + uint32_t components = get_execution_mode_bitset().get(ExecutionModeTriangles) ? 3 : 4; + for (uint32_t i = 0; i < components; i++) + { + statement(builtin_to_glsl(BuiltInTessLevelOuter, StorageClassOutput), "[", i, "] = ", + "half(", to_expression(c->subconstants[i]), ");"); + } + }); + } + + if (inner_factor_initializer_id && (c = maybe_get(inner_factor_initializer_id))) + { + auto &entry_func = get(ir.default_entry_point); + if (get_execution_mode_bitset().get(ExecutionModeTriangles)) + { + entry_func.fixup_hooks_in.push_back([=]() { + statement(builtin_to_glsl(BuiltInTessLevelInner, StorageClassOutput), " = ", "half(", + to_expression(c->subconstants[0]), ");"); + }); + } + else + { + entry_func.fixup_hooks_in.push_back([=]() { + for (uint32_t i = 0; i < 2; i++) + { + statement(builtin_to_glsl(BuiltInTessLevelInner, StorageClassOutput), "[", i, "] = ", + "half(", to_expression(c->subconstants[i]), ");"); + } + }); + } + } + if (stage_in_var_id) { if (!ep_args.empty()) @@ -10786,12 +10906,7 @@ void CompilerMSL::entry_point_args_discrete_descriptors(string &ep_args) { auto &type = get_variable_data_type(var); - // Very specifically, image load-store in argument buffers are disallowed on MSL on iOS. - // But we won't know when the argument buffer is encoded whether this image will have - // a NonWritable decoration. So just use discrete arguments for all storage images - // on iOS. - if (!(msl_options.is_ios() && type.basetype == SPIRType::Image && type.image.sampled == 2) && - var.storage != StorageClassPushConstant) + if (is_supported_argument_buffer_type(type) && var.storage != StorageClassPushConstant) { uint32_t desc_set = get_decoration(var_id, DecorationDescriptorSet); if (descriptor_set_is_argument_buffer(desc_set)) @@ -11059,6 +11174,11 @@ void CompilerMSL::fix_up_shader_inputs_outputs() uint32_t var_id = var.self; BuiltIn bi_type = ir.meta[var_id].decoration.builtin_type; + if (var.storage != StorageClassInput && var.storage != StorageClassOutput) + return; + if (!interface_variable_exists_in_entry_point(var.self)) + return; + if (var.storage == StorageClassInput && is_builtin_variable(var) && active_input_builtins.get(bi_type)) { switch (bi_type) @@ -11598,10 +11718,18 @@ uint32_t CompilerMSL::get_metal_resource_index(SPIRVariable &var, SPIRType::Base if (has_extended_decoration(var.self, resource_decoration)) return get_extended_decoration(var.self, resource_decoration); - // Allow user to enable decoration binding - if (msl_options.enable_decoration_binding) + auto &type = get(var.basetype); + + if (type_is_msl_framebuffer_fetch(type)) { - // If there is no explicit mapping of bindings to MSL, use the declared binding. + // Frame-buffer fetch gets its fallback resource index from the input attachment index, + // which is then treated as color index. + return get_decoration(var.self, DecorationInputAttachmentIndex); + } + else if (msl_options.enable_decoration_binding) + { + // Allow user to enable decoration binding. + // If there is no explicit mapping of bindings to MSL, use the declared binding as a fallback. if (has_decoration(var.self, DecorationBinding)) { var_binding = get_decoration(var.self, DecorationBinding); @@ -11620,7 +11748,6 @@ uint32_t CompilerMSL::get_metal_resource_index(SPIRVariable &var, SPIRType::Base allocate_argument_buffer_ids = descriptor_set_is_argument_buffer(var_desc_set); uint32_t binding_stride = 1; - auto &type = get(var.basetype); for (uint32_t i = 0; i < uint32_t(type.array.size()); i++) binding_stride *= to_array_size_literal(type, i); @@ -11629,13 +11756,7 @@ uint32_t CompilerMSL::get_metal_resource_index(SPIRVariable &var, SPIRType::Base // If a binding has not been specified, revert to incrementing resource indices. uint32_t resource_index; - if (type_is_msl_framebuffer_fetch(type)) - { - // Frame-buffer fetch gets its fallback resource index from the input attachment index, - // which is then treated as color index. - resource_index = get_decoration(var.self, DecorationInputAttachmentIndex); - } - else if (allocate_argument_buffer_ids) + if (allocate_argument_buffer_ids) { // Allocate from a flat ID binding space. resource_index = next_metal_resource_ids[var_desc_set]; @@ -11884,11 +12005,8 @@ string CompilerMSL::ensure_valid_name(string name, string pfx) return (name.size() >= 2 && name[0] == '_' && isdigit(name[1])) ? (pfx + name) : name; } -// Replace all names that match MSL keywords or Metal Standard Library functions. -void CompilerMSL::replace_illegal_names() +const std::unordered_set &CompilerMSL::get_reserved_keyword_set() { - // FIXME: MSL and GLSL are doing two different things here. - // Agree on convention and remove this override. static const unordered_set keywords = { "kernel", "vertex", @@ -12022,6 +12140,11 @@ void CompilerMSL::replace_illegal_names() "quad_broadcast", }; + return keywords; +} + +const std::unordered_set &CompilerMSL::get_illegal_func_names() +{ static const unordered_set illegal_func_names = { "main", "saturate", @@ -12148,6 +12271,17 @@ void CompilerMSL::replace_illegal_names() "M_SQRT1_2", }; + return illegal_func_names; +} + +// Replace all names that match MSL keywords or Metal Standard Library functions. +void CompilerMSL::replace_illegal_names() +{ + // FIXME: MSL and GLSL are doing two different things here. + // Agree on convention and remove this override. + auto &keywords = get_reserved_keyword_set(); + auto &illegal_func_names = get_illegal_func_names(); + ir.for_each_typed_id([&](uint32_t self, SPIRVariable &) { auto *meta = ir.find_meta(self); if (!meta) @@ -12178,6 +12312,16 @@ void CompilerMSL::replace_illegal_names() mbr_dec.alias += "0"; }); + CompilerGLSL::replace_illegal_names(); +} + +void CompilerMSL::replace_illegal_entry_point_names() +{ + auto &illegal_func_names = get_illegal_func_names(); + + // It is important to this before we fixup identifiers, + // since if ep_name is reserved, we will need to fix that up, + // and then copy alias back into entry.name after the fixup. for (auto &entry : ir.entry_points) { // Change both the entry point name and the alias, to keep them synced. @@ -12185,11 +12329,14 @@ void CompilerMSL::replace_illegal_names() if (illegal_func_names.find(ep_name) != end(illegal_func_names)) ep_name += "0"; - // Always write this because entry point might have been renamed earlier. ir.meta[entry.first].decoration.alias = ep_name; } +} - CompilerGLSL::replace_illegal_names(); +void CompilerMSL::sync_entry_point_aliases_and_names() +{ + for (auto &entry : ir.entry_points) + entry.second.name = ir.meta[entry.first].decoration.alias; } string CompilerMSL::to_member_reference(uint32_t base, const SPIRType &type, uint32_t index, bool ptr_chain) @@ -14445,6 +14592,17 @@ bool CompilerMSL::descriptor_set_is_argument_buffer(uint32_t desc_set) const return (argument_buffer_discrete_mask & (1u << desc_set)) == 0; } +bool CompilerMSL::is_supported_argument_buffer_type(const SPIRType &type) const +{ + // Very specifically, image load-store in argument buffers are disallowed on MSL on iOS. + // But we won't know when the argument buffer is encoded whether this image will have + // a NonWritable decoration. So just use discrete arguments for all storage images + // on iOS. + bool is_storage_image = type.basetype == SPIRType::Image && type.image.sampled == 2; + bool is_supported_type = !msl_options.is_ios() || !is_storage_image; + return !type_is_msl_framebuffer_fetch(type) && is_supported_type; +} + void CompilerMSL::analyze_argument_buffers() { // Gather all used resources and sort them out into argument buffers. @@ -14527,23 +14685,20 @@ void CompilerMSL::analyze_argument_buffers() { inline_block_vars.push_back(var_id); } - else if (!constexpr_sampler) + else if (!constexpr_sampler && is_supported_argument_buffer_type(type)) { // constexpr samplers are not declared as resources. // Inline uniform blocks are always emitted at the end. - if (!msl_options.is_ios() || type.basetype != SPIRType::Image || type.image.sampled != 2) - { - add_resource_name(var_id); - resources_in_set[desc_set].push_back( - { &var, to_name(var_id), type.basetype, get_metal_resource_index(var, type.basetype), 0 }); + add_resource_name(var_id); + resources_in_set[desc_set].push_back( + { &var, to_name(var_id), type.basetype, get_metal_resource_index(var, type.basetype), 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 }); - } + // 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 }); } } diff --git a/3rdparty/spirv-cross/spirv_msl.hpp b/3rdparty/spirv-cross/spirv_msl.hpp index 003fb83ae..8c7861236 100644 --- a/3rdparty/spirv-cross/spirv_msl.hpp +++ b/3rdparty/spirv-cross/spirv_msl.hpp @@ -735,6 +735,12 @@ protected: void declare_undefined_values() override; void declare_constant_arrays(); + void replace_illegal_entry_point_names(); + void sync_entry_point_aliases_and_names(); + + static const std::unordered_set &get_reserved_keyword_set(); + static const std::unordered_set &get_illegal_func_names(); + // Constant arrays of non-primitive types (i.e. matrices) won't link properly into Metal libraries void declare_complex_constant_arrays(); @@ -1020,6 +1026,7 @@ protected: void activate_argument_buffer_resources(); bool type_is_msl_framebuffer_fetch(const SPIRType &type) const; + bool is_supported_argument_buffer_type(const SPIRType &type) const; // OpcodeHandler that handles several MSL preprocessing operations. struct OpCodePreprocessor : OpcodeHandler diff --git a/3rdparty/spirv-cross/spirv_parser.cpp b/3rdparty/spirv-cross/spirv_parser.cpp index b7144c910..eff50dd30 100644 --- a/3rdparty/spirv-cross/spirv_parser.cpp +++ b/3rdparty/spirv-cross/spirv_parser.cpp @@ -727,7 +727,7 @@ void Parser::parse(const Instruction &instruction) break; } - case OpTypeRayQueryProvisionalKHR: + case OpTypeRayQueryKHR: { uint32_t id = ops[0]; auto &type = set(id); @@ -992,6 +992,22 @@ void Parser::parse(const Instruction &instruction) break; } + case OpTerminateRayKHR: + // NV variant is not a terminator. + if (!current_block) + SPIRV_CROSS_THROW("Trying to end a non-existing block."); + current_block->terminator = SPIRBlock::TerminateRay; + current_block = nullptr; + break; + + case OpIgnoreIntersectionKHR: + // NV variant is not a terminator. + if (!current_block) + SPIRV_CROSS_THROW("Trying to end a non-existing block."); + current_block->terminator = SPIRBlock::IgnoreIntersection; + current_block = nullptr; + break; + case OpReturn: { if (!current_block)