LOVE/love
1
0
Fork 0
mirror of https://github.com/love2d/love.git synced 2026-08-12 08:30:56 +02:00
Code Issues Packages Projects Releases Wiki Activity

update SPIRV-Cross

Browse Source
This commit is contained in:
Sasha Szpakowski
2024-01-07 21:17:43 -04:00
parent a99d5afc57
commit 421cfd1a5e
22 changed files with 9184 additions and 1946 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/libraries/spirv_cross/spirv.h
+384 -8
View File
@@ -31,7 +31,7 @@
/*
** Enumeration tokens for SPIR-V, in various styles:
** C, C++, C++11, JSON, Lua, Python, C#, D
** C, C++, C++11, JSON, Lua, Python, C#, D, Beef
**
** - C will have tokens with a "Spv" prefix, e.g.: SpvSourceLanguageGLSL
** - C++ will have tokens in the "spv" name space, e.g.: spv::SourceLanguageGLSL
@@ -41,6 +41,8 @@
** - C# will use enum classes in the Specification class located in the "Spv" namespace,
** e.g.: Spv.Specification.SourceLanguage.GLSL
** - D will have tokens under the "spv" module, e.g: spv.SourceLanguage.GLSL
** - Beef will use enum classes in the Specification class located in the "Spv" namespace,
** e.g.: Spv.Specification.SourceLanguage.GLSL
**
** Some tokens act like mask values, which can be OR'd together,
** while others are mutually exclusive. The mask-like ones have
@@ -53,12 +55,12 @@
typedef unsigned int SpvId;
#define SPV_VERSION 0x10500
#define SPV_REVISION 4
#define SPV_VERSION 0x10600
#define SPV_REVISION 1
static const unsigned int SpvMagicNumber = 0x07230203;
static const unsigned int SpvVersion = 0x00010500;
static const unsigned int SpvRevision = 4;
static const unsigned int SpvVersion = 0x00010600;
static const unsigned int SpvRevision = 1;
static const unsigned int SpvOpCodeMask = 0xffff;
static const unsigned int SpvWordCountShift = 16;
@@ -69,6 +71,8 @@ typedef enum SpvSourceLanguage_ {
SpvSourceLanguageOpenCL_C = 3,
SpvSourceLanguageOpenCL_CPP = 4,
SpvSourceLanguageHLSL = 5,
SpvSourceLanguageCPP_for_OpenCL = 6,
SpvSourceLanguageSYCL = 7,
SpvSourceLanguageMax = 0x7fffffff,
} SpvSourceLanguage;
@@ -94,6 +98,8 @@ typedef enum SpvExecutionModel_ {
SpvExecutionModelMissNV = 5317,
SpvExecutionModelCallableKHR = 5318,
SpvExecutionModelCallableNV = 5318,
SpvExecutionModelTaskEXT = 5364,
SpvExecutionModelMeshEXT = 5365,
SpvExecutionModelMax = 0x7fffffff,
} SpvExecutionModel;
@@ -154,17 +160,28 @@ typedef enum SpvExecutionMode_ {
SpvExecutionModeSubgroupsPerWorkgroupId = 37,
SpvExecutionModeLocalSizeId = 38,
SpvExecutionModeLocalSizeHintId = 39,
SpvExecutionModeSubgroupUniformControlFlowKHR = 4421,
SpvExecutionModePostDepthCoverage = 4446,
SpvExecutionModeDenormPreserve = 4459,
SpvExecutionModeDenormFlushToZero = 4460,
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,
@@ -172,10 +189,17 @@ typedef enum SpvExecutionMode_ {
SpvExecutionModeSampleInterlockUnorderedEXT = 5369,
SpvExecutionModeShadingRateInterlockOrderedEXT = 5370,
SpvExecutionModeShadingRateInterlockUnorderedEXT = 5371,
SpvExecutionModeSharedLocalMemorySizeINTEL = 5618,
SpvExecutionModeRoundingModeRTPINTEL = 5620,
SpvExecutionModeRoundingModeRTNINTEL = 5621,
SpvExecutionModeFloatingPointModeALTINTEL = 5622,
SpvExecutionModeFloatingPointModeIEEEINTEL = 5623,
SpvExecutionModeMaxWorkgroupSizeINTEL = 5893,
SpvExecutionModeMaxWorkDimINTEL = 5894,
SpvExecutionModeNoGlobalOffsetINTEL = 5895,
SpvExecutionModeNumSIMDWorkitemsINTEL = 5896,
SpvExecutionModeSchedulerTargetFmaxMhzINTEL = 5903,
SpvExecutionModeNamedBarrierCountINTEL = 6417,
SpvExecutionModeMax = 0x7fffffff,
} SpvExecutionMode;
@@ -207,7 +231,10 @@ typedef enum SpvStorageClass_ {
SpvStorageClassShaderRecordBufferNV = 5343,
SpvStorageClassPhysicalStorageBuffer = 5349,
SpvStorageClassPhysicalStorageBufferEXT = 5349,
SpvStorageClassTaskPayloadWorkgroupEXT = 5402,
SpvStorageClassCodeSectionINTEL = 5605,
SpvStorageClassDeviceOnlyINTEL = 5936,
SpvStorageClassHostOnlyINTEL = 5937,
SpvStorageClassMax = 0x7fffffff,
} SpvStorageClass;
@@ -347,6 +374,8 @@ typedef enum SpvImageOperandsShift_ {
SpvImageOperandsVolatileTexelKHRShift = 11,
SpvImageOperandsSignExtendShift = 12,
SpvImageOperandsZeroExtendShift = 13,
SpvImageOperandsNontemporalShift = 14,
SpvImageOperandsOffsetsShift = 16,
SpvImageOperandsMax = 0x7fffffff,
} SpvImageOperandsShift;
@@ -370,6 +399,8 @@ typedef enum SpvImageOperandsMask_ {
SpvImageOperandsVolatileTexelKHRMask = 0x00000800,
SpvImageOperandsSignExtendMask = 0x00001000,
SpvImageOperandsZeroExtendMask = 0x00002000,
SpvImageOperandsNontemporalMask = 0x00004000,
SpvImageOperandsOffsetsMask = 0x00010000,
} SpvImageOperandsMask;
typedef enum SpvFPFastMathModeShift_ {
@@ -378,6 +409,8 @@ typedef enum SpvFPFastMathModeShift_ {
SpvFPFastMathModeNSZShift = 2,
SpvFPFastMathModeAllowRecipShift = 3,
SpvFPFastMathModeFastShift = 4,
SpvFPFastMathModeAllowContractFastINTELShift = 16,
SpvFPFastMathModeAllowReassocINTELShift = 17,
SpvFPFastMathModeMax = 0x7fffffff,
} SpvFPFastMathModeShift;
@@ -388,6 +421,8 @@ typedef enum SpvFPFastMathModeMask_ {
SpvFPFastMathModeNSZMask = 0x00000004,
SpvFPFastMathModeAllowRecipMask = 0x00000008,
SpvFPFastMathModeFastMask = 0x00000010,
SpvFPFastMathModeAllowContractFastINTELMask = 0x00010000,
SpvFPFastMathModeAllowReassocINTELMask = 0x00020000,
} SpvFPFastMathModeMask;
typedef enum SpvFPRoundingMode_ {
@@ -401,6 +436,7 @@ typedef enum SpvFPRoundingMode_ {
typedef enum SpvLinkageType_ {
SpvLinkageTypeExport = 0,
SpvLinkageTypeImport = 1,
SpvLinkageTypeLinkOnceODR = 2,
SpvLinkageTypeMax = 0x7fffffff,
} SpvLinkageType;
@@ -478,9 +514,11 @@ typedef enum SpvDecoration_ {
SpvDecorationPassthroughNV = 5250,
SpvDecorationViewportRelativeNV = 5252,
SpvDecorationSecondaryViewportRelativeNV = 5256,
SpvDecorationPerPrimitiveEXT = 5271,
SpvDecorationPerPrimitiveNV = 5271,
SpvDecorationPerViewNV = 5272,
SpvDecorationPerTaskNV = 5273,
SpvDecorationPerVertexKHR = 5285,
SpvDecorationPerVertexNV = 5285,
SpvDecorationNonUniform = 5300,
SpvDecorationNonUniformEXT = 5300,
@@ -488,12 +526,26 @@ typedef enum SpvDecoration_ {
SpvDecorationRestrictPointerEXT = 5355,
SpvDecorationAliasedPointer = 5356,
SpvDecorationAliasedPointerEXT = 5356,
SpvDecorationBindlessSamplerNV = 5398,
SpvDecorationBindlessImageNV = 5399,
SpvDecorationBoundSamplerNV = 5400,
SpvDecorationBoundImageNV = 5401,
SpvDecorationSIMTCallINTEL = 5599,
SpvDecorationReferencedIndirectlyINTEL = 5602,
SpvDecorationClobberINTEL = 5607,
SpvDecorationSideEffectsINTEL = 5608,
SpvDecorationVectorComputeVariableINTEL = 5624,
SpvDecorationFuncParamIOKindINTEL = 5625,
SpvDecorationVectorComputeFunctionINTEL = 5626,
SpvDecorationStackCallINTEL = 5627,
SpvDecorationGlobalVariableOffsetINTEL = 5628,
SpvDecorationCounterBuffer = 5634,
SpvDecorationHlslCounterBufferGOOGLE = 5634,
SpvDecorationHlslSemanticGOOGLE = 5635,
SpvDecorationUserSemantic = 5635,
SpvDecorationUserTypeGOOGLE = 5636,
SpvDecorationFunctionRoundingModeINTEL = 5822,
SpvDecorationFunctionDenormModeINTEL = 5823,
SpvDecorationRegisterINTEL = 5825,
SpvDecorationMemoryINTEL = 5826,
SpvDecorationNumbanksINTEL = 5827,
@@ -506,6 +558,20 @@ typedef enum SpvDecoration_ {
SpvDecorationMergeINTEL = 5834,
SpvDecorationBankBitsINTEL = 5835,
SpvDecorationForcePow2DepthINTEL = 5836,
SpvDecorationBurstCoalesceINTEL = 5899,
SpvDecorationCacheSizeINTEL = 5900,
SpvDecorationDontStaticallyCoalesceINTEL = 5901,
SpvDecorationPrefetchINTEL = 5902,
SpvDecorationStallEnableINTEL = 5905,
SpvDecorationFuseLoopsInFunctionINTEL = 5907,
SpvDecorationAliasScopeINTEL = 5914,
SpvDecorationNoAliasINTEL = 5915,
SpvDecorationBufferLocationINTEL = 5921,
SpvDecorationIOPipeStorageINTEL = 5944,
SpvDecorationFunctionFloatingPointModeINTEL = 6080,
SpvDecorationSingleElementVectorINTEL = 6085,
SpvDecorationVectorComputeCallableFunctionINTEL = 6087,
SpvDecorationMediaBlockIOINTEL = 6140,
SpvDecorationMax = 0x7fffffff,
} SpvDecoration;
@@ -590,12 +656,18 @@ typedef enum SpvBuiltIn_ {
SpvBuiltInLayerPerViewNV = 5279,
SpvBuiltInMeshViewCountNV = 5280,
SpvBuiltInMeshViewIndicesNV = 5281,
SpvBuiltInBaryCoordKHR = 5286,
SpvBuiltInBaryCoordNV = 5286,
SpvBuiltInBaryCoordNoPerspKHR = 5287,
SpvBuiltInBaryCoordNoPerspNV = 5287,
SpvBuiltInFragSizeEXT = 5292,
SpvBuiltInFragmentSizeNV = 5292,
SpvBuiltInFragInvocationCountEXT = 5293,
SpvBuiltInInvocationsPerPixelNV = 5293,
SpvBuiltInPrimitivePointIndicesEXT = 5294,
SpvBuiltInPrimitiveLineIndicesEXT = 5295,
SpvBuiltInPrimitiveTriangleIndicesEXT = 5296,
SpvBuiltInCullPrimitiveEXT = 5299,
SpvBuiltInLaunchIdKHR = 5319,
SpvBuiltInLaunchIdNV = 5319,
SpvBuiltInLaunchSizeKHR = 5320,
@@ -621,6 +693,7 @@ typedef enum SpvBuiltIn_ {
SpvBuiltInHitTNV = 5332,
SpvBuiltInHitKindKHR = 5333,
SpvBuiltInHitKindNV = 5333,
SpvBuiltInCurrentRayTimeNV = 5334,
SpvBuiltInIncomingRayFlagsKHR = 5351,
SpvBuiltInIncomingRayFlagsNV = 5351,
SpvBuiltInRayGeometryIndexKHR = 5352,
@@ -628,6 +701,7 @@ typedef enum SpvBuiltIn_ {
SpvBuiltInSMCountNV = 5375,
SpvBuiltInWarpIDNV = 5376,
SpvBuiltInSMIDNV = 5377,
SpvBuiltInCullMaskKHR = 6021,
SpvBuiltInMax = 0x7fffffff,
} SpvBuiltIn;
@@ -660,6 +734,7 @@ typedef enum SpvLoopControlShift_ {
SpvLoopControlLoopCoalesceINTELShift = 20,
SpvLoopControlMaxInterleavingINTELShift = 21,
SpvLoopControlSpeculatedIterationsINTELShift = 22,
SpvLoopControlNoFusionINTELShift = 23,
SpvLoopControlMax = 0x7fffffff,
} SpvLoopControlShift;
@@ -681,6 +756,7 @@ typedef enum SpvLoopControlMask_ {
SpvLoopControlLoopCoalesceINTELMask = 0x00100000,
SpvLoopControlMaxInterleavingINTELMask = 0x00200000,
SpvLoopControlSpeculatedIterationsINTELMask = 0x00400000,
SpvLoopControlNoFusionINTELMask = 0x00800000,
} SpvLoopControlMask;
typedef enum SpvFunctionControlShift_ {
@@ -688,6 +764,7 @@ typedef enum SpvFunctionControlShift_ {
SpvFunctionControlDontInlineShift = 1,
SpvFunctionControlPureShift = 2,
SpvFunctionControlConstShift = 3,
SpvFunctionControlOptNoneINTELShift = 16,
SpvFunctionControlMax = 0x7fffffff,
} SpvFunctionControlShift;
@@ -697,6 +774,7 @@ typedef enum SpvFunctionControlMask_ {
SpvFunctionControlDontInlineMask = 0x00000002,
SpvFunctionControlPureMask = 0x00000004,
SpvFunctionControlConstMask = 0x00000008,
SpvFunctionControlOptNoneINTELMask = 0x00010000,
} SpvFunctionControlMask;
typedef enum SpvMemorySemanticsShift_ {
@@ -751,6 +829,8 @@ typedef enum SpvMemoryAccessShift_ {
SpvMemoryAccessMakePointerVisibleKHRShift = 4,
SpvMemoryAccessNonPrivatePointerShift = 5,
SpvMemoryAccessNonPrivatePointerKHRShift = 5,
SpvMemoryAccessAliasScopeINTELMaskShift = 16,
SpvMemoryAccessNoAliasINTELMaskShift = 17,
SpvMemoryAccessMax = 0x7fffffff,
} SpvMemoryAccessShift;
@@ -765,6 +845,8 @@ typedef enum SpvMemoryAccessMask_ {
SpvMemoryAccessMakePointerVisibleKHRMask = 0x00000010,
SpvMemoryAccessNonPrivatePointerMask = 0x00000020,
SpvMemoryAccessNonPrivatePointerKHRMask = 0x00000020,
SpvMemoryAccessAliasScopeINTELMaskMask = 0x00010000,
SpvMemoryAccessNoAliasINTELMaskMask = 0x00020000,
} SpvMemoryAccessMask;
typedef enum SpvScope_ {
@@ -877,9 +959,13 @@ typedef enum SpvCapability_ {
SpvCapabilityGroupNonUniformQuad = 68,
SpvCapabilityShaderLayer = 69,
SpvCapabilityShaderViewportIndex = 70,
SpvCapabilityUniformDecoration = 71,
SpvCapabilityFragmentShadingRateKHR = 4422,
SpvCapabilitySubgroupBallotKHR = 4423,
SpvCapabilityDrawParameters = 4427,
SpvCapabilityWorkgroupMemoryExplicitLayoutKHR = 4428,
SpvCapabilityWorkgroupMemoryExplicitLayout8BitAccessKHR = 4429,
SpvCapabilityWorkgroupMemoryExplicitLayout16BitAccessKHR = 4430,
SpvCapabilitySubgroupVoteKHR = 4431,
SpvCapabilityStorageBuffer16BitAccess = 4433,
SpvCapabilityStorageUniformBufferBlock16 = 4433,
@@ -922,6 +1008,8 @@ typedef enum SpvCapability_ {
SpvCapabilityFragmentFullyCoveredEXT = 5265,
SpvCapabilityMeshShadingNV = 5266,
SpvCapabilityImageFootprintNV = 5282,
SpvCapabilityMeshShadingEXT = 5283,
SpvCapabilityFragmentBarycentricKHR = 5284,
SpvCapabilityFragmentBarycentricNV = 5284,
SpvCapabilityComputeDerivativeGroupQuadsNV = 5288,
SpvCapabilityFragmentDensityEXT = 5291,
@@ -952,6 +1040,7 @@ typedef enum SpvCapability_ {
SpvCapabilityStorageTexelBufferArrayNonUniformIndexing = 5312,
SpvCapabilityStorageTexelBufferArrayNonUniformIndexingEXT = 5312,
SpvCapabilityRayTracingNV = 5340,
SpvCapabilityRayTracingMotionBlurNV = 5341,
SpvCapabilityVulkanMemoryModel = 5345,
SpvCapabilityVulkanMemoryModelKHR = 5345,
SpvCapabilityVulkanMemoryModelDeviceScope = 5346,
@@ -965,26 +1054,67 @@ typedef enum SpvCapability_ {
SpvCapabilityFragmentShaderShadingRateInterlockEXT = 5372,
SpvCapabilityShaderSMBuiltinsNV = 5373,
SpvCapabilityFragmentShaderPixelInterlockEXT = 5378,
SpvCapabilityDemoteToHelperInvocation = 5379,
SpvCapabilityDemoteToHelperInvocationEXT = 5379,
SpvCapabilityBindlessTextureNV = 5390,
SpvCapabilitySubgroupShuffleINTEL = 5568,
SpvCapabilitySubgroupBufferBlockIOINTEL = 5569,
SpvCapabilitySubgroupImageBlockIOINTEL = 5570,
SpvCapabilitySubgroupImageMediaBlockIOINTEL = 5579,
SpvCapabilityRoundToInfinityINTEL = 5582,
SpvCapabilityFloatingPointModeINTEL = 5583,
SpvCapabilityIntegerFunctions2INTEL = 5584,
SpvCapabilityFunctionPointersINTEL = 5603,
SpvCapabilityIndirectReferencesINTEL = 5604,
SpvCapabilityAsmINTEL = 5606,
SpvCapabilityAtomicFloat32MinMaxEXT = 5612,
SpvCapabilityAtomicFloat64MinMaxEXT = 5613,
SpvCapabilityAtomicFloat16MinMaxEXT = 5616,
SpvCapabilityVectorComputeINTEL = 5617,
SpvCapabilityVectorAnyINTEL = 5619,
SpvCapabilityExpectAssumeKHR = 5629,
SpvCapabilitySubgroupAvcMotionEstimationINTEL = 5696,
SpvCapabilitySubgroupAvcMotionEstimationIntraINTEL = 5697,
SpvCapabilitySubgroupAvcMotionEstimationChromaINTEL = 5698,
SpvCapabilityVariableLengthArrayINTEL = 5817,
SpvCapabilityFunctionFloatControlINTEL = 5821,
SpvCapabilityFPGAMemoryAttributesINTEL = 5824,
SpvCapabilityFPFastMathModeINTEL = 5837,
SpvCapabilityArbitraryPrecisionIntegersINTEL = 5844,
SpvCapabilityArbitraryPrecisionFloatingPointINTEL = 5845,
SpvCapabilityUnstructuredLoopControlsINTEL = 5886,
SpvCapabilityFPGALoopControlsINTEL = 5888,
SpvCapabilityKernelAttributesINTEL = 5892,
SpvCapabilityFPGAKernelAttributesINTEL = 5897,
SpvCapabilityFPGAMemoryAccessesINTEL = 5898,
SpvCapabilityFPGAClusterAttributesINTEL = 5904,
SpvCapabilityLoopFuseINTEL = 5906,
SpvCapabilityMemoryAccessAliasingINTEL = 5910,
SpvCapabilityFPGABufferLocationINTEL = 5920,
SpvCapabilityArbitraryPrecisionFixedPointINTEL = 5922,
SpvCapabilityUSMStorageClassesINTEL = 5935,
SpvCapabilityIOPipesINTEL = 5943,
SpvCapabilityBlockingPipesINTEL = 5945,
SpvCapabilityFPGARegINTEL = 5948,
SpvCapabilityDotProductInputAll = 6016,
SpvCapabilityDotProductInputAllKHR = 6016,
SpvCapabilityDotProductInput4x8Bit = 6017,
SpvCapabilityDotProductInput4x8BitKHR = 6017,
SpvCapabilityDotProductInput4x8BitPacked = 6018,
SpvCapabilityDotProductInput4x8BitPackedKHR = 6018,
SpvCapabilityDotProduct = 6019,
SpvCapabilityDotProductKHR = 6019,
SpvCapabilityRayCullMaskKHR = 6020,
SpvCapabilityBitInstructions = 6025,
SpvCapabilityGroupNonUniformRotateKHR = 6026,
SpvCapabilityAtomicFloat32AddEXT = 6033,
SpvCapabilityAtomicFloat64AddEXT = 6034,
SpvCapabilityLongConstantCompositeINTEL = 6089,
SpvCapabilityOptNoneINTEL = 6094,
SpvCapabilityAtomicFloat16AddEXT = 6095,
SpvCapabilityDebugInfoModuleINTEL = 6114,
SpvCapabilitySplitBarrierINTEL = 6141,
SpvCapabilityGroupUniformArithmeticKHR = 6400,
SpvCapabilityMax = 0x7fffffff,
} SpvCapability;
@@ -1051,6 +1181,44 @@ typedef enum SpvFragmentShadingRateMask_ {
SpvFragmentShadingRateHorizontal4PixelsMask = 0x00000008,
} SpvFragmentShadingRateMask;
typedef enum SpvFPDenormMode_ {
SpvFPDenormModePreserve = 0,
SpvFPDenormModeFlushToZero = 1,
SpvFPDenormModeMax = 0x7fffffff,
} SpvFPDenormMode;
typedef enum SpvFPOperationMode_ {
SpvFPOperationModeIEEE = 0,
SpvFPOperationModeALT = 1,
SpvFPOperationModeMax = 0x7fffffff,
} SpvFPOperationMode;
typedef enum SpvQuantizationModes_ {
SpvQuantizationModesTRN = 0,
SpvQuantizationModesTRN_ZERO = 1,
SpvQuantizationModesRND = 2,
SpvQuantizationModesRND_ZERO = 3,
SpvQuantizationModesRND_INF = 4,
SpvQuantizationModesRND_MIN_INF = 5,
SpvQuantizationModesRND_CONV = 6,
SpvQuantizationModesRND_CONV_ODD = 7,
SpvQuantizationModesMax = 0x7fffffff,
} SpvQuantizationModes;
typedef enum SpvOverflowModes_ {
SpvOverflowModesWRAP = 0,
SpvOverflowModesSAT = 1,
SpvOverflowModesSAT_ZERO = 2,
SpvOverflowModesSAT_SYM = 3,
SpvOverflowModesMax = 0x7fffffff,
} SpvOverflowModes;
typedef enum SpvPackedVectorFormat_ {
SpvPackedVectorFormatPackedVectorFormat4x8Bit = 0,
SpvPackedVectorFormatPackedVectorFormat4x8BitKHR = 0,
SpvPackedVectorFormatMax = 0x7fffffff,
} SpvPackedVectorFormat;
typedef enum SpvOp_ {
SpvOpNop = 0,
SpvOpUndef = 1,
@@ -1402,12 +1570,25 @@ typedef enum SpvOp_ {
SpvOpSubgroupAllKHR = 4428,
SpvOpSubgroupAnyKHR = 4429,
SpvOpSubgroupAllEqualKHR = 4430,
SpvOpGroupNonUniformRotateKHR = 4431,
SpvOpSubgroupReadInvocationKHR = 4432,
SpvOpTraceRayKHR = 4445,
SpvOpExecuteCallableKHR = 4446,
SpvOpConvertUToAccelerationStructureKHR = 4447,
SpvOpIgnoreIntersectionKHR = 4448,
SpvOpTerminateRayKHR = 4449,
SpvOpSDot = 4450,
SpvOpSDotKHR = 4450,
SpvOpUDot = 4451,
SpvOpUDotKHR = 4451,
SpvOpSUDot = 4452,
SpvOpSUDotKHR = 4452,
SpvOpSDotAccSat = 4453,
SpvOpSDotAccSatKHR = 4453,
SpvOpUDotAccSat = 4454,
SpvOpUDotAccSatKHR = 4454,
SpvOpSUDotAccSat = 4455,
SpvOpSUDotAccSatKHR = 4455,
SpvOpTypeRayQueryKHR = 4472,
SpvOpRayQueryInitializeKHR = 4473,
SpvOpRayQueryTerminateKHR = 4474,
@@ -1427,6 +1608,8 @@ typedef enum SpvOp_ {
SpvOpFragmentFetchAMD = 5012,
SpvOpReadClockKHR = 5056,
SpvOpImageSampleFootprintNV = 5283,
SpvOpEmitMeshTasksEXT = 5294,
SpvOpSetMeshOutputsEXT = 5295,
SpvOpGroupNonUniformPartitionNV = 5296,
SpvOpWritePackedPrimitiveIndices4x8NV = 5299,
SpvOpReportIntersectionKHR = 5334,
@@ -1434,6 +1617,8 @@ typedef enum SpvOp_ {
SpvOpIgnoreIntersectionNV = 5335,
SpvOpTerminateRayNV = 5336,
SpvOpTraceNV = 5337,
SpvOpTraceMotionNV = 5338,
SpvOpTraceRayMotionNV = 5339,
SpvOpTypeAccelerationStructureKHR = 5341,
SpvOpTypeAccelerationStructureNV = 5341,
SpvOpExecuteCallableNV = 5344,
@@ -1444,8 +1629,16 @@ typedef enum SpvOp_ {
SpvOpCooperativeMatrixLengthNV = 5362,
SpvOpBeginInvocationInterlockEXT = 5364,
SpvOpEndInvocationInterlockEXT = 5365,
SpvOpDemoteToHelperInvocation = 5380,
SpvOpDemoteToHelperInvocationEXT = 5380,
SpvOpIsHelperInvocationEXT = 5381,
SpvOpConvertUToImageNV = 5391,
SpvOpConvertUToSamplerNV = 5392,
SpvOpConvertImageToUNV = 5393,
SpvOpConvertSamplerToUNV = 5394,
SpvOpConvertUToSampledImageNV = 5395,
SpvOpConvertSampledImageToUNV = 5396,
SpvOpSamplerImageAddressingModeNV = 5397,
SpvOpSubgroupShuffleINTEL = 5571,
SpvOpSubgroupShuffleDownINTEL = 5572,
SpvOpSubgroupShuffleUpINTEL = 5573,
@@ -1470,8 +1663,15 @@ typedef enum SpvOp_ {
SpvOpUSubSatINTEL = 5596,
SpvOpIMul32x16INTEL = 5597,
SpvOpUMul32x16INTEL = 5598,
SpvOpFunctionPointerINTEL = 5600,
SpvOpConstantFunctionPointerINTEL = 5600,
SpvOpFunctionPointerCallINTEL = 5601,
SpvOpAsmTargetINTEL = 5609,
SpvOpAsmINTEL = 5610,
SpvOpAsmCallINTEL = 5611,
SpvOpAtomicFMinEXT = 5614,
SpvOpAtomicFMaxEXT = 5615,
SpvOpAssumeTrueKHR = 5630,
SpvOpExpectKHR = 5631,
SpvOpDecorateString = 5632,
SpvOpDecorateStringGOOGLE = 5632,
SpvOpMemberDecorateString = 5633,
@@ -1594,7 +1794,67 @@ typedef enum SpvOp_ {
SpvOpSubgroupAvcSicGetPackedSkcLumaCountThresholdINTEL = 5814,
SpvOpSubgroupAvcSicGetPackedSkcLumaSumThresholdINTEL = 5815,
SpvOpSubgroupAvcSicGetInterRawSadsINTEL = 5816,
SpvOpVariableLengthArrayINTEL = 5818,
SpvOpSaveMemoryINTEL = 5819,
SpvOpRestoreMemoryINTEL = 5820,
SpvOpArbitraryFloatSinCosPiINTEL = 5840,
SpvOpArbitraryFloatCastINTEL = 5841,
SpvOpArbitraryFloatCastFromIntINTEL = 5842,
SpvOpArbitraryFloatCastToIntINTEL = 5843,
SpvOpArbitraryFloatAddINTEL = 5846,
SpvOpArbitraryFloatSubINTEL = 5847,
SpvOpArbitraryFloatMulINTEL = 5848,
SpvOpArbitraryFloatDivINTEL = 5849,
SpvOpArbitraryFloatGTINTEL = 5850,
SpvOpArbitraryFloatGEINTEL = 5851,
SpvOpArbitraryFloatLTINTEL = 5852,
SpvOpArbitraryFloatLEINTEL = 5853,
SpvOpArbitraryFloatEQINTEL = 5854,
SpvOpArbitraryFloatRecipINTEL = 5855,
SpvOpArbitraryFloatRSqrtINTEL = 5856,
SpvOpArbitraryFloatCbrtINTEL = 5857,
SpvOpArbitraryFloatHypotINTEL = 5858,
SpvOpArbitraryFloatSqrtINTEL = 5859,
SpvOpArbitraryFloatLogINTEL = 5860,
SpvOpArbitraryFloatLog2INTEL = 5861,
SpvOpArbitraryFloatLog10INTEL = 5862,
SpvOpArbitraryFloatLog1pINTEL = 5863,
SpvOpArbitraryFloatExpINTEL = 5864,
SpvOpArbitraryFloatExp2INTEL = 5865,
SpvOpArbitraryFloatExp10INTEL = 5866,
SpvOpArbitraryFloatExpm1INTEL = 5867,
SpvOpArbitraryFloatSinINTEL = 5868,
SpvOpArbitraryFloatCosINTEL = 5869,
SpvOpArbitraryFloatSinCosINTEL = 5870,
SpvOpArbitraryFloatSinPiINTEL = 5871,
SpvOpArbitraryFloatCosPiINTEL = 5872,
SpvOpArbitraryFloatASinINTEL = 5873,
SpvOpArbitraryFloatASinPiINTEL = 5874,
SpvOpArbitraryFloatACosINTEL = 5875,
SpvOpArbitraryFloatACosPiINTEL = 5876,
SpvOpArbitraryFloatATanINTEL = 5877,
SpvOpArbitraryFloatATanPiINTEL = 5878,
SpvOpArbitraryFloatATan2INTEL = 5879,
SpvOpArbitraryFloatPowINTEL = 5880,
SpvOpArbitraryFloatPowRINTEL = 5881,
SpvOpArbitraryFloatPowNINTEL = 5882,
SpvOpLoopControlINTEL = 5887,
SpvOpAliasDomainDeclINTEL = 5911,
SpvOpAliasScopeDeclINTEL = 5912,
SpvOpAliasScopeListDeclINTEL = 5913,
SpvOpFixedSqrtINTEL = 5923,
SpvOpFixedRecipINTEL = 5924,
SpvOpFixedRsqrtINTEL = 5925,
SpvOpFixedSinINTEL = 5926,
SpvOpFixedCosINTEL = 5927,
SpvOpFixedSinCosINTEL = 5928,
SpvOpFixedSinPiINTEL = 5929,
SpvOpFixedCosPiINTEL = 5930,
SpvOpFixedSinCosPiINTEL = 5931,
SpvOpFixedLogINTEL = 5932,
SpvOpFixedExpINTEL = 5933,
SpvOpPtrCastToCrossWorkgroupINTEL = 5934,
SpvOpCrossWorkgroupCastToPtrINTEL = 5938,
SpvOpReadPipeBlockingINTEL = 5946,
SpvOpWritePipeBlockingINTEL = 5947,
SpvOpFPGARegINTEL = 5949,
@@ -1616,10 +1876,27 @@ typedef enum SpvOp_ {
SpvOpRayQueryGetIntersectionObjectToWorldKHR = 6031,
SpvOpRayQueryGetIntersectionWorldToObjectKHR = 6032,
SpvOpAtomicFAddEXT = 6035,
SpvOpTypeBufferSurfaceINTEL = 6086,
SpvOpTypeStructContinuedINTEL = 6090,
SpvOpConstantCompositeContinuedINTEL = 6091,
SpvOpSpecConstantCompositeContinuedINTEL = 6092,
SpvOpControlBarrierArriveINTEL = 6142,
SpvOpControlBarrierWaitINTEL = 6143,
SpvOpGroupIMulKHR = 6401,
SpvOpGroupFMulKHR = 6402,
SpvOpGroupBitwiseAndKHR = 6403,
SpvOpGroupBitwiseOrKHR = 6404,
SpvOpGroupBitwiseXorKHR = 6405,
SpvOpGroupLogicalAndKHR = 6406,
SpvOpGroupLogicalOrKHR = 6407,
SpvOpGroupLogicalXorKHR = 6408,
SpvOpMax = 0x7fffffff,
} SpvOp;
#ifdef SPV_ENABLE_UTILITY_CODE
#ifndef __cplusplus
#include <stdbool.h>
#endif
inline void SpvHasResultAndType(SpvOp opcode, bool *hasResult, bool *hasResultType) {
*hasResult = *hasResultType = false;
switch (opcode) {
@@ -1974,12 +2251,19 @@ inline void SpvHasResultAndType(SpvOp opcode, bool *hasResult, bool *hasResultTy
case SpvOpSubgroupAllKHR: *hasResult = true; *hasResultType = true; break;
case SpvOpSubgroupAnyKHR: *hasResult = true; *hasResultType = true; break;
case SpvOpSubgroupAllEqualKHR: *hasResult = true; *hasResultType = true; break;
case SpvOpGroupNonUniformRotateKHR: *hasResult = true; *hasResultType = true; break;
case SpvOpSubgroupReadInvocationKHR: *hasResult = true; *hasResultType = true; 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 SpvOpSDot: *hasResult = true; *hasResultType = true; break;
case SpvOpUDot: *hasResult = true; *hasResultType = true; break;
case SpvOpSUDot: *hasResult = true; *hasResultType = true; break;
case SpvOpSDotAccSat: *hasResult = true; *hasResultType = true; break;
case SpvOpUDotAccSat: *hasResult = true; *hasResultType = true; break;
case SpvOpSUDotAccSat: *hasResult = true; *hasResultType = true; break;
case SpvOpTypeRayQueryKHR: *hasResult = true; *hasResultType = false; break;
case SpvOpRayQueryInitializeKHR: *hasResult = false; *hasResultType = false; break;
case SpvOpRayQueryTerminateKHR: *hasResult = false; *hasResultType = false; break;
@@ -1999,12 +2283,16 @@ 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;
case SpvOpIgnoreIntersectionNV: *hasResult = false; *hasResultType = false; break;
case SpvOpTerminateRayNV: *hasResult = false; *hasResultType = false; break;
case SpvOpTraceNV: *hasResult = false; *hasResultType = false; break;
case SpvOpTraceMotionNV: *hasResult = false; *hasResultType = false; break;
case SpvOpTraceRayMotionNV: *hasResult = false; *hasResultType = false; break;
case SpvOpTypeAccelerationStructureNV: *hasResult = true; *hasResultType = false; break;
case SpvOpExecuteCallableNV: *hasResult = false; *hasResultType = false; break;
case SpvOpTypeCooperativeMatrixNV: *hasResult = true; *hasResultType = false; break;
@@ -2014,8 +2302,15 @@ inline void SpvHasResultAndType(SpvOp opcode, bool *hasResult, bool *hasResultTy
case SpvOpCooperativeMatrixLengthNV: *hasResult = true; *hasResultType = true; break;
case SpvOpBeginInvocationInterlockEXT: *hasResult = false; *hasResultType = false; break;
case SpvOpEndInvocationInterlockEXT: *hasResult = false; *hasResultType = false; break;
case SpvOpDemoteToHelperInvocationEXT: *hasResult = false; *hasResultType = false; break;
case SpvOpDemoteToHelperInvocation: *hasResult = false; *hasResultType = false; break;
case SpvOpIsHelperInvocationEXT: *hasResult = true; *hasResultType = true; break;
case SpvOpConvertUToImageNV: *hasResult = true; *hasResultType = true; break;
case SpvOpConvertUToSamplerNV: *hasResult = true; *hasResultType = true; break;
case SpvOpConvertImageToUNV: *hasResult = true; *hasResultType = true; break;
case SpvOpConvertSamplerToUNV: *hasResult = true; *hasResultType = true; break;
case SpvOpConvertUToSampledImageNV: *hasResult = true; *hasResultType = true; break;
case SpvOpConvertSampledImageToUNV: *hasResult = true; *hasResultType = true; break;
case SpvOpSamplerImageAddressingModeNV: *hasResult = false; *hasResultType = false; break;
case SpvOpSubgroupShuffleINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpSubgroupShuffleDownINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpSubgroupShuffleUpINTEL: *hasResult = true; *hasResultType = true; break;
@@ -2040,8 +2335,15 @@ 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 SpvOpConstantFunctionPointerINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpFunctionPointerCallINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpAsmTargetINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpAsmINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpAsmCallINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpAtomicFMinEXT: *hasResult = true; *hasResultType = true; break;
case SpvOpAtomicFMaxEXT: *hasResult = true; *hasResultType = true; break;
case SpvOpAssumeTrueKHR: *hasResult = false; *hasResultType = false; break;
case SpvOpExpectKHR: *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;
@@ -2162,7 +2464,67 @@ 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 SpvOpVariableLengthArrayINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpSaveMemoryINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpRestoreMemoryINTEL: *hasResult = false; *hasResultType = false; break;
case SpvOpArbitraryFloatSinCosPiINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatCastINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatCastFromIntINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatCastToIntINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatAddINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatSubINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatMulINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatDivINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatGTINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatGEINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatLTINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatLEINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatEQINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatRecipINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatRSqrtINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatCbrtINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatHypotINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatSqrtINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatLogINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatLog2INTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatLog10INTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatLog1pINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatExpINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatExp2INTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatExp10INTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatExpm1INTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatSinINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatCosINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatSinCosINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatSinPiINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatCosPiINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatASinINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatASinPiINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatACosINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatACosPiINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatATanINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatATanPiINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatATan2INTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatPowINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatPowRINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpArbitraryFloatPowNINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpLoopControlINTEL: *hasResult = false; *hasResultType = false; break;
case SpvOpAliasDomainDeclINTEL: *hasResult = true; *hasResultType = false; break;
case SpvOpAliasScopeDeclINTEL: *hasResult = true; *hasResultType = false; break;
case SpvOpAliasScopeListDeclINTEL: *hasResult = true; *hasResultType = false; break;
case SpvOpFixedSqrtINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpFixedRecipINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpFixedRsqrtINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpFixedSinINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpFixedCosINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpFixedSinCosINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpFixedSinPiINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpFixedCosPiINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpFixedSinCosPiINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpFixedLogINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpFixedExpINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpPtrCastToCrossWorkgroupINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpCrossWorkgroupCastToPtrINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpReadPipeBlockingINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpWritePipeBlockingINTEL: *hasResult = true; *hasResultType = true; break;
case SpvOpFPGARegINTEL: *hasResult = true; *hasResultType = true; break;
@@ -2184,6 +2546,20 @@ inline void SpvHasResultAndType(SpvOp opcode, bool *hasResult, bool *hasResultTy
case SpvOpRayQueryGetIntersectionObjectToWorldKHR: *hasResult = true; *hasResultType = true; break;
case SpvOpRayQueryGetIntersectionWorldToObjectKHR: *hasResult = true; *hasResultType = true; break;
case SpvOpAtomicFAddEXT: *hasResult = true; *hasResultType = true; break;
case SpvOpTypeBufferSurfaceINTEL: *hasResult = true; *hasResultType = false; break;
case SpvOpTypeStructContinuedINTEL: *hasResult = false; *hasResultType = false; break;
case SpvOpConstantCompositeContinuedINTEL: *hasResult = false; *hasResultType = false; break;
case SpvOpSpecConstantCompositeContinuedINTEL: *hasResult = false; *hasResultType = false; break;
case SpvOpControlBarrierArriveINTEL: *hasResult = false; *hasResultType = false; break;
case SpvOpControlBarrierWaitINTEL: *hasResult = false; *hasResultType = false; break;
case SpvOpGroupIMulKHR: *hasResult = true; *hasResultType = true; break;
case SpvOpGroupFMulKHR: *hasResult = true; *hasResultType = true; break;
case SpvOpGroupBitwiseAndKHR: *hasResult = true; *hasResultType = true; break;
case SpvOpGroupBitwiseOrKHR: *hasResult = true; *hasResultType = true; break;
case SpvOpGroupBitwiseXorKHR: *hasResult = true; *hasResultType = true; break;
case SpvOpGroupLogicalAndKHR: *hasResult = true; *hasResultType = true; break;
case SpvOpGroupLogicalOrKHR: *hasResult = true; *hasResultType = true; break;
case SpvOpGroupLogicalXorKHR: *hasResult = true; *hasResultType = true; break;
}
}
#endif /* SPV_ENABLE_UTILITY_CODE */
src/libraries/spirv_cross/spirv.hpp
+384 -8
View File
@@ -26,7 +26,7 @@
// the Binary Section of the SPIR-V specification.
// Enumeration tokens for SPIR-V, in various styles:
// C, C++, C++11, JSON, Lua, Python, C#, D
// C, C++, C++11, JSON, Lua, Python, C#, D, Beef
//
// - C will have tokens with a "Spv" prefix, e.g.: SpvSourceLanguageGLSL
// - C++ will have tokens in the "spv" name space, e.g.: spv::SourceLanguageGLSL
@@ -36,6 +36,8 @@
// - C# will use enum classes in the Specification class located in the "Spv" namespace,
// e.g.: Spv.Specification.SourceLanguage.GLSL
// - D will have tokens under the "spv" module, e.g: spv.SourceLanguage.GLSL
// - Beef will use enum classes in the Specification class located in the "Spv" namespace,
// e.g.: Spv.Specification.SourceLanguage.GLSL
//
// Some tokens act like mask values, which can be OR'd together,
// while others are mutually exclusive. The mask-like ones have
@@ -49,12 +51,12 @@ namespace spv {
typedef unsigned int Id;
#define SPV_VERSION 0x10500
#define SPV_REVISION 4
#define SPV_VERSION 0x10600
#define SPV_REVISION 1
static const unsigned int MagicNumber = 0x07230203;
static const unsigned int Version = 0x00010500;
static const unsigned int Revision = 4;
static const unsigned int Version = 0x00010600;
static const unsigned int Revision = 1;
static const unsigned int OpCodeMask = 0xffff;
static const unsigned int WordCountShift = 16;
@@ -65,6 +67,8 @@ enum SourceLanguage {
SourceLanguageOpenCL_C = 3,
SourceLanguageOpenCL_CPP = 4,
SourceLanguageHLSL = 5,
SourceLanguageCPP_for_OpenCL = 6,
SourceLanguageSYCL = 7,
SourceLanguageMax = 0x7fffffff,
};
@@ -90,6 +94,8 @@ enum ExecutionModel {
ExecutionModelMissNV = 5317,
ExecutionModelCallableKHR = 5318,
ExecutionModelCallableNV = 5318,
ExecutionModelTaskEXT = 5364,
ExecutionModelMeshEXT = 5365,
ExecutionModelMax = 0x7fffffff,
};
@@ -150,17 +156,28 @@ enum ExecutionMode {
ExecutionModeSubgroupsPerWorkgroupId = 37,
ExecutionModeLocalSizeId = 38,
ExecutionModeLocalSizeHintId = 39,
ExecutionModeSubgroupUniformControlFlowKHR = 4421,
ExecutionModePostDepthCoverage = 4446,
ExecutionModeDenormPreserve = 4459,
ExecutionModeDenormFlushToZero = 4460,
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,
@@ -168,10 +185,17 @@ enum ExecutionMode {
ExecutionModeSampleInterlockUnorderedEXT = 5369,
ExecutionModeShadingRateInterlockOrderedEXT = 5370,
ExecutionModeShadingRateInterlockUnorderedEXT = 5371,
ExecutionModeSharedLocalMemorySizeINTEL = 5618,
ExecutionModeRoundingModeRTPINTEL = 5620,
ExecutionModeRoundingModeRTNINTEL = 5621,
ExecutionModeFloatingPointModeALTINTEL = 5622,
ExecutionModeFloatingPointModeIEEEINTEL = 5623,
ExecutionModeMaxWorkgroupSizeINTEL = 5893,
ExecutionModeMaxWorkDimINTEL = 5894,
ExecutionModeNoGlobalOffsetINTEL = 5895,
ExecutionModeNumSIMDWorkitemsINTEL = 5896,
ExecutionModeSchedulerTargetFmaxMhzINTEL = 5903,
ExecutionModeNamedBarrierCountINTEL = 6417,
ExecutionModeMax = 0x7fffffff,
};
@@ -203,7 +227,10 @@ enum StorageClass {
StorageClassShaderRecordBufferNV = 5343,
StorageClassPhysicalStorageBuffer = 5349,
StorageClassPhysicalStorageBufferEXT = 5349,
StorageClassTaskPayloadWorkgroupEXT = 5402,
StorageClassCodeSectionINTEL = 5605,
StorageClassDeviceOnlyINTEL = 5936,
StorageClassHostOnlyINTEL = 5937,
StorageClassMax = 0x7fffffff,
};
@@ -343,6 +370,8 @@ enum ImageOperandsShift {
ImageOperandsVolatileTexelKHRShift = 11,
ImageOperandsSignExtendShift = 12,
ImageOperandsZeroExtendShift = 13,
ImageOperandsNontemporalShift = 14,
ImageOperandsOffsetsShift = 16,
ImageOperandsMax = 0x7fffffff,
};
@@ -366,6 +395,8 @@ enum ImageOperandsMask {
ImageOperandsVolatileTexelKHRMask = 0x00000800,
ImageOperandsSignExtendMask = 0x00001000,
ImageOperandsZeroExtendMask = 0x00002000,
ImageOperandsNontemporalMask = 0x00004000,
ImageOperandsOffsetsMask = 0x00010000,
};
enum FPFastMathModeShift {
@@ -374,6 +405,8 @@ enum FPFastMathModeShift {
FPFastMathModeNSZShift = 2,
FPFastMathModeAllowRecipShift = 3,
FPFastMathModeFastShift = 4,
FPFastMathModeAllowContractFastINTELShift = 16,
FPFastMathModeAllowReassocINTELShift = 17,
FPFastMathModeMax = 0x7fffffff,
};
@@ -384,6 +417,8 @@ enum FPFastMathModeMask {
FPFastMathModeNSZMask = 0x00000004,
FPFastMathModeAllowRecipMask = 0x00000008,
FPFastMathModeFastMask = 0x00000010,
FPFastMathModeAllowContractFastINTELMask = 0x00010000,
FPFastMathModeAllowReassocINTELMask = 0x00020000,
};
enum FPRoundingMode {
@@ -397,6 +432,7 @@ enum FPRoundingMode {
enum LinkageType {
LinkageTypeExport = 0,
LinkageTypeImport = 1,
LinkageTypeLinkOnceODR = 2,
LinkageTypeMax = 0x7fffffff,
};
@@ -474,9 +510,11 @@ enum Decoration {
DecorationPassthroughNV = 5250,
DecorationViewportRelativeNV = 5252,
DecorationSecondaryViewportRelativeNV = 5256,
DecorationPerPrimitiveEXT = 5271,
DecorationPerPrimitiveNV = 5271,
DecorationPerViewNV = 5272,
DecorationPerTaskNV = 5273,
DecorationPerVertexKHR = 5285,
DecorationPerVertexNV = 5285,
DecorationNonUniform = 5300,
DecorationNonUniformEXT = 5300,
@@ -484,12 +522,26 @@ enum Decoration {
DecorationRestrictPointerEXT = 5355,
DecorationAliasedPointer = 5356,
DecorationAliasedPointerEXT = 5356,
DecorationBindlessSamplerNV = 5398,
DecorationBindlessImageNV = 5399,
DecorationBoundSamplerNV = 5400,
DecorationBoundImageNV = 5401,
DecorationSIMTCallINTEL = 5599,
DecorationReferencedIndirectlyINTEL = 5602,
DecorationClobberINTEL = 5607,
DecorationSideEffectsINTEL = 5608,
DecorationVectorComputeVariableINTEL = 5624,
DecorationFuncParamIOKindINTEL = 5625,
DecorationVectorComputeFunctionINTEL = 5626,
DecorationStackCallINTEL = 5627,
DecorationGlobalVariableOffsetINTEL = 5628,
DecorationCounterBuffer = 5634,
DecorationHlslCounterBufferGOOGLE = 5634,
DecorationHlslSemanticGOOGLE = 5635,
DecorationUserSemantic = 5635,
DecorationUserTypeGOOGLE = 5636,
DecorationFunctionRoundingModeINTEL = 5822,
DecorationFunctionDenormModeINTEL = 5823,
DecorationRegisterINTEL = 5825,
DecorationMemoryINTEL = 5826,
DecorationNumbanksINTEL = 5827,
@@ -502,6 +554,20 @@ enum Decoration {
DecorationMergeINTEL = 5834,
DecorationBankBitsINTEL = 5835,
DecorationForcePow2DepthINTEL = 5836,
DecorationBurstCoalesceINTEL = 5899,
DecorationCacheSizeINTEL = 5900,
DecorationDontStaticallyCoalesceINTEL = 5901,
DecorationPrefetchINTEL = 5902,
DecorationStallEnableINTEL = 5905,
DecorationFuseLoopsInFunctionINTEL = 5907,
DecorationAliasScopeINTEL = 5914,
DecorationNoAliasINTEL = 5915,
DecorationBufferLocationINTEL = 5921,
DecorationIOPipeStorageINTEL = 5944,
DecorationFunctionFloatingPointModeINTEL = 6080,
DecorationSingleElementVectorINTEL = 6085,
DecorationVectorComputeCallableFunctionINTEL = 6087,
DecorationMediaBlockIOINTEL = 6140,
DecorationMax = 0x7fffffff,
};
@@ -586,12 +652,18 @@ enum BuiltIn {
BuiltInLayerPerViewNV = 5279,
BuiltInMeshViewCountNV = 5280,
BuiltInMeshViewIndicesNV = 5281,
BuiltInBaryCoordKHR = 5286,
BuiltInBaryCoordNV = 5286,
BuiltInBaryCoordNoPerspKHR = 5287,
BuiltInBaryCoordNoPerspNV = 5287,
BuiltInFragSizeEXT = 5292,
BuiltInFragmentSizeNV = 5292,
BuiltInFragInvocationCountEXT = 5293,
BuiltInInvocationsPerPixelNV = 5293,
BuiltInPrimitivePointIndicesEXT = 5294,
BuiltInPrimitiveLineIndicesEXT = 5295,
BuiltInPrimitiveTriangleIndicesEXT = 5296,
BuiltInCullPrimitiveEXT = 5299,
BuiltInLaunchIdKHR = 5319,
BuiltInLaunchIdNV = 5319,
BuiltInLaunchSizeKHR = 5320,
@@ -617,6 +689,7 @@ enum BuiltIn {
BuiltInHitTNV = 5332,
BuiltInHitKindKHR = 5333,
BuiltInHitKindNV = 5333,
BuiltInCurrentRayTimeNV = 5334,
BuiltInIncomingRayFlagsKHR = 5351,
BuiltInIncomingRayFlagsNV = 5351,
BuiltInRayGeometryIndexKHR = 5352,
@@ -624,6 +697,7 @@ enum BuiltIn {
BuiltInSMCountNV = 5375,
BuiltInWarpIDNV = 5376,
BuiltInSMIDNV = 5377,
BuiltInCullMaskKHR = 6021,
BuiltInMax = 0x7fffffff,
};
@@ -656,6 +730,7 @@ enum LoopControlShift {
LoopControlLoopCoalesceINTELShift = 20,
LoopControlMaxInterleavingINTELShift = 21,
LoopControlSpeculatedIterationsINTELShift = 22,
LoopControlNoFusionINTELShift = 23,
LoopControlMax = 0x7fffffff,
};
@@ -677,6 +752,7 @@ enum LoopControlMask {
LoopControlLoopCoalesceINTELMask = 0x00100000,
LoopControlMaxInterleavingINTELMask = 0x00200000,
LoopControlSpeculatedIterationsINTELMask = 0x00400000,
LoopControlNoFusionINTELMask = 0x00800000,
};
enum FunctionControlShift {
@@ -684,6 +760,7 @@ enum FunctionControlShift {
FunctionControlDontInlineShift = 1,
FunctionControlPureShift = 2,
FunctionControlConstShift = 3,
FunctionControlOptNoneINTELShift = 16,
FunctionControlMax = 0x7fffffff,
};
@@ -693,6 +770,7 @@ enum FunctionControlMask {
FunctionControlDontInlineMask = 0x00000002,
FunctionControlPureMask = 0x00000004,
FunctionControlConstMask = 0x00000008,
FunctionControlOptNoneINTELMask = 0x00010000,
};
enum MemorySemanticsShift {
@@ -747,6 +825,8 @@ enum MemoryAccessShift {
MemoryAccessMakePointerVisibleKHRShift = 4,
MemoryAccessNonPrivatePointerShift = 5,
MemoryAccessNonPrivatePointerKHRShift = 5,
MemoryAccessAliasScopeINTELMaskShift = 16,
MemoryAccessNoAliasINTELMaskShift = 17,
MemoryAccessMax = 0x7fffffff,
};
@@ -761,6 +841,8 @@ enum MemoryAccessMask {
MemoryAccessMakePointerVisibleKHRMask = 0x00000010,
MemoryAccessNonPrivatePointerMask = 0x00000020,
MemoryAccessNonPrivatePointerKHRMask = 0x00000020,
MemoryAccessAliasScopeINTELMaskMask = 0x00010000,
MemoryAccessNoAliasINTELMaskMask = 0x00020000,
};
enum Scope {
@@ -873,9 +955,13 @@ enum Capability {
CapabilityGroupNonUniformQuad = 68,
CapabilityShaderLayer = 69,
CapabilityShaderViewportIndex = 70,
CapabilityUniformDecoration = 71,
CapabilityFragmentShadingRateKHR = 4422,
CapabilitySubgroupBallotKHR = 4423,
CapabilityDrawParameters = 4427,
CapabilityWorkgroupMemoryExplicitLayoutKHR = 4428,
CapabilityWorkgroupMemoryExplicitLayout8BitAccessKHR = 4429,
CapabilityWorkgroupMemoryExplicitLayout16BitAccessKHR = 4430,
CapabilitySubgroupVoteKHR = 4431,
CapabilityStorageBuffer16BitAccess = 4433,
CapabilityStorageUniformBufferBlock16 = 4433,
@@ -918,6 +1004,8 @@ enum Capability {
CapabilityFragmentFullyCoveredEXT = 5265,
CapabilityMeshShadingNV = 5266,
CapabilityImageFootprintNV = 5282,
CapabilityMeshShadingEXT = 5283,
CapabilityFragmentBarycentricKHR = 5284,
CapabilityFragmentBarycentricNV = 5284,
CapabilityComputeDerivativeGroupQuadsNV = 5288,
CapabilityFragmentDensityEXT = 5291,
@@ -948,6 +1036,7 @@ enum Capability {
CapabilityStorageTexelBufferArrayNonUniformIndexing = 5312,
CapabilityStorageTexelBufferArrayNonUniformIndexingEXT = 5312,
CapabilityRayTracingNV = 5340,
CapabilityRayTracingMotionBlurNV = 5341,
CapabilityVulkanMemoryModel = 5345,
CapabilityVulkanMemoryModelKHR = 5345,
CapabilityVulkanMemoryModelDeviceScope = 5346,
@@ -961,26 +1050,67 @@ enum Capability {
CapabilityFragmentShaderShadingRateInterlockEXT = 5372,
CapabilityShaderSMBuiltinsNV = 5373,
CapabilityFragmentShaderPixelInterlockEXT = 5378,
CapabilityDemoteToHelperInvocation = 5379,
CapabilityDemoteToHelperInvocationEXT = 5379,
CapabilityBindlessTextureNV = 5390,
CapabilitySubgroupShuffleINTEL = 5568,
CapabilitySubgroupBufferBlockIOINTEL = 5569,
CapabilitySubgroupImageBlockIOINTEL = 5570,
CapabilitySubgroupImageMediaBlockIOINTEL = 5579,
CapabilityRoundToInfinityINTEL = 5582,
CapabilityFloatingPointModeINTEL = 5583,
CapabilityIntegerFunctions2INTEL = 5584,
CapabilityFunctionPointersINTEL = 5603,
CapabilityIndirectReferencesINTEL = 5604,
CapabilityAsmINTEL = 5606,
CapabilityAtomicFloat32MinMaxEXT = 5612,
CapabilityAtomicFloat64MinMaxEXT = 5613,
CapabilityAtomicFloat16MinMaxEXT = 5616,
CapabilityVectorComputeINTEL = 5617,
CapabilityVectorAnyINTEL = 5619,
CapabilityExpectAssumeKHR = 5629,
CapabilitySubgroupAvcMotionEstimationINTEL = 5696,
CapabilitySubgroupAvcMotionEstimationIntraINTEL = 5697,
CapabilitySubgroupAvcMotionEstimationChromaINTEL = 5698,
CapabilityVariableLengthArrayINTEL = 5817,
CapabilityFunctionFloatControlINTEL = 5821,
CapabilityFPGAMemoryAttributesINTEL = 5824,
CapabilityFPFastMathModeINTEL = 5837,
CapabilityArbitraryPrecisionIntegersINTEL = 5844,
CapabilityArbitraryPrecisionFloatingPointINTEL = 5845,
CapabilityUnstructuredLoopControlsINTEL = 5886,
CapabilityFPGALoopControlsINTEL = 5888,
CapabilityKernelAttributesINTEL = 5892,
CapabilityFPGAKernelAttributesINTEL = 5897,
CapabilityFPGAMemoryAccessesINTEL = 5898,
CapabilityFPGAClusterAttributesINTEL = 5904,
CapabilityLoopFuseINTEL = 5906,
CapabilityMemoryAccessAliasingINTEL = 5910,
CapabilityFPGABufferLocationINTEL = 5920,
CapabilityArbitraryPrecisionFixedPointINTEL = 5922,
CapabilityUSMStorageClassesINTEL = 5935,
CapabilityIOPipesINTEL = 5943,
CapabilityBlockingPipesINTEL = 5945,
CapabilityFPGARegINTEL = 5948,
CapabilityDotProductInputAll = 6016,
CapabilityDotProductInputAllKHR = 6016,
CapabilityDotProductInput4x8Bit = 6017,
CapabilityDotProductInput4x8BitKHR = 6017,
CapabilityDotProductInput4x8BitPacked = 6018,
CapabilityDotProductInput4x8BitPackedKHR = 6018,
CapabilityDotProduct = 6019,
CapabilityDotProductKHR = 6019,
CapabilityRayCullMaskKHR = 6020,
CapabilityBitInstructions = 6025,
CapabilityGroupNonUniformRotateKHR = 6026,
CapabilityAtomicFloat32AddEXT = 6033,
CapabilityAtomicFloat64AddEXT = 6034,
CapabilityLongConstantCompositeINTEL = 6089,
CapabilityOptNoneINTEL = 6094,
CapabilityAtomicFloat16AddEXT = 6095,
CapabilityDebugInfoModuleINTEL = 6114,
CapabilitySplitBarrierINTEL = 6141,
CapabilityGroupUniformArithmeticKHR = 6400,
CapabilityMax = 0x7fffffff,
};
@@ -1047,6 +1177,44 @@ enum FragmentShadingRateMask {
FragmentShadingRateHorizontal4PixelsMask = 0x00000008,
};
enum FPDenormMode {
FPDenormModePreserve = 0,
FPDenormModeFlushToZero = 1,
FPDenormModeMax = 0x7fffffff,
};
enum FPOperationMode {
FPOperationModeIEEE = 0,
FPOperationModeALT = 1,
FPOperationModeMax = 0x7fffffff,
};
enum QuantizationModes {
QuantizationModesTRN = 0,
QuantizationModesTRN_ZERO = 1,
QuantizationModesRND = 2,
QuantizationModesRND_ZERO = 3,
QuantizationModesRND_INF = 4,
QuantizationModesRND_MIN_INF = 5,
QuantizationModesRND_CONV = 6,
QuantizationModesRND_CONV_ODD = 7,
QuantizationModesMax = 0x7fffffff,
};
enum OverflowModes {
OverflowModesWRAP = 0,
OverflowModesSAT = 1,
OverflowModesSAT_ZERO = 2,
OverflowModesSAT_SYM = 3,
OverflowModesMax = 0x7fffffff,
};
enum PackedVectorFormat {
PackedVectorFormatPackedVectorFormat4x8Bit = 0,
PackedVectorFormatPackedVectorFormat4x8BitKHR = 0,
PackedVectorFormatMax = 0x7fffffff,
};
enum Op {
OpNop = 0,
OpUndef = 1,
@@ -1398,12 +1566,25 @@ enum Op {
OpSubgroupAllKHR = 4428,
OpSubgroupAnyKHR = 4429,
OpSubgroupAllEqualKHR = 4430,
OpGroupNonUniformRotateKHR = 4431,
OpSubgroupReadInvocationKHR = 4432,
OpTraceRayKHR = 4445,
OpExecuteCallableKHR = 4446,
OpConvertUToAccelerationStructureKHR = 4447,
OpIgnoreIntersectionKHR = 4448,
OpTerminateRayKHR = 4449,
OpSDot = 4450,
OpSDotKHR = 4450,
OpUDot = 4451,
OpUDotKHR = 4451,
OpSUDot = 4452,
OpSUDotKHR = 4452,
OpSDotAccSat = 4453,
OpSDotAccSatKHR = 4453,
OpUDotAccSat = 4454,
OpUDotAccSatKHR = 4454,
OpSUDotAccSat = 4455,
OpSUDotAccSatKHR = 4455,
OpTypeRayQueryKHR = 4472,
OpRayQueryInitializeKHR = 4473,
OpRayQueryTerminateKHR = 4474,
@@ -1423,6 +1604,8 @@ enum Op {
OpFragmentFetchAMD = 5012,
OpReadClockKHR = 5056,
OpImageSampleFootprintNV = 5283,
OpEmitMeshTasksEXT = 5294,
OpSetMeshOutputsEXT = 5295,
OpGroupNonUniformPartitionNV = 5296,
OpWritePackedPrimitiveIndices4x8NV = 5299,
OpReportIntersectionKHR = 5334,
@@ -1430,6 +1613,8 @@ enum Op {
OpIgnoreIntersectionNV = 5335,
OpTerminateRayNV = 5336,
OpTraceNV = 5337,
OpTraceMotionNV = 5338,
OpTraceRayMotionNV = 5339,
OpTypeAccelerationStructureKHR = 5341,
OpTypeAccelerationStructureNV = 5341,
OpExecuteCallableNV = 5344,
@@ -1440,8 +1625,16 @@ enum Op {
OpCooperativeMatrixLengthNV = 5362,
OpBeginInvocationInterlockEXT = 5364,
OpEndInvocationInterlockEXT = 5365,
OpDemoteToHelperInvocation = 5380,
OpDemoteToHelperInvocationEXT = 5380,
OpIsHelperInvocationEXT = 5381,
OpConvertUToImageNV = 5391,
OpConvertUToSamplerNV = 5392,
OpConvertImageToUNV = 5393,
OpConvertSamplerToUNV = 5394,
OpConvertUToSampledImageNV = 5395,
OpConvertSampledImageToUNV = 5396,
OpSamplerImageAddressingModeNV = 5397,
OpSubgroupShuffleINTEL = 5571,
OpSubgroupShuffleDownINTEL = 5572,
OpSubgroupShuffleUpINTEL = 5573,
@@ -1466,8 +1659,15 @@ enum Op {
OpUSubSatINTEL = 5596,
OpIMul32x16INTEL = 5597,
OpUMul32x16INTEL = 5598,
OpFunctionPointerINTEL = 5600,
OpConstantFunctionPointerINTEL = 5600,
OpFunctionPointerCallINTEL = 5601,
OpAsmTargetINTEL = 5609,
OpAsmINTEL = 5610,
OpAsmCallINTEL = 5611,
OpAtomicFMinEXT = 5614,
OpAtomicFMaxEXT = 5615,
OpAssumeTrueKHR = 5630,
OpExpectKHR = 5631,
OpDecorateString = 5632,
OpDecorateStringGOOGLE = 5632,
OpMemberDecorateString = 5633,
@@ -1590,7 +1790,67 @@ enum Op {
OpSubgroupAvcSicGetPackedSkcLumaCountThresholdINTEL = 5814,
OpSubgroupAvcSicGetPackedSkcLumaSumThresholdINTEL = 5815,
OpSubgroupAvcSicGetInterRawSadsINTEL = 5816,
OpVariableLengthArrayINTEL = 5818,
OpSaveMemoryINTEL = 5819,
OpRestoreMemoryINTEL = 5820,
OpArbitraryFloatSinCosPiINTEL = 5840,
OpArbitraryFloatCastINTEL = 5841,
OpArbitraryFloatCastFromIntINTEL = 5842,
OpArbitraryFloatCastToIntINTEL = 5843,
OpArbitraryFloatAddINTEL = 5846,
OpArbitraryFloatSubINTEL = 5847,
OpArbitraryFloatMulINTEL = 5848,
OpArbitraryFloatDivINTEL = 5849,
OpArbitraryFloatGTINTEL = 5850,
OpArbitraryFloatGEINTEL = 5851,
OpArbitraryFloatLTINTEL = 5852,
OpArbitraryFloatLEINTEL = 5853,
OpArbitraryFloatEQINTEL = 5854,
OpArbitraryFloatRecipINTEL = 5855,
OpArbitraryFloatRSqrtINTEL = 5856,
OpArbitraryFloatCbrtINTEL = 5857,
OpArbitraryFloatHypotINTEL = 5858,
OpArbitraryFloatSqrtINTEL = 5859,
OpArbitraryFloatLogINTEL = 5860,
OpArbitraryFloatLog2INTEL = 5861,
OpArbitraryFloatLog10INTEL = 5862,
OpArbitraryFloatLog1pINTEL = 5863,
OpArbitraryFloatExpINTEL = 5864,
OpArbitraryFloatExp2INTEL = 5865,
OpArbitraryFloatExp10INTEL = 5866,
OpArbitraryFloatExpm1INTEL = 5867,
OpArbitraryFloatSinINTEL = 5868,
OpArbitraryFloatCosINTEL = 5869,
OpArbitraryFloatSinCosINTEL = 5870,
OpArbitraryFloatSinPiINTEL = 5871,
OpArbitraryFloatCosPiINTEL = 5872,
OpArbitraryFloatASinINTEL = 5873,
OpArbitraryFloatASinPiINTEL = 5874,
OpArbitraryFloatACosINTEL = 5875,
OpArbitraryFloatACosPiINTEL = 5876,
OpArbitraryFloatATanINTEL = 5877,
OpArbitraryFloatATanPiINTEL = 5878,
OpArbitraryFloatATan2INTEL = 5879,
OpArbitraryFloatPowINTEL = 5880,
OpArbitraryFloatPowRINTEL = 5881,
OpArbitraryFloatPowNINTEL = 5882,
OpLoopControlINTEL = 5887,
OpAliasDomainDeclINTEL = 5911,
OpAliasScopeDeclINTEL = 5912,
OpAliasScopeListDeclINTEL = 5913,
OpFixedSqrtINTEL = 5923,
OpFixedRecipINTEL = 5924,
OpFixedRsqrtINTEL = 5925,
OpFixedSinINTEL = 5926,
OpFixedCosINTEL = 5927,
OpFixedSinCosINTEL = 5928,
OpFixedSinPiINTEL = 5929,
OpFixedCosPiINTEL = 5930,
OpFixedSinCosPiINTEL = 5931,
OpFixedLogINTEL = 5932,
OpFixedExpINTEL = 5933,
OpPtrCastToCrossWorkgroupINTEL = 5934,
OpCrossWorkgroupCastToPtrINTEL = 5938,
OpReadPipeBlockingINTEL = 5946,
OpWritePipeBlockingINTEL = 5947,
OpFPGARegINTEL = 5949,
@@ -1612,10 +1872,27 @@ enum Op {
OpRayQueryGetIntersectionObjectToWorldKHR = 6031,
OpRayQueryGetIntersectionWorldToObjectKHR = 6032,
OpAtomicFAddEXT = 6035,
OpTypeBufferSurfaceINTEL = 6086,
OpTypeStructContinuedINTEL = 6090,
OpConstantCompositeContinuedINTEL = 6091,
OpSpecConstantCompositeContinuedINTEL = 6092,
OpControlBarrierArriveINTEL = 6142,
OpControlBarrierWaitINTEL = 6143,
OpGroupIMulKHR = 6401,
OpGroupFMulKHR = 6402,
OpGroupBitwiseAndKHR = 6403,
OpGroupBitwiseOrKHR = 6404,
OpGroupBitwiseXorKHR = 6405,
OpGroupLogicalAndKHR = 6406,
OpGroupLogicalOrKHR = 6407,
OpGroupLogicalXorKHR = 6408,
OpMax = 0x7fffffff,
};
#ifdef SPV_ENABLE_UTILITY_CODE
#ifndef __cplusplus
#include <stdbool.h>
#endif
inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) {
*hasResult = *hasResultType = false;
switch (opcode) {
@@ -1970,12 +2247,19 @@ inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) {
case OpSubgroupAllKHR: *hasResult = true; *hasResultType = true; break;
case OpSubgroupAnyKHR: *hasResult = true; *hasResultType = true; break;
case OpSubgroupAllEqualKHR: *hasResult = true; *hasResultType = true; break;
case OpGroupNonUniformRotateKHR: *hasResult = true; *hasResultType = true; break;
case OpSubgroupReadInvocationKHR: *hasResult = true; *hasResultType = true; 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 OpSDot: *hasResult = true; *hasResultType = true; break;
case OpUDot: *hasResult = true; *hasResultType = true; break;
case OpSUDot: *hasResult = true; *hasResultType = true; break;
case OpSDotAccSat: *hasResult = true; *hasResultType = true; break;
case OpUDotAccSat: *hasResult = true; *hasResultType = true; break;
case OpSUDotAccSat: *hasResult = true; *hasResultType = true; break;
case OpTypeRayQueryKHR: *hasResult = true; *hasResultType = false; break;
case OpRayQueryInitializeKHR: *hasResult = false; *hasResultType = false; break;
case OpRayQueryTerminateKHR: *hasResult = false; *hasResultType = false; break;
@@ -1995,12 +2279,16 @@ 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;
case OpIgnoreIntersectionNV: *hasResult = false; *hasResultType = false; break;
case OpTerminateRayNV: *hasResult = false; *hasResultType = false; break;
case OpTraceNV: *hasResult = false; *hasResultType = false; break;
case OpTraceMotionNV: *hasResult = false; *hasResultType = false; break;
case OpTraceRayMotionNV: *hasResult = false; *hasResultType = false; break;
case OpTypeAccelerationStructureNV: *hasResult = true; *hasResultType = false; break;
case OpExecuteCallableNV: *hasResult = false; *hasResultType = false; break;
case OpTypeCooperativeMatrixNV: *hasResult = true; *hasResultType = false; break;
@@ -2010,8 +2298,15 @@ inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) {
case OpCooperativeMatrixLengthNV: *hasResult = true; *hasResultType = true; break;
case OpBeginInvocationInterlockEXT: *hasResult = false; *hasResultType = false; break;
case OpEndInvocationInterlockEXT: *hasResult = false; *hasResultType = false; break;
case OpDemoteToHelperInvocationEXT: *hasResult = false; *hasResultType = false; break;
case OpDemoteToHelperInvocation: *hasResult = false; *hasResultType = false; break;
case OpIsHelperInvocationEXT: *hasResult = true; *hasResultType = true; break;
case OpConvertUToImageNV: *hasResult = true; *hasResultType = true; break;
case OpConvertUToSamplerNV: *hasResult = true; *hasResultType = true; break;
case OpConvertImageToUNV: *hasResult = true; *hasResultType = true; break;
case OpConvertSamplerToUNV: *hasResult = true; *hasResultType = true; break;
case OpConvertUToSampledImageNV: *hasResult = true; *hasResultType = true; break;
case OpConvertSampledImageToUNV: *hasResult = true; *hasResultType = true; break;
case OpSamplerImageAddressingModeNV: *hasResult = false; *hasResultType = false; break;
case OpSubgroupShuffleINTEL: *hasResult = true; *hasResultType = true; break;
case OpSubgroupShuffleDownINTEL: *hasResult = true; *hasResultType = true; break;
case OpSubgroupShuffleUpINTEL: *hasResult = true; *hasResultType = true; break;
@@ -2036,8 +2331,15 @@ 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 OpConstantFunctionPointerINTEL: *hasResult = true; *hasResultType = true; break;
case OpFunctionPointerCallINTEL: *hasResult = true; *hasResultType = true; break;
case OpAsmTargetINTEL: *hasResult = true; *hasResultType = true; break;
case OpAsmINTEL: *hasResult = true; *hasResultType = true; break;
case OpAsmCallINTEL: *hasResult = true; *hasResultType = true; break;
case OpAtomicFMinEXT: *hasResult = true; *hasResultType = true; break;
case OpAtomicFMaxEXT: *hasResult = true; *hasResultType = true; break;
case OpAssumeTrueKHR: *hasResult = false; *hasResultType = false; break;
case OpExpectKHR: *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;
@@ -2158,7 +2460,67 @@ 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 OpVariableLengthArrayINTEL: *hasResult = true; *hasResultType = true; break;
case OpSaveMemoryINTEL: *hasResult = true; *hasResultType = true; break;
case OpRestoreMemoryINTEL: *hasResult = false; *hasResultType = false; break;
case OpArbitraryFloatSinCosPiINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatCastINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatCastFromIntINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatCastToIntINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatAddINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatSubINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatMulINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatDivINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatGTINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatGEINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatLTINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatLEINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatEQINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatRecipINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatRSqrtINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatCbrtINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatHypotINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatSqrtINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatLogINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatLog2INTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatLog10INTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatLog1pINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatExpINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatExp2INTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatExp10INTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatExpm1INTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatSinINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatCosINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatSinCosINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatSinPiINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatCosPiINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatASinINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatASinPiINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatACosINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatACosPiINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatATanINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatATanPiINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatATan2INTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatPowINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatPowRINTEL: *hasResult = true; *hasResultType = true; break;
case OpArbitraryFloatPowNINTEL: *hasResult = true; *hasResultType = true; break;
case OpLoopControlINTEL: *hasResult = false; *hasResultType = false; break;
case OpAliasDomainDeclINTEL: *hasResult = true; *hasResultType = false; break;
case OpAliasScopeDeclINTEL: *hasResult = true; *hasResultType = false; break;
case OpAliasScopeListDeclINTEL: *hasResult = true; *hasResultType = false; break;
case OpFixedSqrtINTEL: *hasResult = true; *hasResultType = true; break;
case OpFixedRecipINTEL: *hasResult = true; *hasResultType = true; break;
case OpFixedRsqrtINTEL: *hasResult = true; *hasResultType = true; break;
case OpFixedSinINTEL: *hasResult = true; *hasResultType = true; break;
case OpFixedCosINTEL: *hasResult = true; *hasResultType = true; break;
case OpFixedSinCosINTEL: *hasResult = true; *hasResultType = true; break;
case OpFixedSinPiINTEL: *hasResult = true; *hasResultType = true; break;
case OpFixedCosPiINTEL: *hasResult = true; *hasResultType = true; break;
case OpFixedSinCosPiINTEL: *hasResult = true; *hasResultType = true; break;
case OpFixedLogINTEL: *hasResult = true; *hasResultType = true; break;
case OpFixedExpINTEL: *hasResult = true; *hasResultType = true; break;
case OpPtrCastToCrossWorkgroupINTEL: *hasResult = true; *hasResultType = true; break;
case OpCrossWorkgroupCastToPtrINTEL: *hasResult = true; *hasResultType = true; break;
case OpReadPipeBlockingINTEL: *hasResult = true; *hasResultType = true; break;
case OpWritePipeBlockingINTEL: *hasResult = true; *hasResultType = true; break;
case OpFPGARegINTEL: *hasResult = true; *hasResultType = true; break;
@@ -2180,6 +2542,20 @@ inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) {
case OpRayQueryGetIntersectionObjectToWorldKHR: *hasResult = true; *hasResultType = true; break;
case OpRayQueryGetIntersectionWorldToObjectKHR: *hasResult = true; *hasResultType = true; break;
case OpAtomicFAddEXT: *hasResult = true; *hasResultType = true; break;
case OpTypeBufferSurfaceINTEL: *hasResult = true; *hasResultType = false; break;
case OpTypeStructContinuedINTEL: *hasResult = false; *hasResultType = false; break;
case OpConstantCompositeContinuedINTEL: *hasResult = false; *hasResultType = false; break;
case OpSpecConstantCompositeContinuedINTEL: *hasResult = false; *hasResultType = false; break;
case OpControlBarrierArriveINTEL: *hasResult = false; *hasResultType = false; break;
case OpControlBarrierWaitINTEL: *hasResult = false; *hasResultType = false; break;
case OpGroupIMulKHR: *hasResult = true; *hasResultType = true; break;
case OpGroupFMulKHR: *hasResult = true; *hasResultType = true; break;
case OpGroupBitwiseAndKHR: *hasResult = true; *hasResultType = true; break;
case OpGroupBitwiseOrKHR: *hasResult = true; *hasResultType = true; break;
case OpGroupBitwiseXorKHR: *hasResult = true; *hasResultType = true; break;
case OpGroupLogicalAndKHR: *hasResult = true; *hasResultType = true; break;
case OpGroupLogicalOrKHR: *hasResult = true; *hasResultType = true; break;
case OpGroupLogicalXorKHR: *hasResult = true; *hasResultType = true; break;
}
}
#endif /* SPV_ENABLE_UTILITY_CODE */
src/libraries/spirv_cross/spirv_cfg.cpp
+27 -6
View File
@@ -306,15 +306,36 @@ bool CFG::node_terminates_control_flow_in_sub_graph(BlockID from, BlockID to) co
bool true_path_ignore = false;
bool false_path_ignore = false;
if (ignore_block_id && dom.terminator == SPIRBlock::Select)
bool merges_to_nothing = dom.merge == SPIRBlock::MergeNone ||
(dom.merge == SPIRBlock::MergeSelection && dom.next_block &&
compiler.get<SPIRBlock>(dom.next_block).terminator == SPIRBlock::Unreachable) ||
(dom.merge == SPIRBlock::MergeLoop && dom.merge_block &&
compiler.get<SPIRBlock>(dom.merge_block).terminator == SPIRBlock::Unreachable);
if (dom.self == from || merges_to_nothing)
{
auto &true_block = compiler.get<SPIRBlock>(dom.true_block);
auto &false_block = compiler.get<SPIRBlock>(dom.false_block);
auto &ignore_block = compiler.get<SPIRBlock>(ignore_block_id);
true_path_ignore = compiler.execution_is_branchless(true_block, ignore_block);
false_path_ignore = compiler.execution_is_branchless(false_block, ignore_block);
// We can only ignore inner branchy paths if there is no merge,
// i.e. no code is generated afterwards. E.g. this allows us to elide continue:
// for (;;) { if (cond) { continue; } else { break; } }.
// Codegen here in SPIR-V will be something like either no merge if one path directly breaks, or
// we merge to Unreachable.
if (ignore_block_id && dom.terminator == SPIRBlock::Select)
{
auto &true_block = compiler.get<SPIRBlock>(dom.true_block);
auto &false_block = compiler.get<SPIRBlock>(dom.false_block);
auto &ignore_block = compiler.get<SPIRBlock>(ignore_block_id);
true_path_ignore = compiler.execution_is_branchless(true_block, ignore_block);
false_path_ignore = compiler.execution_is_branchless(false_block, ignore_block);
}
}
// Cases where we allow traversal. This serves as a proxy for post-dominance in a loop body.
// TODO: Might want to do full post-dominance analysis, but it's a lot of churn for something like this ...
// - We're the merge block of a selection construct. Jump to header.
// - We're the merge block of a loop. Jump to header.
// - Direct branch. Trivial.
// - Allow cases inside a branch if the header cannot merge execution before loop exit.
if ((dom.merge == SPIRBlock::MergeSelection && dom.next_block == to) ||
(dom.merge == SPIRBlock::MergeLoop && dom.merge_block == to) ||
(dom.terminator == SPIRBlock::Direct && dom.next_block == to) ||
src/libraries/spirv_cross/spirv_cfg.hpp
+5
View File
@@ -59,6 +59,11 @@ public:
return 0;
}
bool is_reachable(uint32_t block) const
{
return visit_order.count(block) != 0;
}
uint32_t get_visit_order(uint32_t block) const
{
auto itr = visit_order.find(block);
src/libraries/spirv_cross/spirv_common.hpp
+77 -11
View File
@@ -24,7 +24,11 @@
#ifndef SPIRV_CROSS_COMMON_HPP
#define SPIRV_CROSS_COMMON_HPP
#ifndef SPV_ENABLE_UTILITY_CODE
#define SPV_ENABLE_UTILITY_CODE
#endif
#include "spirv.hpp"
#include "spirv_cross_containers.hpp"
#include "spirv_cross_error_handling.hpp"
#include <functional>
@@ -291,6 +295,20 @@ inline std::string convert_to_string(double t, char locale_radix_point)
return buf;
}
#if defined(__clang__) || defined(__GNUC__)
#pragma GCC diagnostic pop
#elif defined(_MSC_VER)
#pragma warning(pop)
#endif
class FloatFormatter
{
public:
virtual ~FloatFormatter() = default;
virtual std::string format_float(float value) = 0;
virtual std::string format_double(double value) = 0;
};
template <typename T>
struct ValueSaver
{
@@ -314,12 +332,6 @@ struct ValueSaver
T saved;
};
#if defined(__clang__) || defined(__GNUC__)
#pragma GCC diagnostic pop
#elif defined(_MSC_VER)
#pragma warning(pop)
#endif
struct Instruction
{
uint16_t op = 0;
@@ -536,6 +548,9 @@ struct SPIRType : IVariant
type = TypeType
};
spv::Op op = spv::Op::OpNop;
explicit SPIRType(spv::Op op_) : op(op_) {}
enum BaseType
{
Unknown,
@@ -606,7 +621,7 @@ struct SPIRType : IVariant
uint32_t sampled;
spv::ImageFormat format;
spv::AccessQualifier access;
} image;
} image = {};
// Structs can be declared multiple times if they are used as part of interface blocks.
// We want to detect this so that we only emit the struct definition once.
@@ -638,7 +653,10 @@ struct SPIRExtension : IVariant
SPV_AMD_shader_ballot,
SPV_AMD_shader_explicit_vertex_parameter,
SPV_AMD_shader_trinary_minmax,
SPV_AMD_gcn_shader
SPV_AMD_gcn_shader,
NonSemanticDebugPrintf,
NonSemanticShaderDebugInfo,
NonSemanticGeneric
};
explicit SPIRExtension(Extension ext_)
@@ -672,10 +690,12 @@ struct SPIREntryPoint
struct WorkgroupSize
{
uint32_t x = 0, y = 0, z = 0;
uint32_t id_x = 0, id_y = 0, id_z = 0;
uint32_t constant = 0; // Workgroup size can be expressed as a constant/spec-constant instead.
} 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;
};
@@ -689,7 +709,7 @@ struct SPIRExpression : IVariant
// Only created by the backend target to avoid creating tons of temporaries.
SPIRExpression(std::string expr, TypeID expression_type_, bool immutable_)
: expression(move(expr))
: expression(std::move(expr))
, expression_type(expression_type_)
, immutable(immutable_)
{
@@ -720,6 +740,9 @@ struct SPIRExpression : IVariant
// Whether or not this is an access chain expression.
bool access_chain = false;
// Whether or not gl_MeshVerticesEXT[].gl_Position (as a whole or .y) is referenced
bool access_meshlet_position_y = false;
// A list of expressions which this expression depends on.
SmallVector<ID> expression_dependencies;
@@ -770,7 +793,8 @@ struct SPIRBlock : IVariant
Unreachable, // Noop
Kill, // Discard
IgnoreIntersection, // Ray Tracing
TerminateRay // Ray Tracing
TerminateRay, // Ray Tracing
EmitMeshTasks // Mesh shaders
};
enum Merge
@@ -832,6 +856,13 @@ struct SPIRBlock : IVariant
BlockID false_block = 0;
BlockID default_block = 0;
// If terminator is EmitMeshTasksEXT.
struct
{
ID groups[3];
ID payload;
} mesh = {};
SmallVector<Instruction> ops;
struct Phi
@@ -1071,7 +1102,6 @@ struct SPIRVariable : IVariant
// Temporaries which can remain forwarded as long as this variable is not modified.
SmallVector<ID> dependees;
bool forwardable = true;
bool deferred_declaration = false;
bool phi_variable = false;
@@ -1563,6 +1593,8 @@ struct AccessChainMeta
bool storage_is_packed = false;
bool storage_is_invariant = false;
bool flattened_struct = false;
bool relaxed_precision = false;
bool access_meshlet_position_y = false;
};
enum ExtendedDecorations
@@ -1630,6 +1662,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
};
@@ -1640,6 +1678,7 @@ struct Meta
std::string alias;
std::string qualified_alias;
std::string hlsl_semantic;
std::string user_type;
Bitset decoration_flags;
spv::BuiltIn builtin_type = spv::BuiltInMax;
uint32_t location = 0;
@@ -1775,6 +1814,33 @@ static inline bool opcode_is_sign_invariant(spv::Op opcode)
}
}
static inline bool opcode_can_promote_integer_implicitly(spv::Op opcode)
{
switch (opcode)
{
case spv::OpSNegate:
case spv::OpNot:
case spv::OpBitwiseAnd:
case spv::OpBitwiseOr:
case spv::OpBitwiseXor:
case spv::OpShiftLeftLogical:
case spv::OpShiftRightLogical:
case spv::OpShiftRightArithmetic:
case spv::OpIAdd:
case spv::OpISub:
case spv::OpIMul:
case spv::OpSDiv:
case spv::OpUDiv:
case spv::OpSRem:
case spv::OpUMod:
case spv::OpSMod:
return true;
default:
return false;
}
}
struct SetBindingPair
{
uint32_t desc_set;
src/libraries/spirv_cross/spirv_cpp.cpp
+1 -6
View File
@@ -274,8 +274,6 @@ void CompilerCPP::emit_resources()
if (emitted)
statement("");
declare_undefined_values();
statement("inline void init(spirv_cross_shader& s)");
begin_scope();
statement(resource_type, "::init(s);");
@@ -338,11 +336,8 @@ string CompilerCPP::compile()
uint32_t pass_count = 0;
do
{
if (pass_count >= 3)
SPIRV_CROSS_THROW("Over 3 compilation loops detected. Must be a bug!");
resource_registrations.clear();
reset();
reset(pass_count);
// Move constructor for this type is broken on GCC 4.9 ...
buffer.reset();
src/libraries/spirv_cross/spirv_cross.cpp
+357 -125
View File
@@ -36,16 +36,16 @@ using namespace SPIRV_CROSS_NAMESPACE;
Compiler::Compiler(vector<uint32_t> ir_)
{
Parser parser(move(ir_));
Parser parser(std::move(ir_));
parser.parse();
set_ir(move(parser.get_parsed_ir()));
set_ir(std::move(parser.get_parsed_ir()));
}
Compiler::Compiler(const uint32_t *ir_, size_t word_count)
{
Parser parser(ir_, word_count);
parser.parse();
set_ir(move(parser.get_parsed_ir()));
set_ir(std::move(parser.get_parsed_ir()));
}
Compiler::Compiler(const ParsedIR &ir_)
@@ -55,12 +55,12 @@ Compiler::Compiler(const ParsedIR &ir_)
Compiler::Compiler(ParsedIR &&ir_)
{
set_ir(move(ir_));
set_ir(std::move(ir_));
}
void Compiler::set_ir(ParsedIR &&ir_)
{
ir = move(ir_);
ir = std::move(ir_);
parse_fixup();
}
@@ -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:
@@ -621,7 +627,22 @@ bool Compiler::is_matrix(const SPIRType &type) const
bool Compiler::is_array(const SPIRType &type) const
{
return !type.array.empty();
return type.op == OpTypeArray || type.op == OpTypeRuntimeArray;
}
bool Compiler::is_pointer(const SPIRType &type) const
{
return type.op == OpTypePointer && type.basetype != SPIRType::Unknown; // Ignore function pointers.
}
bool Compiler::is_physical_pointer(const SPIRType &type) const
{
return type.op == OpTypePointer && type.storage == StorageClassPhysicalStorageBuffer;
}
bool Compiler::is_runtime_size_array(const SPIRType &type)
{
return type.op == OpTypeRuntimeArray;
}
ShaderResources Compiler::get_shader_resources() const
@@ -719,7 +740,7 @@ bool Compiler::InterfaceVariableAccessHandler::handle(Op opcode, const uint32_t
case OpExtInst:
{
if (length < 5)
if (length < 3)
return false;
auto &extension_set = compiler.get<SPIRExtension>(args[2]);
switch (extension_set.ext)
@@ -852,7 +873,7 @@ unordered_set<VariableID> Compiler::get_active_interface_variables() const
void Compiler::set_enabled_interface_variables(std::unordered_set<VariableID> active_variables)
{
active_interface_variables = move(active_variables);
active_interface_variables = std::move(active_variables);
check_active_interface_variables = true;
}
@@ -985,57 +1006,74 @@ ShaderResources Compiler::get_shader_resources(const unordered_set<VariableID> *
// in the future.
res.push_constant_buffers.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
}
// Images
else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Image &&
type.image.sampled == 2)
else if (type.storage == StorageClassShaderRecordBufferKHR)
{
res.storage_images.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
}
// Separate images
else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Image &&
type.image.sampled == 1)
{
res.separate_images.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
}
// Separate samplers
else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::Sampler)
{
res.separate_samplers.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
}
// Textures
else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::SampledImage)
{
res.sampled_images.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
res.shader_record_buffers.push_back({ var.self, var.basetype, type.self, get_remapped_declared_block_name(var.self, ssbo_instance_name) });
}
// Atomic counters
else if (type.storage == StorageClassAtomicCounter)
{
res.atomic_counters.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
}
// Acceleration structures
else if (type.storage == StorageClassUniformConstant && type.basetype == SPIRType::AccelerationStructure)
else if (type.storage == StorageClassUniformConstant)
{
res.acceleration_structures.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
if (type.basetype == SPIRType::Image)
{
// Images
if (type.image.sampled == 2)
{
res.storage_images.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
}
// Separate images
else if (type.image.sampled == 1)
{
res.separate_images.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
}
}
// Separate samplers
else if (type.basetype == SPIRType::Sampler)
{
res.separate_samplers.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
}
// Textures
else if (type.basetype == SPIRType::SampledImage)
{
res.sampled_images.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
}
// Acceleration structures
else if (type.basetype == SPIRType::AccelerationStructure)
{
res.acceleration_structures.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
}
else
{
res.gl_plain_uniforms.push_back({ var.self, var.basetype, type.self, get_name(var.self) });
}
}
});
return res;
}
bool Compiler::type_is_block_like(const SPIRType &type) const
bool Compiler::type_is_top_level_block(const SPIRType &type) const
{
if (type.basetype != SPIRType::Struct)
return false;
return has_decoration(type.self, DecorationBlock) || has_decoration(type.self, DecorationBufferBlock);
}
if (has_decoration(type.self, DecorationBlock) || has_decoration(type.self, DecorationBufferBlock))
{
bool Compiler::type_is_block_like(const SPIRType &type) const
{
if (type_is_top_level_block(type))
return true;
}
// Block-like types may have Offset decorations.
for (uint32_t i = 0; i < uint32_t(type.member_types.size()); i++)
if (has_member_decoration(type.self, i, DecorationOffset))
return true;
if (type.basetype == SPIRType::Struct)
{
// Block-like types may have Offset decorations.
for (uint32_t i = 0; i < uint32_t(type.member_types.size()); i++)
if (has_member_decoration(type.self, i, DecorationOffset))
return true;
}
return false;
}
@@ -1050,10 +1088,12 @@ void Compiler::parse_fixup()
if (id.get_type() == TypeConstant)
{
auto &c = id.get<SPIRConstant>();
if (ir.meta[c.self].decoration.builtin && ir.meta[c.self].decoration.builtin_type == BuiltInWorkgroupSize)
if (has_decoration(c.self, DecorationBuiltIn) &&
BuiltIn(get_decoration(c.self, DecorationBuiltIn)) == BuiltInWorkgroupSize)
{
// In current SPIR-V, there can be just one constant like this.
// All entry points will receive the constant value.
// WorkgroupSize take precedence over LocalSizeId.
for (auto &entry : ir.entry_points)
{
entry.second.workgroup_size.constant = c.self;
@@ -1067,8 +1107,11 @@ void Compiler::parse_fixup()
{
auto &var = id.get<SPIRVariable>();
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);
}
@@ -1453,6 +1496,58 @@ bool Compiler::get_binary_offset_for_decoration(VariableID id, spv::Decoration d
return true;
}
bool Compiler::block_is_noop(const SPIRBlock &block) const
{
if (block.terminator != SPIRBlock::Direct)
return false;
auto &child = get<SPIRBlock>(block.next_block);
// If this block participates in PHI, the block isn't really noop.
for (auto &phi : block.phi_variables)
if (phi.parent == block.self || phi.parent == child.self)
return false;
for (auto &phi : child.phi_variables)
if (phi.parent == block.self)
return false;
// Verify all instructions have no semantic impact.
for (auto &i : block.ops)
{
auto op = static_cast<Op>(i.op);
switch (op)
{
// Non-Semantic instructions.
case OpLine:
case OpNoLine:
break;
case OpExtInst:
{
auto *ops = stream(i);
auto ext = get<SPIRExtension>(ops[2]).ext;
bool ext_is_nonsemantic_only =
ext == SPIRExtension::NonSemanticShaderDebugInfo ||
ext == SPIRExtension::SPV_debug_info ||
ext == SPIRExtension::NonSemanticGeneric;
if (!ext_is_nonsemantic_only)
return false;
break;
}
default:
return false;
}
}
return true;
}
bool Compiler::block_is_loop_candidate(const SPIRBlock &block, SPIRBlock::Method method) const
{
// Tried and failed.
@@ -1510,7 +1605,7 @@ bool Compiler::block_is_loop_candidate(const SPIRBlock &block, SPIRBlock::Method
{
// Empty loop header that just sets up merge target
// and branches to loop body.
bool ret = block.terminator == SPIRBlock::Direct && block.merge == SPIRBlock::MergeLoop && block.ops.empty();
bool ret = block.terminator == SPIRBlock::Direct && block.merge == SPIRBlock::MergeLoop && block_is_noop(block);
if (!ret)
return false;
@@ -1536,19 +1631,8 @@ bool Compiler::block_is_loop_candidate(const SPIRBlock &block, SPIRBlock::Method
ret = child.terminator == SPIRBlock::Select && child.merge == SPIRBlock::MergeNone &&
(positive_candidate || negative_candidate);
// If we have OpPhi which depends on branches which came from our own block,
// we need to flush phi variables in else block instead of a trivial break,
// so we cannot assume this is a for loop candidate.
if (ret)
{
for (auto &phi : block.phi_variables)
if (phi.parent == block.self || phi.parent == child.self)
return false;
for (auto &phi : child.phi_variables)
if (phi.parent == block.self)
return false;
auto *merge = maybe_get<SPIRBlock>(block.merge_block);
if (merge)
for (auto &phi : merge->phi_variables)
@@ -1573,15 +1657,10 @@ bool Compiler::execution_is_noop(const SPIRBlock &from, const SPIRBlock &to) con
if (start->self == to.self)
return true;
if (!start->ops.empty())
if (!block_is_noop(*start))
return false;
auto &next = get<SPIRBlock>(start->next_block);
// Flushing phi variables does not count as noop.
for (auto &phi : next.phi_variables)
if (phi.parent == start->self)
return false;
start = &next;
}
}
@@ -1675,6 +1754,11 @@ const SmallVector<SPIRBlock::Case> &Compiler::get_case_list(const SPIRBlock &blo
const auto &type = get<SPIRType>(var->basetype);
width = type.width;
}
else if (const auto *undef = maybe_get<SPIRUndef>(block.condition))
{
const auto &type = get<SPIRType>(undef->basetype);
width = type.width;
}
else
{
auto search = ir.load_type_width.find(block.condition);
@@ -2156,6 +2240,12 @@ void Compiler::set_execution_mode(ExecutionMode mode, uint32_t arg0, uint32_t ar
execution.workgroup_size.z = arg2;
break;
case ExecutionModeLocalSizeId:
execution.workgroup_size.id_x = arg0;
execution.workgroup_size.id_y = arg1;
execution.workgroup_size.id_z = arg2;
break;
case ExecutionModeInvocations:
execution.invocations = arg0;
break;
@@ -2164,6 +2254,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;
}
@@ -2183,6 +2277,7 @@ uint32_t Compiler::get_work_group_size_specialization_constants(SpecializationCo
y = { 0, 0 };
z = { 0, 0 };
// WorkgroupSize builtin takes precedence over LocalSize / LocalSizeId.
if (execution.workgroup_size.constant != 0)
{
auto &c = get<SPIRConstant>(execution.workgroup_size.constant);
@@ -2205,6 +2300,29 @@ uint32_t Compiler::get_work_group_size_specialization_constants(SpecializationCo
z.constant_id = get_decoration(c.m.c[0].id[2], DecorationSpecId);
}
}
else if (execution.flags.get(ExecutionModeLocalSizeId))
{
auto &cx = get<SPIRConstant>(execution.workgroup_size.id_x);
if (cx.specialization)
{
x.id = execution.workgroup_size.id_x;
x.constant_id = get_decoration(execution.workgroup_size.id_x, DecorationSpecId);
}
auto &cy = get<SPIRConstant>(execution.workgroup_size.id_y);
if (cy.specialization)
{
y.id = execution.workgroup_size.id_y;
y.constant_id = get_decoration(execution.workgroup_size.id_y, DecorationSpecId);
}
auto &cz = get<SPIRConstant>(execution.workgroup_size.id_z);
if (cz.specialization)
{
z.id = execution.workgroup_size.id_z;
z.constant_id = get_decoration(execution.workgroup_size.id_z, DecorationSpecId);
}
}
return execution.workgroup_size.constant;
}
@@ -2214,15 +2332,42 @@ uint32_t Compiler::get_execution_mode_argument(spv::ExecutionMode mode, uint32_t
auto &execution = get_entry_point();
switch (mode)
{
case ExecutionModeLocalSizeId:
if (execution.flags.get(ExecutionModeLocalSizeId))
{
switch (index)
{
case 0:
return execution.workgroup_size.id_x;
case 1:
return execution.workgroup_size.id_y;
case 2:
return execution.workgroup_size.id_z;
default:
return 0;
}
}
else
return 0;
case ExecutionModeLocalSize:
switch (index)
{
case 0:
return execution.workgroup_size.x;
if (execution.flags.get(ExecutionModeLocalSizeId) && execution.workgroup_size.id_x != 0)
return get<SPIRConstant>(execution.workgroup_size.id_x).scalar();
else
return execution.workgroup_size.x;
case 1:
return execution.workgroup_size.y;
if (execution.flags.get(ExecutionModeLocalSizeId) && execution.workgroup_size.id_y != 0)
return get<SPIRConstant>(execution.workgroup_size.id_y).scalar();
else
return execution.workgroup_size.y;
case 2:
return execution.workgroup_size.z;
if (execution.flags.get(ExecutionModeLocalSizeId) && execution.workgroup_size.id_z != 0)
return get<SPIRConstant>(execution.workgroup_size.id_z).scalar();
else
return execution.workgroup_size.z;
default:
return 0;
}
@@ -2233,6 +2378,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;
}
@@ -2261,6 +2409,11 @@ bool Compiler::is_tessellation_shader() const
return is_tessellation_shader(get_execution_model());
}
bool Compiler::is_tessellating_triangles() const
{
return get_execution_mode_bitset().get(ExecutionModeTriangles);
}
void Compiler::set_remapped_variable_state(VariableID id, bool remap_enable)
{
get<SPIRVariable>(id).remapped_variable = remap_enable;
@@ -2295,6 +2448,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
@@ -2449,7 +2615,7 @@ void Compiler::CombinedImageSamplerHandler::push_remap_parameters(const SPIRFunc
unordered_map<uint32_t, uint32_t> remapping;
for (uint32_t i = 0; i < length; i++)
remapping[func.arguments[i].id] = remap_parameter(args[i]);
parameter_remapping.push(move(remapping));
parameter_remapping.push(std::move(remapping));
}
void Compiler::CombinedImageSamplerHandler::pop_remap_parameters()
@@ -2579,8 +2745,8 @@ void Compiler::CombinedImageSamplerHandler::register_combined_image_sampler(SPIR
auto ptr_type_id = id + 1;
auto combined_id = id + 2;
auto &base = compiler.expression_type(image_id);
auto &type = compiler.set<SPIRType>(type_id);
auto &ptr_type = compiler.set<SPIRType>(ptr_type_id);
auto &type = compiler.set<SPIRType>(type_id, OpTypeSampledImage);
auto &ptr_type = compiler.set<SPIRType>(ptr_type_id, OpTypePointer);
type = base;
type.self = type_id;
@@ -2839,7 +3005,7 @@ bool Compiler::CombinedImageSamplerHandler::handle(Op opcode, const uint32_t *ar
{
// Have to invent the sampled image type.
sampled_type = compiler.ir.increase_bound_by(1);
auto &type = compiler.set<SPIRType>(sampled_type);
auto &type = compiler.set<SPIRType>(sampled_type, OpTypeSampledImage);
type = compiler.expression_type(args[2]);
type.self = sampled_type;
type.basetype = SPIRType::SampledImage;
@@ -2858,7 +3024,7 @@ bool Compiler::CombinedImageSamplerHandler::handle(Op opcode, const uint32_t *ar
// Make a new type, pointer to OpTypeSampledImage, so we can make a variable of this type.
// We will probably have this type lying around, but it doesn't hurt to make duplicates for internal purposes.
auto &type = compiler.set<SPIRType>(type_id);
auto &type = compiler.set<SPIRType>(type_id, OpTypePointer);
auto &base = compiler.get<SPIRType>(sampled_type);
type = base;
type.pointer = true;
@@ -2904,11 +3070,10 @@ VariableID Compiler::build_dummy_sampler_for_combined_images()
auto ptr_type_id = offset + 1;
auto var_id = offset + 2;
SPIRType sampler_type;
auto &sampler = set<SPIRType>(type_id);
auto &sampler = set<SPIRType>(type_id, OpTypeSampler);
sampler.basetype = SPIRType::Sampler;
auto &ptr_sampler = set<SPIRType>(ptr_type_id);
auto &ptr_sampler = set<SPIRType>(ptr_type_id, OpTypePointer);
ptr_sampler = sampler;
ptr_sampler.self = type_id;
ptr_sampler.storage = StorageClassUniformConstant;
@@ -3111,8 +3276,8 @@ void Compiler::AnalyzeVariableScopeAccessHandler::notify_variable_access(uint32_
return;
// Access chains used in multiple blocks mean hoisting all the variables used to construct the access chain as not all backends can use pointers.
auto itr = access_chain_children.find(id);
if (itr != end(access_chain_children))
auto itr = rvalue_forward_children.find(id);
if (itr != end(rvalue_forward_children))
for (auto child_id : itr->second)
notify_variable_access(child_id, block);
@@ -3165,7 +3330,20 @@ bool Compiler::AnalyzeVariableScopeAccessHandler::handle(spv::Op op, const uint3
// Keep track of the types of temporaries, so we can hoist them out as necessary.
uint32_t result_type, result_id;
if (compiler.instruction_to_result_type(result_type, result_id, op, args, length))
{
// For some opcodes, we will need to override the result id.
// If we need to hoist the temporary, the temporary type is the input, not the result.
// FIXME: This will likely break with OpCopyObject + hoisting, but we'll have to
// solve it if we ever get there ...
if (op == OpConvertUToAccelerationStructureKHR)
{
auto itr = result_id_to_type.find(args[2]);
if (itr != result_id_to_type.end())
result_type = itr->second;
}
result_id_to_type[result_id] = result_type;
}
switch (op)
{
@@ -3207,14 +3385,14 @@ bool Compiler::AnalyzeVariableScopeAccessHandler::handle(spv::Op op, const uint3
if (var)
{
accessed_variables_to_block[var->self].insert(current_block->self);
access_chain_children[args[1]].insert(var->self);
rvalue_forward_children[args[1]].insert(var->self);
}
// args[2] might be another access chain we have to track use of.
for (uint32_t i = 2; i < length; i++)
{
notify_variable_access(args[i], current_block->self);
access_chain_children[args[1]].insert(args[i]);
rvalue_forward_children[args[1]].insert(args[i]);
}
// Also keep track of the access chain pointer itself.
@@ -3296,6 +3474,12 @@ bool Compiler::AnalyzeVariableScopeAccessHandler::handle(spv::Op op, const uint3
// Might be an access chain we have to track use of.
notify_variable_access(args[2], current_block->self);
// If we're loading an opaque type we cannot lower it to a temporary,
// we must defer access of args[2] until it's used.
auto &type = compiler.get<SPIRType>(args[0]);
if (compiler.type_is_opaque_value(type))
rvalue_forward_children[args[1]].insert(args[2]);
break;
}
@@ -3667,6 +3851,7 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry, AnalyzeVariableScopeA
DominatorBuilder builder(cfg);
auto &blocks = var.second;
auto &type = expression_type(var.first);
BlockID potential_continue_block = 0;
// Figure out which block is dominating all accesses of those variables.
for (auto &block : blocks)
@@ -3688,14 +3873,13 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry, AnalyzeVariableScopeA
{
// The variable is used in multiple continue blocks, this is not a loop
// candidate, signal that by setting block to -1u.
auto &potential = potential_loop_variables[var.first];
if (potential == 0)
potential = block;
if (potential_continue_block == 0)
potential_continue_block = block;
else
potential = ~(0u);
potential_continue_block = ~(0u);
}
}
builder.add_block(block);
}
@@ -3704,6 +3888,34 @@ void Compiler::analyze_variable_scope(SPIRFunction &entry, AnalyzeVariableScopeA
// Add it to a per-block list of variables.
BlockID dominating_block = builder.get_dominator();
if (dominating_block && potential_continue_block != 0 && potential_continue_block != ~0u)
{
auto &inner_block = get<SPIRBlock>(dominating_block);
BlockID merge_candidate = 0;
// Analyze the dominator. If it lives in a different loop scope than the candidate continue
// block, reject the loop variable candidate.
if (inner_block.merge == SPIRBlock::MergeLoop)
merge_candidate = inner_block.merge_block;
else if (inner_block.loop_dominator != SPIRBlock::NoDominator)
merge_candidate = get<SPIRBlock>(inner_block.loop_dominator).merge_block;
if (merge_candidate != 0 && cfg.is_reachable(merge_candidate))
{
// If the merge block has a higher post-visit order, we know that continue candidate
// cannot reach the merge block, and we have two separate scopes.
if (!cfg.is_reachable(potential_continue_block) ||
cfg.get_visit_order(merge_candidate) > cfg.get_visit_order(potential_continue_block))
{
potential_continue_block = 0;
}
}
}
if (potential_continue_block != 0 && potential_continue_block != ~0u)
potential_loop_variables[var.first] = potential_continue_block;
// For variables whose dominating block is inside a loop, there is a risk that these variables
// actually need to be preserved across loop iterations. We can express this by adding
// a "read" access to the loop header.
@@ -4303,8 +4515,9 @@ void Compiler::analyze_image_and_sampler_usage()
handler.dependency_hierarchy.clear();
traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), handler);
comparison_ids = move(handler.comparison_ids);
comparison_ids = std::move(handler.comparison_ids);
need_subpass_input = handler.need_subpass_input;
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)
@@ -4356,7 +4569,7 @@ void Compiler::build_function_control_flow_graphs_and_analyze()
CFGBuilder handler(*this);
handler.function_cfgs[ir.default_entry_point].reset(new CFG(*this, get<SPIRFunction>(ir.default_entry_point)));
traverse_all_reachable_opcodes(get<SPIRFunction>(ir.default_entry_point), handler);
function_cfgs = move(handler.function_cfgs);
function_cfgs = std::move(handler.function_cfgs);
bool single_function = function_cfgs.size() <= 1;
for (auto &f : function_cfgs)
@@ -4471,7 +4684,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<SPIRType>(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)
@@ -4646,46 +4863,22 @@ bool Compiler::reflection_ssbo_instance_name_is_significant() const
return aliased_ssbo_types;
}
bool Compiler::instruction_to_result_type(uint32_t &result_type, uint32_t &result_id, spv::Op op, const uint32_t *args,
uint32_t length)
bool Compiler::instruction_to_result_type(uint32_t &result_type, uint32_t &result_id, spv::Op op,
const uint32_t *args, uint32_t length)
{
// Most instructions follow the pattern of <result-type> <result-id> <arguments>.
// There are some exceptions.
switch (op)
{
case OpStore:
case OpCopyMemory:
case OpCopyMemorySized:
case OpImageWrite:
case OpAtomicStore:
case OpAtomicFlagClear:
case OpEmitStreamVertex:
case OpEndStreamPrimitive:
case OpControlBarrier:
case OpMemoryBarrier:
case OpGroupWaitEvents:
case OpRetainEvent:
case OpReleaseEvent:
case OpSetUserEventStatus:
case OpCaptureEventProfilingInfo:
case OpCommitReadPipe:
case OpCommitWritePipe:
case OpGroupCommitReadPipe:
case OpGroupCommitWritePipe:
case OpLine:
case OpNoLine:
if (length < 2)
return false;
default:
if (length > 1 && maybe_get<SPIRType>(args[0]) != nullptr)
{
result_type = args[0];
result_id = args[1];
return true;
}
else
return false;
bool has_result_id = false, has_result_type = false;
HasResultAndType(op, &has_result_id, &has_result_type);
if (has_result_id && has_result_type)
{
result_type = args[0];
result_id = args[1];
return true;
}
else
return false;
}
Bitset Compiler::combined_decoration_for_member(const SPIRType &type, uint32_t index) const
@@ -4769,6 +4962,12 @@ void Compiler::force_recompile()
is_force_recompile = true;
}
void Compiler::force_recompile_guarantee_forward_progress()
{
force_recompile();
is_force_recompile_forward_progress = true;
}
bool Compiler::is_forcing_recompilation() const
{
return is_force_recompile;
@@ -4777,6 +4976,7 @@ bool Compiler::is_forcing_recompilation() const
void Compiler::clear_force_recompile()
{
is_force_recompile = false;
is_force_recompile_forward_progress = false;
}
Compiler::PhysicalStorageBufferPointerHandler::PhysicalStorageBufferPointerHandler(Compiler &compiler_)
@@ -4958,7 +5158,7 @@ void Compiler::analyze_non_block_pointer_types()
for (auto type : handler.non_block_types)
physical_storage_non_block_pointer_types.push_back(type);
sort(begin(physical_storage_non_block_pointer_types), end(physical_storage_non_block_pointer_types));
physical_storage_type_to_alignment = move(handler.physical_block_type_meta);
physical_storage_type_to_alignment = std::move(handler.physical_block_type_meta);
}
bool Compiler::InterlockedResourceAccessPrepassHandler::handle(Op op, const uint32_t *, uint32_t)
@@ -5271,19 +5471,51 @@ void Compiler::analyze_interlocked_resource_usage()
}
}
bool Compiler::type_is_array_of_pointers(const SPIRType &type) const
// Helper function
bool Compiler::check_internal_recursion(const SPIRType &type, std::unordered_set<uint32_t> &checked_ids)
{
if (!type.pointer)
if (type.basetype != SPIRType::Struct)
return false;
// If parent type has same pointer depth, we must have an array of pointers.
return type.pointer_depth == get<SPIRType>(type.parent_type).pointer_depth;
if (checked_ids.count(type.self))
return true;
// Recurse into struct members
bool is_recursive = false;
checked_ids.insert(type.self);
uint32_t mbr_cnt = uint32_t(type.member_types.size());
for (uint32_t mbr_idx = 0; !is_recursive && mbr_idx < mbr_cnt; mbr_idx++)
{
uint32_t mbr_type_id = type.member_types[mbr_idx];
auto &mbr_type = get<SPIRType>(mbr_type_id);
is_recursive |= check_internal_recursion(mbr_type, checked_ids);
}
checked_ids.erase(type.self);
return is_recursive;
}
bool Compiler::type_is_top_level_physical_pointer(const SPIRType &type) const
// Return whether the struct type contains a structural recursion nested somewhere within its content.
bool Compiler::type_contains_recursion(const SPIRType &type)
{
return type.pointer && type.storage == StorageClassPhysicalStorageBuffer &&
type.pointer_depth > get<SPIRType>(type.parent_type).pointer_depth;
std::unordered_set<uint32_t> checked_ids;
return check_internal_recursion(type, checked_ids);
}
bool Compiler::type_is_array_of_pointers(const SPIRType &type) const
{
if (!is_array(type))
return false;
// BDA types must have parent type hierarchy.
if (!type.parent_type)
return false;
// Punch through all array layers.
auto *parent = &get<SPIRType>(type.parent_type);
while (is_array(*parent))
parent = &get<SPIRType>(parent->parent_type);
return is_pointer(*parent);
}
bool Compiler::flush_phi_required(BlockID from, BlockID to) const
src/libraries/spirv_cross/spirv_cross.hpp
+32 -3
View File
@@ -24,6 +24,9 @@
#ifndef SPIRV_CROSS_HPP
#define SPIRV_CROSS_HPP
#ifndef SPV_ENABLE_UTILITY_CODE
#define SPV_ENABLE_UTILITY_CODE
#endif
#include "spirv.hpp"
#include "spirv_cfg.hpp"
#include "spirv_cross_parsed_ir.hpp"
@@ -91,11 +94,14 @@ struct ShaderResources
SmallVector<Resource> sampled_images;
SmallVector<Resource> atomic_counters;
SmallVector<Resource> acceleration_structures;
SmallVector<Resource> gl_plain_uniforms;
// There can only be one push constant block,
// but keep the vector in case this restriction is lifted in the future.
SmallVector<Resource> push_constant_buffers;
SmallVector<Resource> shader_record_buffers;
// For Vulkan GLSL and HLSL source,
// these correspond to separate texture2D and samplers respectively.
SmallVector<Resource> separate_images;
@@ -357,12 +363,16 @@ public:
void set_execution_mode(spv::ExecutionMode mode, uint32_t arg0 = 0, uint32_t arg1 = 0, uint32_t arg2 = 0);
// Gets argument for an execution mode (LocalSize, Invocations, OutputVertices).
// For LocalSize, the index argument is used to select the dimension (X = 0, Y = 1, Z = 2).
// For LocalSize or LocalSizeId, the index argument is used to select the dimension (X = 0, Y = 1, Z = 2).
// For execution modes which do not have arguments, 0 is returned.
// LocalSizeId query returns an ID. If LocalSizeId execution mode is not used, it returns 0.
// LocalSize always returns a literal. If execution mode is LocalSizeId,
// the literal (spec constant or not) is still returned.
uint32_t get_execution_mode_argument(spv::ExecutionMode mode, uint32_t index = 0) const;
spv::ExecutionModel get_execution_model() const;
bool is_tessellation_shader() const;
bool is_tessellating_triangles() const;
// In SPIR-V, the compute work group size can be represented by a constant vector, in which case
// the LocalSize execution mode is ignored.
@@ -380,6 +390,8 @@ public:
// If the component is not a specialization constant, a zeroed out struct will be written.
// The return value is the constant ID of the builtin WorkGroupSize, but this is not expected to be useful
// for most use cases.
// If LocalSizeId is used, there is no uvec3 value representing the workgroup size, so the return value is 0,
// but x, y and z are written as normal if the components are specialization constants.
uint32_t get_work_group_size_specialization_constants(SpecializationConstant &x, SpecializationConstant &y,
SpecializationConstant &z) const;
@@ -551,6 +563,11 @@ protected:
}
}
uint32_t *stream_mutable(const Instruction &instr) const
{
return const_cast<uint32_t *>(stream(instr));
}
ParsedIR ir;
// Marks variables which have global scope and variables which can alias with other variables
// (SSBO, image load store, etc)
@@ -666,6 +683,9 @@ protected:
bool is_vector(const SPIRType &type) const;
bool is_matrix(const SPIRType &type) const;
bool is_array(const SPIRType &type) const;
bool is_pointer(const SPIRType &type) const;
bool is_physical_pointer(const SPIRType &type) const;
static bool is_runtime_size_array(const SPIRType &type);
uint32_t expression_type_id(uint32_t id) const;
const SPIRType &expression_type(uint32_t id) const;
bool expression_is_lvalue(uint32_t id) const;
@@ -730,16 +750,20 @@ protected:
SPIRBlock::ContinueBlockType continue_block_type(const SPIRBlock &continue_block) const;
void force_recompile();
void force_recompile_guarantee_forward_progress();
void clear_force_recompile();
bool is_forcing_recompilation() const;
bool is_force_recompile = false;
bool is_force_recompile_forward_progress = false;
bool block_is_noop(const SPIRBlock &block) const;
bool block_is_loop_candidate(const SPIRBlock &block, SPIRBlock::Method method) const;
bool types_are_logically_equivalent(const SPIRType &a, const SPIRType &b) const;
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.
@@ -915,6 +939,7 @@ protected:
// Similar is implemented for images, as well as if subpass inputs are needed.
std::unordered_set<uint32_t> 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,
@@ -954,6 +979,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);
};
@@ -994,7 +1020,8 @@ protected:
std::unordered_map<uint32_t, std::unordered_set<uint32_t>> partial_write_variables_to_block;
std::unordered_set<uint32_t> access_chain_expressions;
// Access chains used in multiple blocks mean hoisting all the variables used to construct the access chain as not all backends can use pointers.
std::unordered_map<uint32_t, std::unordered_set<uint32_t>> access_chain_children;
// This is also relevant when forwarding opaque objects since we cannot lower these to temporaries.
std::unordered_map<uint32_t, std::unordered_set<uint32_t>> rvalue_forward_children;
const SPIRBlock *current_block = nullptr;
};
@@ -1120,9 +1147,11 @@ protected:
bool has_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration) const;
void unset_extended_member_decoration(uint32_t type, uint32_t index, ExtendedDecorations decoration);
bool check_internal_recursion(const SPIRType &type, std::unordered_set<uint32_t> &checked_ids);
bool type_contains_recursion(const SPIRType &type);
bool type_is_array_of_pointers(const SPIRType &type) const;
bool type_is_top_level_physical_pointer(const SPIRType &type) const;
bool type_is_block_like(const SPIRType &type) const;
bool type_is_top_level_block(const SPIRType &type) const;
bool type_is_opaque_value(const SPIRType &type) const;
bool reflection_ssbo_instance_name_is_significant() const;
src/libraries/spirv_cross/spirv_cross_c.cpp
+262 -20
View File
@@ -194,6 +194,7 @@ struct spvc_resources_s : ScratchMemoryAllocation
SmallVector<spvc_reflected_resource> sampled_images;
SmallVector<spvc_reflected_resource> atomic_counters;
SmallVector<spvc_reflected_resource> push_constant_buffers;
SmallVector<spvc_reflected_resource> shader_record_buffers;
SmallVector<spvc_reflected_resource> separate_images;
SmallVector<spvc_reflected_resource> separate_samplers;
SmallVector<spvc_reflected_resource> acceleration_structures;
@@ -251,7 +252,7 @@ spvc_result spvc_context_parse_spirv(spvc_context context, const SpvId *spirv, s
pir->context = context;
Parser parser(spirv, word_count);
parser.parse();
pir->parsed = move(parser.get_parsed_ir());
pir->parsed = std::move(parser.get_parsed_ir());
*parsed_ir = pir.get();
context->allocations.push_back(std::move(pir));
}
@@ -283,7 +284,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back
{
case SPVC_BACKEND_NONE:
if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
comp->compiler.reset(new Compiler(move(parsed_ir->parsed)));
comp->compiler.reset(new Compiler(std::move(parsed_ir->parsed)));
else if (mode == SPVC_CAPTURE_MODE_COPY)
comp->compiler.reset(new Compiler(parsed_ir->parsed));
break;
@@ -291,7 +292,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back
#if SPIRV_CROSS_C_API_GLSL
case SPVC_BACKEND_GLSL:
if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
comp->compiler.reset(new CompilerGLSL(move(parsed_ir->parsed)));
comp->compiler.reset(new CompilerGLSL(std::move(parsed_ir->parsed)));
else if (mode == SPVC_CAPTURE_MODE_COPY)
comp->compiler.reset(new CompilerGLSL(parsed_ir->parsed));
break;
@@ -300,7 +301,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back
#if SPIRV_CROSS_C_API_HLSL
case SPVC_BACKEND_HLSL:
if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
comp->compiler.reset(new CompilerHLSL(move(parsed_ir->parsed)));
comp->compiler.reset(new CompilerHLSL(std::move(parsed_ir->parsed)));
else if (mode == SPVC_CAPTURE_MODE_COPY)
comp->compiler.reset(new CompilerHLSL(parsed_ir->parsed));
break;
@@ -309,7 +310,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back
#if SPIRV_CROSS_C_API_MSL
case SPVC_BACKEND_MSL:
if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
comp->compiler.reset(new CompilerMSL(move(parsed_ir->parsed)));
comp->compiler.reset(new CompilerMSL(std::move(parsed_ir->parsed)));
else if (mode == SPVC_CAPTURE_MODE_COPY)
comp->compiler.reset(new CompilerMSL(parsed_ir->parsed));
break;
@@ -318,7 +319,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back
#if SPIRV_CROSS_C_API_CPP
case SPVC_BACKEND_CPP:
if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
comp->compiler.reset(new CompilerCPP(move(parsed_ir->parsed)));
comp->compiler.reset(new CompilerCPP(std::move(parsed_ir->parsed)));
else if (mode == SPVC_CAPTURE_MODE_COPY)
comp->compiler.reset(new CompilerCPP(parsed_ir->parsed));
break;
@@ -327,7 +328,7 @@ spvc_result spvc_context_create_compiler(spvc_context context, spvc_backend back
#if SPIRV_CROSS_C_API_REFLECT
case SPVC_BACKEND_JSON:
if (mode == SPVC_CAPTURE_MODE_TAKE_OWNERSHIP)
comp->compiler.reset(new CompilerReflection(move(parsed_ir->parsed)));
comp->compiler.reset(new CompilerReflection(std::move(parsed_ir->parsed)));
else if (mode == SPVC_CAPTURE_MODE_COPY)
comp->compiler.reset(new CompilerReflection(parsed_ir->parsed));
break;
@@ -475,6 +476,12 @@ spvc_result spvc_compiler_options_set_uint(spvc_compiler_options options, spvc_c
case SPVC_COMPILER_OPTION_GLSL_OVR_MULTIVIEW_VIEW_COUNT:
options->glsl.ovr_multiview_view_count = value;
break;
case SPVC_COMPILER_OPTION_RELAX_NAN_CHECKS:
options->glsl.relax_nan_checks = value != 0;
break;
case SPVC_COMPILER_OPTION_GLSL_ENABLE_ROW_MAJOR_LOAD_WORKAROUND:
options->glsl.enable_row_major_load_workaround = value != 0;
break;
#endif
#if SPIRV_CROSS_C_API_HLSL
@@ -711,6 +718,42 @@ spvc_result spvc_compiler_options_set_uint(spvc_compiler_options options, spvc_c
case SPVC_COMPILER_OPTION_MSL_IOS_SUPPORT_BASE_VERTEX_INSTANCE:
options->msl.ios_support_base_vertex_instance = value != 0;
break;
case SPVC_COMPILER_OPTION_MSL_RAW_BUFFER_TESE_INPUT:
options->msl.raw_buffer_tese_input = value != 0;
break;
case SPVC_COMPILER_OPTION_MSL_SHADER_PATCH_INPUT_BUFFER_INDEX:
options->msl.shader_patch_input_buffer_index = value;
break;
case SPVC_COMPILER_OPTION_MSL_MANUAL_HELPER_INVOCATION_UPDATES:
options->msl.manual_helper_invocation_updates = value != 0;
break;
case SPVC_COMPILER_OPTION_MSL_CHECK_DISCARDED_FRAG_STORES:
options->msl.check_discarded_frag_stores = value != 0;
break;
case SPVC_COMPILER_OPTION_MSL_ARGUMENT_BUFFERS_TIER:
options->msl.argument_buffers_tier = static_cast<CompilerMSL::Options::ArgumentBuffersTier>(value);
break;
case SPVC_COMPILER_OPTION_MSL_SAMPLE_DREF_LOD_ARRAY_AS_GRAD:
options->msl.sample_dref_lod_array_as_grad = value != 0;
break;
case SPVC_COMPILER_OPTION_MSL_READWRITE_TEXTURE_FENCES:
options->msl.readwrite_texture_fences = value != 0;
break;
case SPVC_COMPILER_OPTION_MSL_REPLACE_RECURSIVE_INPUTS:
options->msl.replace_recursive_inputs = value != 0;
break;
case SPVC_COMPILER_OPTION_MSL_AGX_MANUAL_CUBE_GRAD_FIXUP:
options->msl.agx_manual_cube_grad_fixup = value != 0;
break;
#endif
default:
@@ -789,6 +832,43 @@ spvc_result spvc_compiler_require_extension(spvc_compiler compiler, const char *
#endif
}
size_t spvc_compiler_get_num_required_extensions(spvc_compiler compiler)
{
#if SPIRV_CROSS_C_API_GLSL
if (compiler->backend != SPVC_BACKEND_GLSL)
{
compiler->context->report_error("Enabled extensions can only be queried on GLSL backend.");
return SPVC_ERROR_INVALID_ARGUMENT;
}
return static_cast<CompilerGLSL *>(compiler->compiler.get())->get_required_extensions().size();
#else
compiler->context->report_error("Enabled extensions can only be queried on GLSL backend.");
return 0;
#endif
}
const char *spvc_compiler_get_required_extension(spvc_compiler compiler, size_t index)
{
#if SPIRV_CROSS_C_API_GLSL
if (compiler->backend != SPVC_BACKEND_GLSL)
{
compiler->context->report_error("Enabled extensions can only be queried on GLSL backend.");
return nullptr;
}
auto &exts = static_cast<CompilerGLSL *>(compiler->compiler.get())->get_required_extensions();
if (index < exts.size())
return exts[index].c_str();
else
return nullptr;
#else
(void)index;
compiler->context->report_error("Enabled extensions can only be queried on GLSL backend.");
return nullptr;
#endif
}
spvc_result spvc_compiler_flatten_buffer_block(spvc_compiler compiler, spvc_variable_id id)
{
#if SPIRV_CROSS_C_API_GLSL
@@ -1133,9 +1213,9 @@ spvc_result spvc_compiler_msl_add_vertex_attribute(spvc_compiler compiler, const
}
auto &msl = *static_cast<CompilerMSL *>(compiler->compiler.get());
MSLShaderInput attr;
MSLShaderInterfaceVariable attr;
attr.location = va->location;
attr.format = static_cast<MSLShaderInputFormat>(va->format);
attr.format = static_cast<MSLShaderVariableFormat>(va->format);
attr.builtin = static_cast<spv::BuiltIn>(va->builtin);
msl.add_msl_shader_input(attr);
return SPVC_SUCCESS;
@@ -1146,7 +1226,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)
@@ -1156,9 +1236,9 @@ spvc_result spvc_compiler_msl_add_shader_input(spvc_compiler compiler, const spv
}
auto &msl = *static_cast<CompilerMSL *>(compiler->compiler.get());
MSLShaderInput input;
MSLShaderInterfaceVariable input;
input.location = si->location;
input.format = static_cast<MSLShaderInputFormat>(si->format);
input.format = static_cast<MSLShaderVariableFormat>(si->format);
input.builtin = static_cast<spv::BuiltIn>(si->builtin);
input.vecsize = si->vecsize;
msl.add_msl_shader_input(input);
@@ -1170,6 +1250,80 @@ spvc_result spvc_compiler_msl_add_shader_input(spvc_compiler compiler, const spv
#endif
}
spvc_result spvc_compiler_msl_add_shader_input_2(spvc_compiler compiler, const spvc_msl_shader_interface_var_2 *si)
{
#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<CompilerMSL *>(compiler->compiler.get());
MSLShaderInterfaceVariable input;
input.location = si->location;
input.format = static_cast<MSLShaderVariableFormat>(si->format);
input.builtin = static_cast<spv::BuiltIn>(si->builtin);
input.vecsize = si->vecsize;
input.rate = static_cast<MSLShaderVariableRate>(si->rate);
msl.add_msl_shader_input(input);
return SPVC_SUCCESS;
#else
(void)si;
compiler->context->report_error("MSL function used on a non-MSL backend.");
return SPVC_ERROR_INVALID_ARGUMENT;
#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<CompilerMSL *>(compiler->compiler.get());
MSLShaderInterfaceVariable output;
output.location = so->location;
output.format = static_cast<MSLShaderVariableFormat>(so->format);
output.builtin = static_cast<spv::BuiltIn>(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_shader_output_2(spvc_compiler compiler, const spvc_msl_shader_interface_var_2 *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<CompilerMSL *>(compiler->compiler.get());
MSLShaderInterfaceVariable output;
output.location = so->location;
output.format = static_cast<MSLShaderVariableFormat>(so->format);
output.builtin = static_cast<spv::BuiltIn>(so->builtin);
output.vecsize = so->vecsize;
output.rate = static_cast<MSLShaderVariableRate>(so->rate);
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)
{
@@ -1295,6 +1449,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<CompilerMSL *>(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);
@@ -1639,6 +1811,8 @@ bool spvc_resources_s::copy_resources(const ShaderResources &resources)
return false;
if (!copy_resources(push_constant_buffers, resources.push_constant_buffers))
return false;
if (!copy_resources(shader_record_buffers, resources.shader_record_buffers))
return false;
if (!copy_resources(separate_images, resources.separate_images))
return false;
if (!copy_resources(separate_samplers, resources.separate_samplers))
@@ -1792,6 +1966,10 @@ spvc_result spvc_resources_get_resource_list_for_type(spvc_resources resources,
list = &resources->acceleration_structures;
break;
case SPVC_RESOURCE_TYPE_SHADER_RECORD_BUFFER:
list = &resources->shader_record_buffers;
break;
default:
break;
}
@@ -2406,6 +2584,51 @@ spvc_type_id spvc_constant_get_type(spvc_constant constant)
return constant->constant_type;
}
void spvc_constant_set_scalar_fp16(spvc_constant constant, unsigned column, unsigned row, unsigned short value)
{
constant->m.c[column].r[row].u32 = value;
}
void spvc_constant_set_scalar_fp32(spvc_constant constant, unsigned column, unsigned row, float value)
{
constant->m.c[column].r[row].f32 = value;
}
void spvc_constant_set_scalar_fp64(spvc_constant constant, unsigned column, unsigned row, double value)
{
constant->m.c[column].r[row].f64 = value;
}
void spvc_constant_set_scalar_u32(spvc_constant constant, unsigned column, unsigned row, unsigned value)
{
constant->m.c[column].r[row].u32 = value;
}
void spvc_constant_set_scalar_i32(spvc_constant constant, unsigned column, unsigned row, int value)
{
constant->m.c[column].r[row].i32 = value;
}
void spvc_constant_set_scalar_u16(spvc_constant constant, unsigned column, unsigned row, unsigned short value)
{
constant->m.c[column].r[row].u32 = uint32_t(value);
}
void spvc_constant_set_scalar_i16(spvc_constant constant, unsigned column, unsigned row, signed short value)
{
constant->m.c[column].r[row].u32 = uint32_t(value);
}
void spvc_constant_set_scalar_u8(spvc_constant constant, unsigned column, unsigned row, unsigned char value)
{
constant->m.c[column].r[row].u32 = uint32_t(value);
}
void spvc_constant_set_scalar_i8(spvc_constant constant, unsigned column, unsigned row, signed char value)
{
constant->m.c[column].r[row].u32 = uint32_t(value);
}
spvc_bool spvc_compiler_get_binary_offset_for_decoration(spvc_compiler compiler, spvc_variable_id id,
SpvDecoration decoration,
unsigned *word_offset)
@@ -2508,7 +2731,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<spvc_msl_vertex_format>(attr_default.format);
attr->builtin = static_cast<SpvBuiltIn>(attr_default.builtin);
@@ -2517,16 +2740,35 @@ 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<spvc_msl_shader_input_format>(input_default.format);
input->builtin = static_cast<SpvBuiltIn>(input_default.builtin);
input->vecsize = input_default.vecsize;
MSLShaderInterfaceVariable var_default;
var->location = var_default.location;
var->format = static_cast<spvc_msl_shader_variable_format>(var_default.format);
var->builtin = static_cast<SpvBuiltIn>(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_shader_interface_var_init_2(spvc_msl_shader_interface_var_2 *var)
{
#if SPIRV_CROSS_C_API_MSL
MSLShaderInterfaceVariable var_default;
var->location = var_default.location;
var->format = static_cast<spvc_msl_shader_variable_format>(var_default.format);
var->builtin = static_cast<SpvBuiltIn>(var_default.builtin);
var->vecsize = var_default.vecsize;
var->rate = static_cast<spvc_msl_shader_variable_rate>(var_default.rate);
#else
memset(var, 0, sizeof(*var));
#endif
}
src/libraries/spirv_cross/spirv_cross_c.h
+92 -16
View File
@@ -40,7 +40,7 @@ extern "C" {
/* Bumped if ABI or API breaks backwards compatibility. */
#define SPVC_C_API_VERSION_MAJOR 0
/* Bumped if APIs or enumerations are added in a backwards compatible way. */
#define SPVC_C_API_VERSION_MINOR 48
#define SPVC_C_API_VERSION_MINOR 58
/* Bumped if internal implementation details change. */
#define SPVC_C_API_VERSION_PATCH 0
@@ -225,6 +225,7 @@ typedef enum spvc_resource_type
SPVC_RESOURCE_TYPE_SEPARATE_SAMPLERS = 11,
SPVC_RESOURCE_TYPE_ACCELERATION_STRUCTURE = 12,
SPVC_RESOURCE_TYPE_RAY_QUERY = 13,
SPVC_RESOURCE_TYPE_SHADER_RECORD_BUFFER = 14,
SPVC_RESOURCE_TYPE_INT_MAX = 0x7fffffff
} spvc_resource_type;
@@ -290,23 +291,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;
@@ -329,20 +336,50 @@ 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
/* Maps to C++ API. Deprecated; use spvc_msl_shader_interface_var_2. */
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.
* Deprecated. Use spvc_msl_shader_interface_var_init_2().
*/
SPVC_PUBLIC_API void spvc_msl_shader_interface_var_init(spvc_msl_shader_interface_var *var);
/*
* Deprecated. Use spvc_msl_shader_interface_var_init_2().
*/
SPVC_PUBLIC_API void spvc_msl_shader_input_init(spvc_msl_shader_input *input);
/* Maps to C++ API. */
typedef enum spvc_msl_shader_variable_rate
{
SPVC_MSL_SHADER_VARIABLE_RATE_PER_VERTEX = 0,
SPVC_MSL_SHADER_VARIABLE_RATE_PER_PRIMITIVE = 1,
SPVC_MSL_SHADER_VARIABLE_RATE_PER_PATCH = 2,
SPVC_MSL_SHADER_VARIABLE_RATE_INT_MAX = 0x7fffffff,
} spvc_msl_shader_variable_rate;
/* Maps to C++ API. */
typedef struct spvc_msl_shader_interface_var_2
{
unsigned location;
spvc_msl_shader_variable_format format;
SpvBuiltIn builtin;
unsigned vecsize;
spvc_msl_shader_variable_rate rate;
} spvc_msl_shader_interface_var_2;
/*
* Initializes the shader interface variable struct.
*/
SPVC_PUBLIC_API void spvc_msl_shader_interface_var_init_2(spvc_msl_shader_interface_var_2 *var);
/* Maps to C++ API. */
typedef struct spvc_msl_resource_binding
{
@@ -677,6 +714,21 @@ typedef enum spvc_compiler_option
SPVC_COMPILER_OPTION_GLSL_OVR_MULTIVIEW_VIEW_COUNT = 77 | SPVC_COMPILER_OPTION_GLSL_BIT,
SPVC_COMPILER_OPTION_RELAX_NAN_CHECKS = 78 | SPVC_COMPILER_OPTION_COMMON_BIT,
SPVC_COMPILER_OPTION_MSL_RAW_BUFFER_TESE_INPUT = 79 | SPVC_COMPILER_OPTION_MSL_BIT,
SPVC_COMPILER_OPTION_MSL_SHADER_PATCH_INPUT_BUFFER_INDEX = 80 | SPVC_COMPILER_OPTION_MSL_BIT,
SPVC_COMPILER_OPTION_MSL_MANUAL_HELPER_INVOCATION_UPDATES = 81 | SPVC_COMPILER_OPTION_MSL_BIT,
SPVC_COMPILER_OPTION_MSL_CHECK_DISCARDED_FRAG_STORES = 82 | SPVC_COMPILER_OPTION_MSL_BIT,
SPVC_COMPILER_OPTION_GLSL_ENABLE_ROW_MAJOR_LOAD_WORKAROUND = 83 | SPVC_COMPILER_OPTION_GLSL_BIT,
SPVC_COMPILER_OPTION_MSL_ARGUMENT_BUFFERS_TIER = 84 | SPVC_COMPILER_OPTION_MSL_BIT,
SPVC_COMPILER_OPTION_MSL_SAMPLE_DREF_LOD_ARRAY_AS_GRAD = 85 | SPVC_COMPILER_OPTION_MSL_BIT,
SPVC_COMPILER_OPTION_MSL_READWRITE_TEXTURE_FENCES = 86 | SPVC_COMPILER_OPTION_MSL_BIT,
SPVC_COMPILER_OPTION_MSL_REPLACE_RECURSIVE_INPUTS = 87 | SPVC_COMPILER_OPTION_MSL_BIT,
SPVC_COMPILER_OPTION_MSL_AGX_MANUAL_CUBE_GRAD_FIXUP = 88 | SPVC_COMPILER_OPTION_MSL_BIT,
SPVC_COMPILER_OPTION_INT_MAX = 0x7fffffff
} spvc_compiler_option;
@@ -735,6 +787,8 @@ SPVC_PUBLIC_API spvc_result spvc_compiler_compile(spvc_compiler compiler, const
/* Maps to C++ API. */
SPVC_PUBLIC_API spvc_result spvc_compiler_add_header_line(spvc_compiler compiler, const char *line);
SPVC_PUBLIC_API spvc_result spvc_compiler_require_extension(spvc_compiler compiler, const char *ext);
SPVC_PUBLIC_API size_t spvc_compiler_get_num_required_extensions(spvc_compiler compiler);
SPVC_PUBLIC_API const char *spvc_compiler_get_required_extension(spvc_compiler compiler, size_t index);
SPVC_PUBLIC_API spvc_result spvc_compiler_flatten_buffer_block(spvc_compiler compiler, spvc_variable_id id);
SPVC_PUBLIC_API spvc_bool spvc_compiler_variable_is_depth_or_compare(spvc_compiler compiler, spvc_variable_id id);
@@ -783,14 +837,23 @@ SPVC_PUBLIC_API spvc_result spvc_compiler_msl_add_vertex_attribute(spvc_compiler
const spvc_msl_vertex_attribute *attrs);
SPVC_PUBLIC_API spvc_result spvc_compiler_msl_add_resource_binding(spvc_compiler compiler,
const spvc_msl_resource_binding *binding);
/* Deprecated; use spvc_compiler_msl_add_shader_input_2(). */
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_input_2(spvc_compiler compiler,
const spvc_msl_shader_interface_var_2 *input);
/* Deprecated; use spvc_compiler_msl_add_shader_output_2(). */
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_shader_output_2(spvc_compiler compiler,
const spvc_msl_shader_interface_var_2 *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,
@@ -991,6 +1054,19 @@ SPVC_PUBLIC_API int spvc_constant_get_scalar_i8(spvc_constant constant, unsigned
SPVC_PUBLIC_API void spvc_constant_get_subconstants(spvc_constant constant, const spvc_constant_id **constituents, size_t *count);
SPVC_PUBLIC_API spvc_type_id spvc_constant_get_type(spvc_constant constant);
/*
* C implementation of the C++ api.
*/
SPVC_PUBLIC_API void spvc_constant_set_scalar_fp16(spvc_constant constant, unsigned column, unsigned row, unsigned short value);
SPVC_PUBLIC_API void spvc_constant_set_scalar_fp32(spvc_constant constant, unsigned column, unsigned row, float value);
SPVC_PUBLIC_API void spvc_constant_set_scalar_fp64(spvc_constant constant, unsigned column, unsigned row, double value);
SPVC_PUBLIC_API void spvc_constant_set_scalar_u32(spvc_constant constant, unsigned column, unsigned row, unsigned value);
SPVC_PUBLIC_API void spvc_constant_set_scalar_i32(spvc_constant constant, unsigned column, unsigned row, int value);
SPVC_PUBLIC_API void spvc_constant_set_scalar_u16(spvc_constant constant, unsigned column, unsigned row, unsigned short value);
SPVC_PUBLIC_API void spvc_constant_set_scalar_i16(spvc_constant constant, unsigned column, unsigned row, signed short value);
SPVC_PUBLIC_API void spvc_constant_set_scalar_u8(spvc_constant constant, unsigned column, unsigned row, unsigned char value);
SPVC_PUBLIC_API void spvc_constant_set_scalar_i8(spvc_constant constant, unsigned column, unsigned row, signed char value);
/*
* Misc reflection
* Maps to C++ API.
src/libraries/spirv_cross/spirv_cross_containers.hpp
+10 -2
View File
@@ -26,6 +26,7 @@
#include "spirv_cross_error_handling.hpp"
#include <algorithm>
#include <exception>
#include <functional>
#include <iterator>
#include <limits>
@@ -210,7 +211,8 @@ public:
buffer_capacity = N;
}
SmallVector(const T *arg_list_begin, const T *arg_list_end) SPIRV_CROSS_NOEXCEPT : SmallVector()
template <typename U>
SmallVector(const U *arg_list_begin, const U *arg_list_end) SPIRV_CROSS_NOEXCEPT : SmallVector()
{
auto count = size_t(arg_list_end - arg_list_begin);
reserve(count);
@@ -219,7 +221,13 @@ public:
this->buffer_size = count;
}
SmallVector(std::initializer_list<T> init) SPIRV_CROSS_NOEXCEPT : SmallVector(init.begin(), init.end())
template <typename U>
SmallVector(std::initializer_list<U> init) SPIRV_CROSS_NOEXCEPT : SmallVector(init.begin(), init.end())
{
}
template <typename U, size_t M>
explicit SmallVector(const U (&init)[M]) SPIRV_CROSS_NOEXCEPT : SmallVector(init, init + M)
{
}
src/libraries/spirv_cross/spirv_cross_parsed_ir.cpp
+35 -21
View File
@@ -54,26 +54,26 @@ ParsedIR::ParsedIR()
// Should have been default-implemented, but need this on MSVC 2013.
ParsedIR::ParsedIR(ParsedIR &&other) SPIRV_CROSS_NOEXCEPT
{
*this = move(other);
*this = std::move(other);
}
ParsedIR &ParsedIR::operator=(ParsedIR &&other) SPIRV_CROSS_NOEXCEPT
{
if (this != &other)
{
pool_group = move(other.pool_group);
spirv = move(other.spirv);
meta = move(other.meta);
pool_group = std::move(other.pool_group);
spirv = std::move(other.spirv);
meta = std::move(other.meta);
for (int i = 0; i < TypeCount; i++)
ids_for_type[i] = move(other.ids_for_type[i]);
ids_for_constant_or_type = move(other.ids_for_constant_or_type);
ids_for_constant_or_variable = move(other.ids_for_constant_or_variable);
declared_capabilities = move(other.declared_capabilities);
declared_extensions = move(other.declared_extensions);
block_meta = move(other.block_meta);
continue_block_to_loop_header = move(other.continue_block_to_loop_header);
entry_points = move(other.entry_points);
ids = move(other.ids);
ids_for_type[i] = std::move(other.ids_for_type[i]);
ids_for_constant_undef_or_type = std::move(other.ids_for_constant_undef_or_type);
ids_for_constant_or_variable = std::move(other.ids_for_constant_or_variable);
declared_capabilities = std::move(other.declared_capabilities);
declared_extensions = std::move(other.declared_extensions);
block_meta = std::move(other.block_meta);
continue_block_to_loop_header = std::move(other.continue_block_to_loop_header);
entry_points = std::move(other.entry_points);
ids = std::move(other.ids);
addressing_model = other.addressing_model;
memory_model = other.memory_model;
@@ -102,7 +102,7 @@ ParsedIR &ParsedIR::operator=(const ParsedIR &other)
meta = other.meta;
for (int i = 0; i < TypeCount; i++)
ids_for_type[i] = other.ids_for_type[i];
ids_for_constant_or_type = other.ids_for_constant_or_type;
ids_for_constant_undef_or_type = other.ids_for_constant_undef_or_type;
ids_for_constant_or_variable = other.ids_for_constant_or_variable;
declared_capabilities = other.declared_capabilities;
declared_extensions = other.declared_extensions;
@@ -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<SPIRVariable>(id).remapped_variable)
continue;
auto &m = meta[id];
sanitize_identifier(m.decoration.alias, false, false);
for (auto &memb : m.members)
@@ -349,7 +353,7 @@ void ParsedIR::set_name(ID id, const string &name)
void ParsedIR::set_member_name(TypeID id, uint32_t index, const string &name)
{
auto &m = meta[id];
m.members.resize(max(meta[id].members.size(), size_t(index) + 1));
m.members.resize(max(m.members.size(), size_t(index) + 1));
m.members[index].alias = name;
if (!is_valid_identifier(name) || is_reserved_identifier(name, true, false))
meta_needing_name_fixup.insert(id);
@@ -366,6 +370,10 @@ void ParsedIR::set_decoration_string(ID id, Decoration decoration, const string
dec.hlsl_semantic = argument;
break;
case DecorationUserTypeGOOGLE:
dec.user_type = argument;
break;
default:
break;
}
@@ -451,8 +459,9 @@ void ParsedIR::set_decoration(ID id, Decoration decoration, uint32_t argument)
void ParsedIR::set_member_decoration(TypeID id, uint32_t index, Decoration decoration, uint32_t argument)
{
meta[id].members.resize(max(meta[id].members.size(), size_t(index) + 1));
auto &dec = meta[id].members[index];
auto &m = meta[id];
m.members.resize(max(m.members.size(), size_t(index) + 1));
auto &dec = m.members[index];
dec.decoration_flags.set(decoration);
switch (decoration)
@@ -654,6 +663,9 @@ const string &ParsedIR::get_decoration_string(ID id, Decoration decoration) cons
case DecorationHlslSemanticGOOGLE:
return dec.hlsl_semantic;
case DecorationUserTypeGOOGLE:
return dec.user_type;
default:
return empty_string;
}
@@ -792,7 +804,8 @@ const Bitset &ParsedIR::get_decoration_bitset(ID id) const
void ParsedIR::set_member_decoration_string(TypeID id, uint32_t index, Decoration decoration, const string &argument)
{
meta[id].members.resize(max(meta[id].members.size(), size_t(index) + 1));
auto &m = meta[id];
m.members.resize(max(m.members.size(), size_t(index) + 1));
auto &dec = meta[id].members[index];
dec.decoration_flags.set(decoration);
@@ -928,7 +941,7 @@ void ParsedIR::add_typed_id(Types type, ID id)
{
case TypeConstant:
ids_for_constant_or_variable.push_back(id);
ids_for_constant_or_type.push_back(id);
ids_for_constant_undef_or_type.push_back(id);
break;
case TypeVariable:
@@ -937,7 +950,8 @@ void ParsedIR::add_typed_id(Types type, ID id)
case TypeType:
case TypeConstantOp:
ids_for_constant_or_type.push_back(id);
case TypeUndef:
ids_for_constant_undef_or_type.push_back(id);
break;
default:
@@ -999,7 +1013,7 @@ ParsedIR::LoopLock::LoopLock(uint32_t *lock_)
ParsedIR::LoopLock::LoopLock(LoopLock &&other) SPIRV_CROSS_NOEXCEPT
{
*this = move(other);
*this = std::move(other);
}
ParsedIR::LoopLock &ParsedIR::LoopLock::operator=(LoopLock &&other) SPIRV_CROSS_NOEXCEPT
src/libraries/spirv_cross/spirv_cross_parsed_ir.hpp
+2 -2
View File
@@ -74,8 +74,8 @@ public:
// Special purpose lists which contain a union of types.
// This is needed so we can declare specialization constants and structs in an interleaved fashion,
// among other things.
// Constants can be of struct type, and struct array sizes can use specialization constants.
SmallVector<ID> ids_for_constant_or_type;
// Constants can be undef or of struct type, and struct array sizes can use specialization constants.
SmallVector<ID> ids_for_constant_undef_or_type;
SmallVector<ID> ids_for_constant_or_variable;
// We need to keep track of the width the Ops that contains a type for the
src/libraries/spirv_cross/spirv_glsl.cpp
+3004 -511
View File
File diff suppressed because it is too large Load Diff
src/libraries/spirv_cross/spirv_glsl.hpp
+132 -21
View File
@@ -83,6 +83,11 @@ public:
// Debug option to always emit temporary variables for all expressions.
bool force_temporary = false;
// Debug option, can be increased in an attempt to workaround SPIRV-Cross bugs temporarily.
// If this limit has to be increased, it points to an implementation bug.
// In certain scenarios, the maximum number of debug iterations may increase beyond this limit
// as long as we can prove we're making certain kinds of forward progress.
uint32_t force_recompile_max_debug_iterations = 3;
// If true, Vulkan GLSL features are used instead of GL-compatible features.
// Mostly useful for debugging SPIR-V files.
@@ -133,6 +138,19 @@ public:
// what happens on legacy GLSL targets for blocks and structs.
bool force_flattened_io_blocks = false;
// For opcodes where we have to perform explicit additional nan checks, very ugly code is generated.
// If we opt-in, ignore these requirements.
// In opcodes like NClamp/NMin/NMax and FP compare, ignore NaN behavior.
// Use FClamp/FMin/FMax semantics for clamps and lets implementation choose ordered or unordered
// compares.
bool relax_nan_checks = false;
// 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;
@@ -240,6 +258,10 @@ public:
// require_extension("GL_KHR_my_extension");
void require_extension(const std::string &ext);
// Returns the list of required extensions. After compilation this will contains any other
// extensions that the compiler used automatically, in addition to the user specified ones.
const SmallVector<std::string> &get_required_extensions() const;
// Legacy GLSL compatibility method.
// Takes a uniform or push constant variable and flattens it into a (i|u)vec4 array[N]; array instead.
// For this to work, all types in the block must be the same basic type, e.g. mixing vec2 and vec4 is fine, but
@@ -265,6 +287,14 @@ public:
void mask_stage_output_by_location(uint32_t location, uint32_t component);
void mask_stage_output_by_builtin(spv::BuiltIn builtin);
// Allow to control how to format float literals in the output.
// Set to "nullptr" to use the default "convert_to_string" function.
// This handle is not owned by SPIRV-Cross and must remain valid until compile() has been called.
void set_float_formatter(FloatFormatter *formatter)
{
float_formatter = formatter;
}
protected:
struct ShaderSubgroupSupportHelper
{
@@ -274,6 +304,7 @@ protected:
KHR_shader_subgroup_ballot,
KHR_shader_subgroup_basic,
KHR_shader_subgroup_vote,
KHR_shader_subgroup_arithmetic,
NV_gpu_shader_5,
NV_shader_thread_group,
NV_shader_thread_shuffle,
@@ -306,7 +337,18 @@ protected:
SubgroupInverseBallot_InclBitCount_ExclBitCout = 13,
SubgroupBallotBitExtract = 14,
SubgroupBallotBitCount = 15,
SubgroupArithmeticIAddReduce = 16,
SubgroupArithmeticIAddExclusiveScan = 17,
SubgroupArithmeticIAddInclusiveScan = 18,
SubgroupArithmeticFAddReduce = 19,
SubgroupArithmeticFAddExclusiveScan = 20,
SubgroupArithmeticFAddInclusiveScan = 21,
SubgroupArithmeticIMulReduce = 22,
SubgroupArithmeticIMulExclusiveScan = 23,
SubgroupArithmeticIMulInclusiveScan = 24,
SubgroupArithmeticFMulReduce = 25,
SubgroupArithmeticFMulExclusiveScan = 26,
SubgroupArithmeticFMulInclusiveScan = 27,
FeatureCount
};
@@ -340,9 +382,9 @@ protected:
};
// TODO remove this function when all subgroup ops are supported (or make it always return true)
static bool is_supported_subgroup_op_in_opengl(spv::Op op);
static bool is_supported_subgroup_op_in_opengl(spv::Op op, const uint32_t *ops);
void reset();
void reset(uint32_t iteration_count);
void emit_function(SPIRFunction &func, const Bitset &return_flags);
bool has_extension(const std::string &ext) const;
@@ -352,11 +394,23 @@ protected:
virtual void emit_function_prototype(SPIRFunction &func, const Bitset &return_flags);
SPIRBlock *current_emitting_block = nullptr;
SPIRBlock *current_emitting_switch = nullptr;
SmallVector<SPIRBlock *> current_emitting_switch_stack;
bool current_emitting_switch_fallthrough = false;
virtual void emit_instruction(const Instruction &instr);
struct TemporaryCopy
{
uint32_t dst_id;
uint32_t src_id;
};
TemporaryCopy handle_instruction_precision(const Instruction &instr);
void emit_block_instructions(SPIRBlock &block);
void emit_block_instructions_with_masked_debug(SPIRBlock &block);
// For relax_nan_checks.
GLSLstd450 get_remapped_glsl_op(GLSLstd450 std450_op) const;
spv::Op get_remapped_spirv_op(spv::Op op) const;
virtual void emit_glsl_op(uint32_t result_type, uint32_t result_id, uint32_t op, const uint32_t *args,
uint32_t count);
virtual void emit_spv_amd_shader_ballot_op(uint32_t result_type, uint32_t result_id, uint32_t op,
@@ -385,7 +439,9 @@ protected:
const std::string &qualifier = "", uint32_t base_offset = 0);
virtual void emit_struct_padding_target(const SPIRType &type);
virtual std::string image_type_glsl(const SPIRType &type, uint32_t id = 0);
std::string constant_expression(const SPIRConstant &c);
std::string constant_expression(const SPIRConstant &c,
bool inside_block_like_struct_scope = false,
bool inside_struct_scope = false);
virtual std::string constant_op_expression(const SPIRConstantOp &cop);
virtual std::string constant_expression_vector(const SPIRConstant &c, uint32_t vector);
virtual void emit_fixup();
@@ -419,7 +475,7 @@ protected:
TextureFunctionArguments() = default;
TextureFunctionBaseArguments base;
uint32_t coord = 0, coord_components = 0, dref = 0;
uint32_t grad_x = 0, grad_y = 0, lod = 0, coffset = 0, offset = 0;
uint32_t grad_x = 0, grad_y = 0, lod = 0, offset = 0;
uint32_t bias = 0, component = 0, sample = 0, sparse_texel = 0, min_lod = 0;
bool nonuniform_expression = false;
};
@@ -436,6 +492,8 @@ protected:
virtual bool builtin_translates_to_nonarray(spv::BuiltIn builtin) const;
virtual bool is_user_type_structured(uint32_t id) const;
void emit_copy_logical_type(uint32_t lhs_id, uint32_t lhs_type_id, uint32_t rhs_id, uint32_t rhs_type_id,
SmallVector<uint32_t> chain);
@@ -495,6 +553,8 @@ protected:
// on a single line separated by comma.
SmallVector<std::string> *redirect_statement = nullptr;
const SPIRBlock *current_continue_block = nullptr;
bool block_temporary_hoisting = false;
bool block_debug_directives = false;
void begin_scope();
void end_scope();
@@ -522,7 +582,8 @@ protected:
bool member_is_remapped_physical_type(const SPIRType &type, uint32_t index) const;
bool member_is_packed_physical_type(const SPIRType &type, uint32_t index) const;
virtual std::string convert_row_major_matrix(std::string exp_str, const SPIRType &exp_type,
uint32_t physical_type_id, bool is_packed);
uint32_t physical_type_id, bool is_packed,
bool relaxed = false);
std::unordered_set<std::string> local_variable_names;
std::unordered_set<std::string> resource_names;
@@ -559,6 +620,7 @@ protected:
const char *uint16_t_literal_suffix = "us";
const char *nonuniform_qualifier = "nonuniformEXT";
const char *boolean_mix_function = "mix";
SPIRType::BaseType boolean_in_struct_remapped_type = SPIRType::Boolean;
bool swizzle_is_function = false;
bool shared_is_implied = false;
bool unsized_array_supported = true;
@@ -572,12 +634,13 @@ protected:
bool allow_precision_qualifiers = false;
bool can_swizzle_scalar = false;
bool force_gl_in_out_block = false;
bool force_merged_mesh_block = false;
bool can_return_array = true;
bool allow_truncated_access_chain = false;
bool supports_extensions = false;
bool supports_empty_struct = false;
bool array_is_value_type = true;
bool buffer_offset_array_is_value_type = true;
bool array_is_value_type_in_buffer_blocks = true;
bool comparison_image_samples_scalar = false;
bool native_pointers = false;
bool support_small_type_sampling_result = false;
@@ -587,16 +650,21 @@ protected:
bool support_pointer_to_pointer = false;
bool support_precise_qualifier = false;
bool support_64bit_switch = false;
bool workgroup_size_is_hidden = false;
bool requires_relaxed_precision_analysis = false;
bool implicit_c_integer_promotion_rules = false;
} backend;
void emit_struct(SPIRType &type);
void emit_resources();
void emit_extension_workarounds(spv::ExecutionModel model);
void emit_subgroup_arithmetic_workaround(const std::string &func, spv::Op op, spv::GroupOperation group_op);
void emit_polyfills(uint32_t polyfills, bool relaxed);
void emit_buffer_block_native(const SPIRVariable &var);
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);
@@ -610,6 +678,7 @@ protected:
void emit_block_chain(SPIRBlock &block);
void emit_hoisted_temporaries(SmallVector<std::pair<TypeID, ID>> &temporaries);
std::string constant_value_macro_name(uint32_t id);
int get_constant_mapping_to_workgroup_component(const SPIRConstant &constant) const;
void emit_constant(const SPIRConstant &constant);
void emit_specialization_constant_op(const SPIRConstantOp &constant);
std::string emit_continue_block(uint32_t continue_block, bool follow_true_block, bool follow_false_block);
@@ -628,6 +697,7 @@ protected:
bool should_suppress_usage_tracking(uint32_t id) const;
void emit_mix_op(uint32_t result_type, uint32_t id, uint32_t left, uint32_t right, uint32_t lerp);
void emit_nminmax_op(uint32_t result_type, uint32_t id, uint32_t op0, uint32_t op1, GLSLstd450 op);
void emit_emulated_ahyper_op(uint32_t result_type, uint32_t result_id, uint32_t op0, GLSLstd450 op);
bool to_trivial_mix_op(const SPIRType &type, std::string &op, uint32_t left, uint32_t right, uint32_t lerp);
void emit_quaternary_func_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, uint32_t op2,
uint32_t op3, const char *op);
@@ -658,7 +728,7 @@ protected:
void emit_unrolled_binary_op(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op,
bool negate, SPIRType::BaseType expected_type);
void emit_binary_op_cast(uint32_t result_type, uint32_t result_id, uint32_t op0, uint32_t op1, const char *op,
SPIRType::BaseType input_type, bool skip_cast_if_equal_type);
SPIRType::BaseType input_type, bool skip_cast_if_equal_type, bool implicit_integer_promotion);
SPIRType binary_op_bitcast_helper(std::string &cast_op0, std::string &cast_op1, SPIRType::BaseType &input_type,
uint32_t op0, uint32_t op1, bool skip_cast_if_equal_type);
@@ -669,6 +739,8 @@ protected:
uint32_t false_value);
void emit_unary_op(uint32_t result_type, uint32_t result_id, uint32_t op0, const char *op);
void emit_unary_op_cast(uint32_t result_type, uint32_t result_id, uint32_t op0, const char *op);
virtual void emit_mesh_tasks(SPIRBlock &block);
bool expression_is_forwarded(uint32_t id) const;
bool expression_suppresses_usage_tracking(uint32_t id) const;
bool expression_read_implies_multiple_reads(uint32_t id) const;
@@ -684,7 +756,8 @@ protected:
spv::StorageClass get_expression_effective_storage_class(uint32_t ptr);
virtual bool access_chain_needs_stage_io_builtin_translation(uint32_t base);
virtual void prepare_access_chain_for_scalar_access(std::string &expr, const SPIRType &type,
virtual void check_physical_type_cast(std::string &expr, const SPIRType *type, uint32_t physical_type);
virtual bool prepare_access_chain_for_scalar_access(std::string &expr, const SPIRType &type,
spv::StorageClass storage, bool &is_packed);
std::string access_chain(uint32_t base, const uint32_t *indices, uint32_t count, const SPIRType &target_type,
@@ -716,8 +789,8 @@ protected:
void append_global_func_args(const SPIRFunction &func, uint32_t index, SmallVector<std::string> &arglist);
std::string to_non_uniform_aware_expression(uint32_t id);
std::string to_expression(uint32_t id, bool register_expression_read = true);
std::string to_composite_constructor_expression(uint32_t id, bool uses_buffer_offset);
std::string to_rerolled_array_expression(const std::string &expr, const SPIRType &type);
std::string to_composite_constructor_expression(const SPIRType &parent_type, uint32_t id, bool block_like_type);
std::string to_rerolled_array_expression(const SPIRType &parent_type, const std::string &expr, const SPIRType &type);
std::string to_enclosed_expression(uint32_t id, bool register_expression_read = true);
std::string to_unpacked_expression(uint32_t id, bool register_expression_read = true);
std::string to_unpacked_row_major_matrix_expression(uint32_t id);
@@ -728,17 +801,18 @@ protected:
std::string to_extract_component_expression(uint32_t id, uint32_t index);
std::string to_extract_constant_composite_expression(uint32_t result_type, const SPIRConstant &c,
const uint32_t *chain, uint32_t length);
static bool needs_enclose_expression(const std::string &expr);
std::string enclose_expression(const std::string &expr);
std::string dereference_expression(const SPIRType &expression_type, const std::string &expr);
std::string address_of_expression(const std::string &expr);
void strip_enclosed_expression(std::string &expr);
std::string to_member_name(const SPIRType &type, uint32_t index);
virtual std::string to_member_reference(uint32_t base, const SPIRType &type, uint32_t index, bool ptr_chain);
virtual std::string to_member_reference(uint32_t base, const SPIRType &type, uint32_t index, bool ptr_chain_is_resolved);
std::string to_multi_member_reference(const SPIRType &type, const SmallVector<uint32_t> &indices);
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);
@@ -749,7 +823,7 @@ protected:
std::string layout_for_variable(const SPIRVariable &variable);
std::string to_combined_image_sampler(VariableID image_id, VariableID samp_id);
virtual bool skip_argument(uint32_t id) const;
virtual void emit_array_copy(const std::string &lhs, uint32_t lhs_id, uint32_t rhs_id,
virtual bool emit_array_copy(const char *expr, uint32_t lhs_id, uint32_t rhs_id,
spv::StorageClass lhs_storage, spv::StorageClass rhs_storage);
virtual void emit_block_hints(const SPIRBlock &block);
virtual std::string to_initializer_expression(const SPIRVariable &var);
@@ -789,6 +863,10 @@ protected:
void replace_fragment_outputs();
std::string legacy_tex_op(const std::string &op, const SPIRType &imgtype, uint32_t id);
void forward_relaxed_precision(uint32_t dst_id, const uint32_t *args, uint32_t length);
void analyze_precision_requirements(uint32_t type_id, uint32_t dst_id, uint32_t *args, uint32_t length);
Options::Precision analyze_expression_precision(const uint32_t *args, uint32_t length) const;
uint32_t indent = 0;
std::unordered_set<uint32_t> emitted_functions;
@@ -842,10 +920,25 @@ protected:
return !options.es && options.version < 130;
}
bool requires_transpose_2x2 = false;
bool requires_transpose_3x3 = false;
bool requires_transpose_4x4 = false;
enum Polyfill : uint32_t
{
PolyfillTranspose2x2 = 1 << 0,
PolyfillTranspose3x3 = 1 << 1,
PolyfillTranspose4x4 = 1 << 2,
PolyfillDeterminant2x2 = 1 << 3,
PolyfillDeterminant3x3 = 1 << 4,
PolyfillDeterminant4x4 = 1 << 5,
PolyfillMatrixInverse2x2 = 1 << 6,
PolyfillMatrixInverse3x3 = 1 << 7,
PolyfillMatrixInverse4x4 = 1 << 8,
};
uint32_t required_polyfills = 0;
uint32_t required_polyfills_relaxed = 0;
void require_polyfill(Polyfill polyfill, bool relaxed);
bool ray_tracing_is_khr = false;
bool barycentric_is_nv = false;
void ray_tracing_khr_fixup_locations();
bool args_will_forward(uint32_t id, const uint32_t *args, uint32_t num_args, bool pure);
@@ -879,8 +972,14 @@ protected:
void check_function_call_constraints(const uint32_t *args, uint32_t length);
void handle_invalid_expression(uint32_t id);
void force_temporary_and_recompile(uint32_t id);
void find_static_extensions();
uint32_t consume_temporary_in_precision_context(uint32_t type_id, uint32_t id, Options::Precision precision);
std::unordered_map<uint32_t, uint32_t> temporary_to_mirror_precision_alias;
std::unordered_set<uint32_t> composite_insert_overwritten;
std::unordered_set<uint32_t> block_composite_insert_overwrite;
std::string emit_for_loop_initializers(const SPIRBlock &block);
void emit_while_loop_initializers(const SPIRBlock &block);
bool for_loop_initializers_are_same_type(const SPIRBlock &block);
@@ -889,8 +988,6 @@ protected:
bool type_is_empty(const SPIRType &type);
virtual void declare_undefined_values();
bool can_use_io_location(spv::StorageClass storage, bool block);
const Instruction *get_next_instruction_in_block(const Instruction &instr);
static uint32_t mask_relevant_memory_semantics(uint32_t semantics);
@@ -904,6 +1001,7 @@ protected:
// Builtins in GLSL are always specific signedness, but the SPIR-V can declare them
// as either unsigned or signed.
// Sometimes we will need to automatically perform casts on load and store to make this work.
virtual SPIRType::BaseType get_builtin_basetype(spv::BuiltIn builtin, SPIRType::BaseType default_type);
virtual void cast_to_variable_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type);
virtual void cast_from_variable_load(uint32_t source_id, std::string &expr, const SPIRType &expr_type);
void unroll_array_from_complex_load(uint32_t target_id, uint32_t source_id, std::string &expr);
@@ -926,6 +1024,8 @@ protected:
void fixup_type_alias();
void reorder_type_alias();
void fixup_anonymous_struct_names();
void fixup_anonymous_struct_names(std::unordered_set<uint32_t> &visited, const SPIRType &type);
static const char *vector_swizzle(int vecsize, int index);
@@ -933,13 +1033,24 @@ protected:
bool is_stage_output_builtin_masked(spv::BuiltIn builtin) const;
bool is_stage_output_variable_masked(const SPIRVariable &var) const;
bool is_stage_output_block_member_masked(const SPIRVariable &var, uint32_t index, bool strip_array) const;
bool is_per_primitive_variable(const SPIRVariable &var) const;
uint32_t get_accumulated_member_location(const SPIRVariable &var, uint32_t mbr_idx, bool strip_array) const;
uint32_t get_declared_member_location(const SPIRVariable &var, uint32_t mbr_idx, bool strip_array) const;
std::unordered_set<LocationComponentPair, InternalHasher> masked_output_locations;
std::unordered_set<uint32_t> masked_output_builtins;
FloatFormatter *float_formatter = nullptr;
std::string format_float(float value) const;
std::string format_double(double value) const;
private:
void init();
SmallVector<ConstantID> get_composite_constant_ids(ConstantID const_id);
void fill_composite_constant(SPIRConstant &constant, TypeID type_id, const SmallVector<ConstantID> &initializers);
void set_composite_constant(ConstantID const_id, TypeID type_id, const SmallVector<ConstantID> &initializers);
TypeID get_composite_member_type(TypeID type_id, uint32_t member_idx);
std::unordered_map<uint32_t, SmallVector<ConstantID>> const_composite_insert_ids;
};
} // namespace SPIRV_CROSS_NAMESPACE
src/libraries/spirv_cross/spirv_hlsl.cpp
+1066 -140
View File
File diff suppressed because it is too large Load Diff
src/libraries/spirv_cross/spirv_hlsl.hpp
+42 -7
View File
@@ -98,6 +98,11 @@ struct HLSLResourceBinding
} cbv, uav, srv, sampler;
};
enum HLSLAuxBinding
{
HLSL_AUX_BINDING_BASE_VERTEX_INSTANCE = 0
};
class CompilerHLSL : public CompilerGLSL
{
public:
@@ -105,7 +110,7 @@ public:
{
uint32_t shader_model = 30; // TODO: map ps_4_0_level_9_0,... somehow
// Allows the PointSize builtin, and ignores it, as PointSize is not supported in HLSL.
// Allows the PointSize builtin in SM 4.0+, and ignores it, as PointSize is not supported in SM 4+.
bool point_size_compat = false;
// Allows the PointCoord builtin, returns float2(0.5, 0.5), as PointCoord is not supported in HLSL.
@@ -121,7 +126,7 @@ public:
// By default, a readonly storage buffer will be declared as ByteAddressBuffer (SRV) instead.
// Alternatively, use set_hlsl_force_storage_buffer_as_uav to specify individually.
bool force_storage_buffer_as_uav = false;
// Forces any storage image type marked as NonWritable to be considered an SRV instead.
// For this to work with function call parameters, NonWritable must be considered to be part of the type system
// so that NonWritable image arguments are also translated to Texture rather than RWTexture.
@@ -137,6 +142,14 @@ public:
// If add_vertex_attribute_remap is used and this feature is used,
// the semantic name will be queried once per active location.
bool flatten_matrix_vertex_input_semantics = false;
// Rather than emitting main() for the entry point, use the name in SPIR-V.
bool use_entry_point_name = false;
// Preserve (RW)StructuredBuffer types if the input source was HLSL.
// This relies on UserTypeGOOGLE to encode the buffer type either as "structuredbuffer" or "rwstructuredbuffer"
// whereas the type can be extended with an optional subtype, e.g. "structuredbuffer:int".
bool preserve_structured_buffers = false;
};
explicit CompilerHLSL(std::vector<uint32_t> spirv_)
@@ -208,6 +221,11 @@ public:
// Controls which storage buffer bindings will be forced to be declared as UAVs.
void set_hlsl_force_storage_buffer_as_uav(uint32_t desc_set, uint32_t binding);
// By default, these magic buffers are not assigned a specific binding.
void set_hlsl_aux_buffer_binding(HLSLAuxBinding binding, uint32_t register_index, uint32_t register_space);
void unset_hlsl_aux_buffer_binding(HLSLAuxBinding binding);
bool is_hlsl_aux_buffer_binding_used(HLSLAuxBinding binding) const;
private:
std::string type_to_glsl(const SPIRType &type, uint32_t id = 0) override;
std::string image_type_hlsl(const SPIRType &type, uint32_t id);
@@ -217,14 +235,13 @@ private:
void emit_hlsl_entry_point();
void emit_header() override;
void emit_resources();
void declare_undefined_values() override;
void emit_interface_block_globally(const SPIRVariable &type);
void emit_interface_block_in_struct(const SPIRVariable &var, std::unordered_set<uint32_t> &active_locations);
void emit_interface_block_member_in_struct(const SPIRVariable &var, uint32_t member_index,
uint32_t location,
void emit_interface_block_member_in_struct(const SPIRVariable &var, uint32_t member_index, uint32_t location,
std::unordered_set<uint32_t> &active_locations);
void emit_builtin_inputs_in_struct();
void emit_builtin_outputs_in_struct();
void emit_builtin_primitive_outputs_in_struct();
void emit_texture_op(const Instruction &i, bool sparse) override;
void emit_instruction(const Instruction &instruction) override;
void emit_glsl_op(uint32_t result_type, uint32_t result_id, uint32_t op, const uint32_t *args,
@@ -267,10 +284,14 @@ private:
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_rayquery_function(const char *commited, const char *candidate, const uint32_t *ops);
void emit_mesh_tasks(SPIRBlock &block) override;
const char *to_storage_qualifiers_glsl(const SPIRVariable &var) override;
void replace_illegal_names() override;
SPIRType::BaseType get_builtin_basetype(spv::BuiltIn builtin, SPIRType::BaseType default_type) override;
bool is_hlsl_force_storage_buffer_as_uav(ID id) const;
Options hlsl_options;
@@ -341,6 +362,10 @@ private:
TypeUnpackUint64
};
void analyze_meshlet_writes();
void analyze_meshlet_writes(uint32_t func_id, uint32_t id_per_vertex, uint32_t id_per_primitive,
std::unordered_set<uint32_t> &processed_func_ids);
BitcastType get_bitcast_type(uint32_t result_type, uint32_t op0);
void emit_builtin_variables();
@@ -369,10 +394,20 @@ private:
std::unordered_set<SetBindingPair, InternalHasher> force_uav_buffer_bindings;
// Returns true for BuiltInSampleMask because gl_SampleMask[] is an array in SPIR-V, but SV_Coverage is a scalar in HLSL.
bool builtin_translates_to_nonarray(spv::BuiltIn builtin) const override;
struct
{
uint32_t register_index = 0;
uint32_t register_space = 0;
bool explicit_binding = false;
bool used = false;
} base_vertex_info;
// Returns true if the specified ID has a UserTypeGOOGLE decoration for StructuredBuffer or RWStructuredBuffer resources.
bool is_user_type_structured(uint32_t id) const override;
std::vector<TypeID> composite_selection_workaround_types;
std::string get_inner_entry_point_name() const;
};
} // namespace SPIRV_CROSS_NAMESPACE
src/libraries/spirv_cross/spirv_msl.cpp
+2936 -939
View File
File diff suppressed because it is too large Load Diff
src/libraries/spirv_cross/spirv_msl.hpp
+207 -45
View File
@@ -34,36 +34,53 @@
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.
// Indicates the rate at which a variable changes value, one of: per-vertex,
// per-primitive, or per-patch.
enum MSLShaderVariableRate
{
MSL_SHADER_VARIABLE_RATE_PER_VERTEX = 0,
MSL_SHADER_VARIABLE_RATE_PER_PRIMITIVE = 1,
MSL_SHADER_VARIABLE_RATE_PER_PATCH = 2,
MSL_SHADER_VARIABLE_RATE_INT_MAX = 0x7fffffff,
};
// 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;
MSLShaderVariableRate rate = MSL_SHADER_VARIABLE_RATE_PER_VERTEX;
};
// Matches the binding index of a MSL resource for a binding within a descriptor set.
@@ -301,6 +318,7 @@ public:
uint32_t dynamic_offsets_buffer_index = 23;
uint32_t shader_input_buffer_index = 22;
uint32_t shader_index_buffer_index = 21;
uint32_t shader_patch_input_buffer_index = 20;
uint32_t shader_input_wg_index = 0;
uint32_t device_index = 0;
uint32_t enable_frag_output_mask = 0xffffffff;
@@ -321,10 +339,28 @@ public:
bool dispatch_base = false;
bool texture_1D_as_2D = false;
// Enable use of MSL 2.0 indirect argument buffers.
// Enable use of Metal argument buffers.
// MSL 2.0 must also be enabled.
bool argument_buffers = false;
// Defines Metal argument buffer tier levels.
// Uses same values as Metal MTLArgumentBuffersTier enumeration.
enum class ArgumentBuffersTier
{
Tier1 = 0,
Tier2 = 1,
};
// When using Metal argument buffers, indicates the Metal argument buffer tier level supported by the Metal platform.
// Ignored when Options::argument_buffers is disabled.
// - Tier1 supports writable images on macOS, but not on iOS.
// - Tier2 supports writable images on macOS and iOS, and higher resource count limits.
// Tier capabilities based on recommendations from Apple engineering.
ArgumentBuffersTier argument_buffers_tier = ArgumentBuffersTier::Tier1;
// Enables specifick argument buffer format with extra information to track SSBO-length
bool runtime_array_rich_descriptor = false;
// Ensures vertex and instance indices start at zero. This reflects the behavior of HLSL with SV_VertexID and SV_InstanceID.
bool enable_base_index_zero = false;
@@ -382,6 +418,11 @@ public:
// builtins are processed, but should result in more efficient usage of the GPU.
bool multi_patch_workgroup = false;
// Use storage buffers instead of vertex-style attributes for tessellation evaluation
// input. This may require conversion of inputs in the generated post-tessellation
// vertex shader, but allows the use of nested arrays.
bool raw_buffer_tese_input = false;
// If set, a vertex shader will be compiled as part of a tessellation pipeline.
// It will be translated as a compute kernel, so it can use the global invocation ID
// to index the output buffer.
@@ -393,7 +434,7 @@ public:
// and will be addressed using the current ViewIndex.
bool arrayed_subpass_input = false;
// Whether to use SIMD-group or quadgroup functions to implement group nnon-uniform
// Whether to use SIMD-group or quadgroup functions to implement group non-uniform
// operations. Some GPUs on iOS do not support the SIMD-group functions, only the
// quadgroup functions.
bool ios_use_simdgroup_functions = false;
@@ -435,6 +476,49 @@ public:
// the extra threads away.
bool force_sample_rate_shading = false;
// If set, gl_HelperInvocation will be set manually whenever a fragment is discarded.
// Some Metal devices have a bug where simd_is_helper_thread() does not return true
// after a fragment has been discarded. This is a workaround that is only expected to be needed
// until the bug is fixed in Metal; it is provided as an option to allow disabling it when that occurs.
bool manual_helper_invocation_updates = true;
// If set, extra checks will be emitted in fragment shaders to prevent writes
// from discarded fragments. Some Metal devices have a bug where writes to storage resources
// from discarded fragment threads continue to occur, despite the fragment being
// discarded. This is a workaround that is only expected to be needed until the
// bug is fixed in Metal; it is provided as an option so it can be enabled
// only when the bug is present.
bool check_discarded_frag_stores = false;
// If set, Lod operands to OpImageSample*DrefExplicitLod for 1D and 2D array images
// will be implemented using a gradient instead of passing the level operand directly.
// Some Metal devices have a bug where the level() argument to depth2d_array<T>::sample_compare()
// in a fragment shader is biased by some unknown amount, possibly dependent on the
// partial derivatives of the texture coordinates. This is a workaround that is only
// expected to be needed until the bug is fixed in Metal; it is provided as an option
// so it can be enabled only when the bug is present.
bool sample_dref_lod_array_as_grad = false;
// MSL doesn't guarantee coherence between writes and subsequent reads of read_write textures.
// This inserts fences before each read of a read_write texture to ensure coherency.
// If you're sure you never rely on this, you can set this to false for a possible performance improvement.
// Note: Only Apple's GPU compiler takes advantage of the lack of coherency, so make sure to test on Apple GPUs if you disable this.
bool readwrite_texture_fences = true;
// Metal 3.1 introduced a Metal regression bug which causes infinite recursion during
// Metal's analysis of an entry point input structure that is itself recursive. Enabling
// this option will replace the recursive input declaration with a alternate variable of
// type void*, and then cast to the correct type at the top of the entry point function.
// The bug has been reported to Apple, and will hopefully be fixed in future releases.
bool replace_recursive_inputs = false;
// If set, manual fixups of gradient vectors for cube texture lookups will be performed.
// All released Apple Silicon GPUs to date behave incorrectly when sampling a cube texture
// with explicit gradients. They will ignore one of the three partial derivatives based
// on the selected major axis, and expect the remaining derivatives to be partially
// transformed.
bool agx_manual_cube_grad_fixup = false;
bool is_ios() const
{
return platform == iOS;
@@ -445,6 +529,11 @@ public:
return platform == macOS;
}
bool use_quadgroup_operation() const
{
return is_ios() && !ios_use_simdgroup_functions;
}
void set_msl_version(uint32_t major, uint32_t minor = 0, uint32_t patch = 0)
{
msl_version = make_msl_version(major, minor, patch);
@@ -494,6 +583,11 @@ public:
return !buffers_requiring_array_length.empty();
}
bool buffer_requires_array_length(VariableID id) const
{
return buffers_requiring_array_length.count(id) != 0;
}
// Provide feedback to calling API to allow it to pass a buffer
// containing the view mask for the current multiview subpass.
bool needs_view_mask_buffer() const
@@ -534,10 +628,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
@@ -572,6 +671,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.
@@ -579,6 +681,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
@@ -633,7 +742,7 @@ public:
protected:
// An enum of SPIR-V functions that are implemented in additional
// source code that is added to the shader if necessary.
enum SPVFuncImpl
enum SPVFuncImpl : uint8_t
{
SPVFuncImplNone,
SPVFuncImplMod,
@@ -654,12 +763,14 @@ protected:
SPVFuncImplArrayOfArrayCopy6Dim = SPVFuncImplArrayCopyMultidimBase + 6,
SPVFuncImplTexelBufferCoords,
SPVFuncImplImage2DAtomicCoords, // Emulate texture2D atomic operations
SPVFuncImplGradientCube,
SPVFuncImplFMul,
SPVFuncImplFAdd,
SPVFuncImplFSub,
SPVFuncImplQuantizeToF16,
SPVFuncImplCubemapTo2DArrayFace,
SPVFuncImplUnsafeArray, // Allow Metal to use the array<T> template to make arrays a value type
SPVFuncImplStorageMatrix, // Allow threadgroup construction of matrices
SPVFuncImplInverse4x4,
SPVFuncImplInverse3x3,
SPVFuncImplInverse2x2,
@@ -708,11 +819,18 @@ protected:
SPVFuncImplConvertYCbCrBT601,
SPVFuncImplConvertYCbCrBT2020,
SPVFuncImplDynamicImageSampler,
SPVFuncImplRayQueryIntersectionParams,
SPVFuncImplVariableDescriptor,
SPVFuncImplVariableSizedDescriptor,
SPVFuncImplVariableDescriptorArray,
SPVFuncImplPaddedStd140
};
// 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,
@@ -731,6 +849,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;
@@ -738,9 +857,6 @@ protected:
std::string type_to_array_glsl(const SPIRType &type) override;
std::string constant_op_expression(const SPIRConstantOp &cop) override;
// Threadgroup arrays can't have a wrapper type
std::string variable_decl(const SPIRVariable &variable) override;
bool variable_decl_is_remapped_storage(const SPIRVariable &variable, spv::StorageClass storage) const override;
// GCC workaround of lambdas calling protected functions (for older GCC versions)
@@ -765,10 +881,9 @@ protected:
std::string bitcast_glsl_op(const SPIRType &result_type, const SPIRType &argument_type) override;
bool emit_complex_bitcast(uint32_t result_id, uint32_t id, uint32_t op0) override;
bool skip_argument(uint32_t id) const override;
std::string to_member_reference(uint32_t base, const SPIRType &type, uint32_t index, bool ptr_chain) override;
std::string to_member_reference(uint32_t base, const SPIRType &type, uint32_t index, bool ptr_chain_is_resolved) override;
std::string to_qualifiers_glsl(uint32_t id) override;
void replace_illegal_names() override;
void declare_undefined_values() override;
void declare_constant_arrays();
void replace_illegal_entry_point_names();
@@ -784,13 +899,17 @@ protected:
bool is_non_native_row_major_matrix(uint32_t id) override;
bool member_is_non_native_row_major_matrix(const SPIRType &type, uint32_t index) override;
std::string convert_row_major_matrix(std::string exp_str, const SPIRType &exp_type, uint32_t physical_type_id,
bool is_packed) override;
bool is_packed, bool relaxed) override;
bool is_tesc_shader() const;
bool is_tese_shader() const;
void preprocess_op_codes();
void localize_global_variables();
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<uint32_t, std::set<uint32_t>> function_global_vars;
void extract_global_variables_from_function(uint32_t func_id, std::set<uint32_t> &added_arg_ids,
@@ -826,13 +945,22 @@ protected:
bool add_component_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
SPIRVariable &var, const SPIRType &type,
InterfaceBlockMeta &meta);
void add_plain_member_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
SPIRType &ib_type, SPIRVariable &var, uint32_t index,
InterfaceBlockMeta &meta);
void add_composite_member_variable_to_interface_block(spv::StorageClass storage, const std::string &ib_var_ref,
SPIRType &ib_type, SPIRVariable &var, uint32_t index,
InterfaceBlockMeta &meta);
void add_plain_member_variable_to_interface_block(spv::StorageClass storage,
const std::string &ib_var_ref, SPIRType &ib_type,
SPIRVariable &var, SPIRType &var_type,
uint32_t mbr_idx, InterfaceBlockMeta &meta,
const std::string &mbr_name_qual,
const std::string &var_chain_qual,
uint32_t &location, uint32_t &var_mbr_idx);
void add_composite_member_variable_to_interface_block(spv::StorageClass storage,
const std::string &ib_var_ref, SPIRType &ib_type,
SPIRVariable &var, SPIRType &var_type,
uint32_t mbr_idx, InterfaceBlockMeta &meta,
const std::string &mbr_name_qual,
const std::string &var_chain_qual,
uint32_t &location, uint32_t &var_mbr_idx);
void add_tess_level_input_to_interface_block(const std::string &ib_var_ref, SPIRType &ib_type, SPIRVariable &var);
void add_tess_level_input(const std::string &base_ref, const std::string &mbr_name, SPIRVariable &var);
void fix_up_interface_member_indices(spv::StorageClass storage, uint32_t ib_type_id);
@@ -848,7 +976,8 @@ protected:
void emit_specialization_constants_and_structs();
void emit_interface_block(uint32_t ib_var_id);
bool maybe_emit_array_assignment(uint32_t id_lhs, uint32_t id_rhs);
uint32_t get_resource_array_size(uint32_t id) const;
bool is_var_runtime_size_array(const SPIRVariable &var) const;
uint32_t get_resource_array_size(const SPIRType &type, uint32_t id) const;
void fix_up_shader_inputs_outputs();
@@ -858,7 +987,7 @@ protected:
std::string entry_point_arg_stage_in();
void entry_point_args_builtin(std::string &args);
void entry_point_args_discrete_descriptors(std::string &args);
std::string to_qualified_member_name(const SPIRType &type, uint32_t index);
std::string append_member_name(const std::string &qualifier, const SPIRType &type, uint32_t index);
std::string ensure_valid_name(std::string name, std::string pfx);
std::string to_sampler_expression(uint32_t id);
std::string to_swizzle_expression(uint32_t id);
@@ -872,11 +1001,14 @@ protected:
std::string member_attribute_qualifier(const SPIRType &type, uint32_t index);
std::string member_location_attribute_qualifier(const SPIRType &type, uint32_t index);
std::string argument_decl(const SPIRFunction::Parameter &arg);
const char *descriptor_address_space(uint32_t id, spv::StorageClass storage, const char *plain_address_space) const;
std::string round_fp_tex_coords(std::string tex_coords, bool coord_is_fp);
uint32_t get_metal_resource_index(SPIRVariable &var, SPIRType::BaseType basetype, uint32_t plane = 0);
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;
@@ -912,7 +1044,7 @@ protected:
bool validate_member_packing_rules_msl(const SPIRType &type, uint32_t index) const;
std::string get_argument_address_space(const SPIRVariable &argument);
std::string get_type_address_space(const SPIRType &type, uint32_t id, bool argument = false);
const char *to_restrict(uint32_t id, bool space = true);
const char *to_restrict(uint32_t id, bool space);
SPIRType &get_stage_in_struct_type();
SPIRType &get_stage_out_struct_type();
SPIRType &get_patch_stage_in_struct_type();
@@ -920,14 +1052,14 @@ protected:
std::string get_tess_factor_struct_name();
SPIRType &get_uint_type();
uint32_t get_uint_type_id();
void emit_atomic_func_op(uint32_t result_type, uint32_t result_id, const char *op, uint32_t mem_order_1,
uint32_t mem_order_2, bool has_mem_order_2, uint32_t op0, uint32_t op1 = 0,
void emit_atomic_func_op(uint32_t result_type, uint32_t result_id, const char *op, spv::Op opcode,
uint32_t mem_order_1, uint32_t mem_order_2, bool has_mem_order_2, uint32_t op0, uint32_t op1 = 0,
bool op1_is_pointer = false, bool op1_is_literal = false, uint32_t op2 = 0);
const char *get_memory_order(uint32_t spv_mem_sem);
void add_pragma_line(const std::string &line);
void add_typedef_line(const std::string &line);
void emit_barrier(uint32_t id_exe_scope, uint32_t id_mem_scope, uint32_t id_mem_sem);
void emit_array_copy(const std::string &lhs, uint32_t lhs_id, uint32_t rhs_id,
bool emit_array_copy(const char *expr, uint32_t lhs_id, uint32_t rhs_id,
spv::StorageClass lhs_storage, spv::StorageClass rhs_storage) override;
void build_implicit_builtins();
uint32_t build_constant_uint_array_pointer();
@@ -937,6 +1069,7 @@ protected:
uint32_t builtin_frag_coord_id = 0;
uint32_t builtin_sample_id_id = 0;
uint32_t builtin_sample_mask_id = 0;
uint32_t builtin_helper_invocation_id = 0;
uint32_t builtin_vertex_idx_id = 0;
uint32_t builtin_base_vertex_id = 0;
uint32_t builtin_instance_idx_id = 0;
@@ -961,6 +1094,7 @@ protected:
uint32_t argument_buffer_padding_sampler_type_id = 0;
bool does_shader_write_sample_mask = false;
bool frag_shader_needs_discard_checks = false;
void cast_to_variable_store(uint32_t target_id, std::string &expr, const SPIRType &expr_type) override;
void cast_from_variable_load(uint32_t source_id, std::string &expr, const SPIRType &expr_type) override;
@@ -969,9 +1103,11 @@ protected:
void analyze_sampled_image_usage();
bool access_chain_needs_stage_io_builtin_translation(uint32_t base) override;
void prepare_access_chain_for_scalar_access(std::string &expr, const SPIRType &type, spv::StorageClass storage,
bool prepare_access_chain_for_scalar_access(std::string &expr, const SPIRType &type, spv::StorageClass storage,
bool &is_packed) override;
void fix_up_interpolant_access_chain(const uint32_t *ops, uint32_t length);
void check_physical_type_cast(std::string &expr, const SPIRType *type, uint32_t physical_type) override;
bool emit_tessellation_access_chain(const uint32_t *ops, uint32_t length);
bool emit_tessellation_io_load(uint32_t result_type, uint32_t id, uint32_t ptr);
bool is_out_of_bounds_tessellation_level(uint32_t id_lhs);
@@ -985,12 +1121,17 @@ protected:
Options msl_options;
std::set<SPVFuncImpl> spv_function_implementations;
// Must be ordered to ensure declarations are in a specific order.
std::map<LocationComponentPair, MSLShaderInput> inputs_by_location;
std::unordered_map<uint32_t, MSLShaderInput> inputs_by_builtin;
std::map<LocationComponentPair, MSLShaderInterfaceVariable> inputs_by_location;
std::unordered_map<uint32_t, MSLShaderInterfaceVariable> inputs_by_builtin;
std::map<LocationComponentPair, MSLShaderInterfaceVariable> outputs_by_location;
std::unordered_map<uint32_t, MSLShaderInterfaceVariable> outputs_by_builtin;
std::unordered_set<uint32_t> location_inputs_in_use;
std::unordered_set<uint32_t> location_inputs_in_use_fallback;
std::unordered_set<uint32_t> location_outputs_in_use;
std::unordered_set<uint32_t> location_outputs_in_use_fallback;
std::unordered_map<uint32_t, uint32_t> fragment_output_components;
std::unordered_map<uint32_t, uint32_t> builtin_to_automatic_input_location;
std::unordered_map<uint32_t, uint32_t> builtin_to_automatic_output_location;
std::set<std::string> pragma_lines;
std::set<std::string> typedef_lines;
SmallVector<uint32_t> vars_needing_early_declaration;
@@ -1010,6 +1151,8 @@ protected:
VariableID patch_stage_out_var_id = 0;
VariableID stage_in_ptr_var_id = 0;
VariableID stage_out_ptr_var_id = 0;
VariableID tess_level_inner_var_id = 0;
VariableID tess_level_outer_var_id = 0;
VariableID stage_out_masked_builtin_type_id = 0;
// Handle HLSL-style 0-based vertex/instance index.
@@ -1036,6 +1179,7 @@ protected:
bool needs_subgroup_invocation_id = false;
bool needs_subgroup_size = false;
bool needs_sample_id = false;
bool needs_helper_invocation = false;
std::string qual_pos_var_name;
std::string stage_in_var_name = "in";
std::string stage_out_var_name = "out";
@@ -1048,6 +1192,7 @@ protected:
std::string input_wg_var_name = "gl_in";
std::string input_buffer_var_name = "spvIn";
std::string output_buffer_var_name = "spvOut";
std::string patch_input_buffer_var_name = "spvPatchIn";
std::string patch_output_buffer_var_name = "spvPatchOut";
std::string tess_factor_buffer_var_name = "spvTessLevel";
std::string index_buffer_var_name = "spvIndices";
@@ -1059,9 +1204,13 @@ protected:
const MSLConstexprSampler *find_constexpr_sampler(uint32_t id) const;
std::unordered_set<uint32_t> buffers_requiring_array_length;
SmallVector<uint32_t> buffer_arrays;
std::unordered_set<uint32_t> atomic_image_vars; // Emulate texture2D atomic operations
SmallVector<std::pair<uint32_t, uint32_t>> buffer_aliases_argument;
SmallVector<uint32_t> buffer_aliases_discrete;
std::unordered_set<uint32_t> atomic_image_vars_emulated; // Emulate texture2D atomic operations
std::unordered_set<uint32_t> pull_model_inputs;
std::unordered_set<uint32_t> recursive_inputs;
SmallVector<SPIRVariable *> entry_point_bindings;
// Must be ordered since array is in a specific order.
std::map<SetBindingPair, std::pair<uint32_t, uint32_t>> buffers_requiring_dynamic_offset;
@@ -1100,6 +1249,16 @@ protected:
bool variable_storage_requires_stage_io(spv::StorageClass storage) const;
bool needs_manual_helper_invocation_updates() const
{
return msl_options.manual_helper_invocation_updates && msl_options.supports_msl_version(2, 3);
}
bool needs_frag_discard_checks() const
{
return get_execution_model() == spv::ExecutionModelFragment && msl_options.supports_msl_version(2, 3) &&
msl_options.check_discarded_frag_stores && frag_shader_needs_discard_checks;
}
bool has_additional_fixed_sample_mask() const { return msl_options.additional_fixed_sample_mask != 0xffffffff; }
std::string additional_fixed_sample_mask_str() const;
@@ -1117,13 +1276,16 @@ protected:
CompilerMSL &compiler;
std::unordered_map<uint32_t, uint32_t> result_types;
std::unordered_map<uint32_t, uint32_t> image_pointers; // Emulate texture2D atomic operations
std::unordered_map<uint32_t, uint32_t> image_pointers_emulated; // Emulate texture2D atomic operations
bool suppress_missing_prototypes = false;
bool uses_atomics = false;
bool uses_resource_write = false;
bool uses_image_write = false;
bool uses_buffer_write = false;
bool uses_discard = false;
bool needs_subgroup_invocation_id = false;
bool needs_subgroup_size = false;
bool needs_sample_id = false;
bool needs_helper_invocation = false;
};
// OpcodeHandler that scans for uses of sampled images
src/libraries/spirv_cross/spirv_parser.cpp
+110 -38
View File
@@ -31,7 +31,7 @@ namespace SPIRV_CROSS_NAMESPACE
{
Parser::Parser(vector<uint32_t> spirv)
{
ir.spirv = move(spirv);
ir.spirv = std::move(spirv);
}
Parser::Parser(const uint32_t *spirv_data, size_t word_count)
@@ -68,6 +68,7 @@ static bool is_valid_spirv_version(uint32_t version)
case 0x10300: // SPIR-V 1.3
case 0x10400: // SPIR-V 1.4
case 0x10500: // SPIR-V 1.5
case 0x10600: // SPIR-V 1.6
return true;
default:
@@ -182,6 +183,15 @@ void Parser::parse(const Instruction &instruction)
auto op = static_cast<Op>(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:
@@ -258,29 +268,37 @@ void Parser::parse(const Instruction &instruction)
case OpExtension:
{
auto ext = extract_string(ir.spirv, instruction.offset);
ir.declared_extensions.push_back(move(ext));
ir.declared_extensions.push_back(std::move(ext));
break;
}
case OpExtInstImport:
{
uint32_t id = ops[0];
SPIRExtension::Extension spirv_ext = SPIRExtension::Unsupported;
auto ext = extract_string(ir.spirv, instruction.offset + 1);
if (ext == "GLSL.std.450")
set<SPIRExtension>(id, SPIRExtension::GLSL);
spirv_ext = SPIRExtension::GLSL;
else if (ext == "DebugInfo")
set<SPIRExtension>(id, SPIRExtension::SPV_debug_info);
spirv_ext = SPIRExtension::SPV_debug_info;
else if (ext == "SPV_AMD_shader_ballot")
set<SPIRExtension>(id, SPIRExtension::SPV_AMD_shader_ballot);
spirv_ext = SPIRExtension::SPV_AMD_shader_ballot;
else if (ext == "SPV_AMD_shader_explicit_vertex_parameter")
set<SPIRExtension>(id, SPIRExtension::SPV_AMD_shader_explicit_vertex_parameter);
spirv_ext = SPIRExtension::SPV_AMD_shader_explicit_vertex_parameter;
else if (ext == "SPV_AMD_shader_trinary_minmax")
set<SPIRExtension>(id, SPIRExtension::SPV_AMD_shader_trinary_minmax);
spirv_ext = SPIRExtension::SPV_AMD_shader_trinary_minmax;
else if (ext == "SPV_AMD_gcn_shader")
set<SPIRExtension>(id, SPIRExtension::SPV_AMD_gcn_shader);
else
set<SPIRExtension>(id, SPIRExtension::Unsupported);
spirv_ext = SPIRExtension::SPV_AMD_gcn_shader;
else if (ext == "NonSemantic.DebugPrintf")
spirv_ext = SPIRExtension::NonSemanticDebugPrintf;
else if (ext == "NonSemantic.Shader.DebugInfo.100")
spirv_ext = SPIRExtension::NonSemanticShaderDebugInfo;
else if (ext.find("NonSemantic.") == 0)
spirv_ext = SPIRExtension::NonSemanticGeneric;
set<SPIRExtension>(id, spirv_ext);
// Other SPIR-V extensions which have ExtInstrs are currently not supported.
break;
@@ -290,7 +308,15 @@ void Parser::parse(const Instruction &instruction)
{
// The SPIR-V debug information extended instructions might come at global scope.
if (current_block)
{
current_block->ops.push_back(instruction);
if (length >= 2)
{
const auto *type = maybe_get<SPIRType>(ops[0]);
if (type)
ir.load_type_width.insert({ ops[1], type->width });
}
}
break;
}
@@ -338,12 +364,32 @@ void Parser::parse(const Instruction &instruction)
execution.output_vertices = ops[2];
break;
case ExecutionModeOutputPrimitivesEXT:
execution.output_primitives = ops[2];
break;
default:
break;
}
break;
}
case OpExecutionModeId:
{
auto &execution = ir.entry_points[ops[0]];
auto mode = static_cast<ExecutionMode>(ops[1]);
execution.flags.set(mode);
if (mode == ExecutionModeLocalSizeId)
{
execution.workgroup_size.id_x = ops[2];
execution.workgroup_size.id_y = ops[3];
execution.workgroup_size.id_z = ops[4];
}
break;
}
case OpName:
{
uint32_t id = ops[0];
@@ -471,7 +517,7 @@ void Parser::parse(const Instruction &instruction)
case OpTypeVoid:
{
uint32_t id = ops[0];
auto &type = set<SPIRType>(id);
auto &type = set<SPIRType>(id, op);
type.basetype = SPIRType::Void;
break;
}
@@ -479,7 +525,7 @@ void Parser::parse(const Instruction &instruction)
case OpTypeBool:
{
uint32_t id = ops[0];
auto &type = set<SPIRType>(id);
auto &type = set<SPIRType>(id, op);
type.basetype = SPIRType::Boolean;
type.width = 1;
break;
@@ -489,7 +535,7 @@ void Parser::parse(const Instruction &instruction)
{
uint32_t id = ops[0];
uint32_t width = ops[1];
auto &type = set<SPIRType>(id);
auto &type = set<SPIRType>(id, op);
if (width == 64)
type.basetype = SPIRType::Double;
else if (width == 32)
@@ -507,7 +553,7 @@ void Parser::parse(const Instruction &instruction)
uint32_t id = ops[0];
uint32_t width = ops[1];
bool signedness = ops[2] != 0;
auto &type = set<SPIRType>(id);
auto &type = set<SPIRType>(id, op);
type.basetype = signedness ? to_signed_basetype(width) : to_unsigned_basetype(width);
type.width = width;
break;
@@ -522,9 +568,9 @@ void Parser::parse(const Instruction &instruction)
uint32_t vecsize = ops[2];
auto &base = get<SPIRType>(ops[1]);
auto &vecbase = set<SPIRType>(id);
auto &vecbase = set<SPIRType>(id, base);
vecbase = base;
vecbase.op = op;
vecbase.vecsize = vecsize;
vecbase.self = id;
vecbase.parent_type = ops[1];
@@ -537,9 +583,9 @@ void Parser::parse(const Instruction &instruction)
uint32_t colcount = ops[2];
auto &base = get<SPIRType>(ops[1]);
auto &matrixbase = set<SPIRType>(id);
auto &matrixbase = set<SPIRType>(id, base);
matrixbase = base;
matrixbase.op = op;
matrixbase.columns = colcount;
matrixbase.self = id;
matrixbase.parent_type = ops[1];
@@ -549,12 +595,11 @@ void Parser::parse(const Instruction &instruction)
case OpTypeArray:
{
uint32_t id = ops[0];
auto &arraybase = set<SPIRType>(id);
uint32_t tid = ops[1];
auto &base = get<SPIRType>(tid);
auto &arraybase = set<SPIRType>(id, base);
arraybase = base;
arraybase.op = op;
arraybase.parent_type = tid;
uint32_t cid = ops[2];
@@ -569,7 +614,9 @@ void Parser::parse(const Instruction &instruction)
arraybase.array_size_literal.push_back(literal);
arraybase.array.push_back(literal ? c->scalar() : cid);
// Do NOT set arraybase.self!
// .self resolves down to non-array/non-pointer type.
arraybase.self = base.self;
break;
}
@@ -578,25 +625,27 @@ void Parser::parse(const Instruction &instruction)
uint32_t id = ops[0];
auto &base = get<SPIRType>(ops[1]);
auto &arraybase = set<SPIRType>(id);
auto &arraybase = set<SPIRType>(id, base);
// We're copying type information into Array types, so we'll need a fixup for any physical pointer
// references.
if (base.forward_pointer)
forward_pointer_fixups.push_back({ id, ops[1] });
arraybase = base;
arraybase.op = op;
arraybase.array.push_back(0);
arraybase.array_size_literal.push_back(true);
arraybase.parent_type = ops[1];
// Do NOT set arraybase.self!
// .self resolves down to non-array/non-pointer type.
arraybase.self = base.self;
break;
}
case OpTypeImage:
{
uint32_t id = ops[0];
auto &type = set<SPIRType>(id);
auto &type = set<SPIRType>(id, op);
type.basetype = SPIRType::Image;
type.image.type = ops[1];
type.image.dim = static_cast<Dim>(ops[2]);
@@ -613,7 +662,7 @@ void Parser::parse(const Instruction &instruction)
{
uint32_t id = ops[0];
uint32_t imagetype = ops[1];
auto &type = set<SPIRType>(id);
auto &type = set<SPIRType>(id, op);
type = get<SPIRType>(imagetype);
type.basetype = SPIRType::SampledImage;
type.self = id;
@@ -623,7 +672,7 @@ void Parser::parse(const Instruction &instruction)
case OpTypeSampler:
{
uint32_t id = ops[0];
auto &type = set<SPIRType>(id);
auto &type = set<SPIRType>(id, op);
type.basetype = SPIRType::Sampler;
break;
}
@@ -636,10 +685,13 @@ void Parser::parse(const Instruction &instruction)
// We won't be able to compile it, but we shouldn't crash when parsing.
// We should be able to reflect.
auto *base = maybe_get<SPIRType>(ops[2]);
auto &ptrbase = set<SPIRType>(id);
auto &ptrbase = set<SPIRType>(id, op);
if (base)
{
ptrbase = *base;
ptrbase.op = op;
}
ptrbase.pointer = true;
ptrbase.pointer_depth++;
@@ -660,7 +712,7 @@ void Parser::parse(const Instruction &instruction)
case OpTypeForwardPointer:
{
uint32_t id = ops[0];
auto &ptrbase = set<SPIRType>(id);
auto &ptrbase = set<SPIRType>(id, op);
ptrbase.pointer = true;
ptrbase.pointer_depth++;
ptrbase.storage = static_cast<StorageClass>(ops[1]);
@@ -675,7 +727,7 @@ void Parser::parse(const Instruction &instruction)
case OpTypeStruct:
{
uint32_t id = ops[0];
auto &type = set<SPIRType>(id);
auto &type = set<SPIRType>(id, op);
type.basetype = SPIRType::Struct;
for (uint32_t i = 1; i < length; i++)
type.member_types.push_back(ops[i]);
@@ -724,7 +776,7 @@ void Parser::parse(const Instruction &instruction)
case OpTypeAccelerationStructureKHR:
{
uint32_t id = ops[0];
auto &type = set<SPIRType>(id);
auto &type = set<SPIRType>(id, op);
type.basetype = SPIRType::AccelerationStructure;
break;
}
@@ -732,7 +784,7 @@ void Parser::parse(const Instruction &instruction)
case OpTypeRayQueryKHR:
{
uint32_t id = ops[0];
auto &type = set<SPIRType>(id);
auto &type = set<SPIRType>(id, op);
type.basetype = SPIRType::RayQuery;
break;
}
@@ -979,10 +1031,9 @@ void Parser::parse(const Instruction &instruction)
{
uint32_t ids = ir.increase_bound_by(2);
SPIRType type;
auto &type = set<SPIRType>(ids, OpTypeInt);
type.basetype = SPIRType::Int;
type.width = 32;
set<SPIRType>(ids, type);
auto &c = set<SPIRConstant>(ids + 1, ids);
current_block->condition = c.self;
@@ -993,7 +1044,16 @@ void Parser::parse(const Instruction &instruction)
}
else
{
ir.block_meta[current_block->next_block] &= ~ParsedIR::BLOCK_META_SELECTION_MERGE_BIT;
// Collapse loops if we have to.
bool collapsed_loop = current_block->true_block == current_block->merge_block &&
current_block->merge == SPIRBlock::MergeLoop;
if (collapsed_loop)
{
ir.block_meta[current_block->merge_block] &= ~ParsedIR::BLOCK_META_LOOP_MERGE_BIT;
ir.block_meta[current_block->continue_block] &= ~ParsedIR::BLOCK_META_CONTINUE_BIT;
}
current_block->next_block = current_block->true_block;
current_block->condition = 0;
current_block->true_block = 0;
@@ -1042,6 +1102,7 @@ void Parser::parse(const Instruction &instruction)
}
case OpKill:
case OpTerminateInvocation:
{
if (!current_block)
SPIRV_CROSS_THROW("Trying to end a non-existing block.");
@@ -1066,6 +1127,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)
@@ -1194,10 +1267,9 @@ void Parser::parse(const Instruction &instruction)
{
const auto *type = maybe_get<SPIRType>(ops[0]);
if (type)
{
ir.load_type_width.insert({ ops[1], type->width });
}
}
if (!current_block)
SPIRV_CROSS_THROW("Currently no block to insert opcode.");
src/libraries/spirv_cross/spirv_parser.hpp
+2
View File
@@ -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;
src/libraries/spirv_cross/spirv_reflect.cpp
+17 -17
View File
@@ -291,8 +291,8 @@ static bool naturally_emit_type(const SPIRType &type)
bool CompilerReflection::type_is_reference(const SPIRType &type) const
{
// Physical pointers and arrays of physical pointers need to refer to the pointee's type.
return type_is_top_level_physical_pointer(type) ||
(!type.array.empty() && type_is_top_level_physical_pointer(get<SPIRType>(type.parent_type)));
return is_physical_pointer(type) ||
(type_is_array_of_pointers(type) && type.storage == StorageClassPhysicalStorageBuffer);
}
void CompilerReflection::emit_types()
@@ -341,7 +341,7 @@ void CompilerReflection::emit_type(uint32_t type_id, bool &emitted_open_tag)
json_stream->emit_json_key_object("_" + std::to_string(type_id));
json_stream->emit_json_key_value("name", name);
if (type_is_top_level_physical_pointer(type))
if (is_physical_pointer(type))
{
json_stream->emit_json_key_value("type", "_" + std::to_string(type.parent_type));
json_stream->emit_json_key_value("physical_pointer", true);
@@ -404,7 +404,7 @@ void CompilerReflection::emit_type_member(const SPIRType &type, uint32_t index)
void CompilerReflection::emit_type_array(const SPIRType &type)
{
if (!type_is_top_level_physical_pointer(type) && !type.array.empty())
if (!is_physical_pointer(type) && !type.array.empty())
{
json_stream->emit_json_key_array("array");
// Note that we emit the zeros here as a means of identifying
@@ -444,7 +444,7 @@ void CompilerReflection::emit_type_member_qualifiers(const SPIRType &type, uint3
if (dec.decoration_flags.get(DecorationRowMajor))
json_stream->emit_json_key_value("row_major", true);
if (type_is_top_level_physical_pointer(membertype))
if (is_physical_pointer(membertype))
json_stream->emit_json_key_value("physical_pointer", true);
}
}
@@ -587,18 +587,18 @@ void CompilerReflection::emit_resources(const char *tag, const SmallVector<Resou
{
bool ssbo_block = type.storage == StorageClassStorageBuffer ||
(type.storage == StorageClassUniform && typeflags.get(DecorationBufferBlock));
if (ssbo_block)
{
auto buffer_flags = get_buffer_block_flags(res.id);
if (buffer_flags.get(DecorationNonReadable))
json_stream->emit_json_key_value("writeonly", true);
if (buffer_flags.get(DecorationNonWritable))
json_stream->emit_json_key_value("readonly", true);
if (buffer_flags.get(DecorationRestrict))
json_stream->emit_json_key_value("restrict", true);
if (buffer_flags.get(DecorationCoherent))
json_stream->emit_json_key_value("coherent", true);
}
Bitset qualifier_mask = ssbo_block ? get_buffer_block_flags(res.id) : mask;
if (qualifier_mask.get(DecorationNonReadable))
json_stream->emit_json_key_value("writeonly", true);
if (qualifier_mask.get(DecorationNonWritable))
json_stream->emit_json_key_value("readonly", true);
if (qualifier_mask.get(DecorationRestrict))
json_stream->emit_json_key_value("restrict", true);
if (qualifier_mask.get(DecorationCoherent))
json_stream->emit_json_key_value("coherent", true);
if (qualifier_mask.get(DecorationVolatile))
json_stream->emit_json_key_value("volatile", true);
}
emit_type_array(type);