From affc050e21e60efebd65fc684096512aaf215d80 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=D0=91=D1=80=D0=B0=D0=BD=D0=B8=D0=BC=D0=B8=D1=80=20=D0=9A?= =?UTF-8?q?=D0=B0=D1=80=D0=B0=D1=9F=D0=B8=D1=9B?= Date: Sun, 14 Jun 2026 08:49:56 -0700 Subject: [PATCH] Updated spirv-cross. --- 3rdparty/spirv-cross/main.cpp | 12 + 3rdparty/spirv-cross/spirv.h | 629 ++++++++++---- 3rdparty/spirv-cross/spirv.hpp | 629 ++++++++++---- 3rdparty/spirv-cross/spirv_common.hpp | 26 +- 3rdparty/spirv-cross/spirv_cross.cpp | 272 +++++- 3rdparty/spirv-cross/spirv_cross.hpp | 42 + .../spirv-cross/spirv_cross_parsed_ir.cpp | 32 + .../spirv-cross/spirv_cross_parsed_ir.hpp | 7 + 3rdparty/spirv-cross/spirv_glsl.cpp | 807 +++++++++++++++--- 3rdparty/spirv-cross/spirv_glsl.hpp | 37 +- 3rdparty/spirv-cross/spirv_hlsl.cpp | 32 +- 3rdparty/spirv-cross/spirv_msl.cpp | 171 ++-- 3rdparty/spirv-cross/spirv_msl.hpp | 2 + 3rdparty/spirv-cross/spirv_parser.cpp | 89 +- 14 files changed, 2270 insertions(+), 517 deletions(-) diff --git a/3rdparty/spirv-cross/main.cpp b/3rdparty/spirv-cross/main.cpp index 7dc5404dd..759d38215 100644 --- a/3rdparty/spirv-cross/main.cpp +++ b/3rdparty/spirv-cross/main.cpp @@ -694,6 +694,8 @@ struct CLIArguments uint32_t glsl_ovr_multiview_view_count = 0; SmallVector> glsl_ext_framebuffer_fetch; bool glsl_ext_framebuffer_fetch_noncoherent = false; + uint32_t glsl_descriptor_heap_set = UINT32_MAX; + uint32_t glsl_descriptor_heap_binding = UINT32_MAX; bool vulkan_glsl_disable_ext_samplerless_texture_functions = false; bool emit_line_directives = false; bool enable_storage_image_qualifier_deduction = true; @@ -821,6 +823,7 @@ static void print_help_glsl() "\t\tPrimary use case is supporting external samplers in ESSL for video rendering on Android where you could remap a texture to a YUV one.\n" "\t[--glsl-force-flattened-io-blocks]:\n\t\tAlways flatten I/O blocks and structs.\n" "\t[--glsl-ovr-multiview-view-count count]:\n\t\tIn GL_OVR_multiview2, specify layout(num_views).\n" + "\t[--glsl-descriptor-heap-set-binding desc_set binding]:\n\t\tInstead of layout(descriptor_heap), emit layout(set = desc_set, binding = binding) instead for compatibility with mapping API.\n" ); // clang-format on } @@ -1458,6 +1461,10 @@ static string compile_iteration(const CLIArguments &args, std::vector opts.force_recompile_max_debug_iterations = args.force_recompile_max_debug_iterations; compiler->set_common_options(opts); + // This is enough for Vulkan mapping API. + if (args.glsl_descriptor_heap_set != UINT32_MAX) + compiler->remap_descriptor_heap(ResourceTypeUnknown, args.glsl_descriptor_heap_set, args.glsl_descriptor_heap_binding); + for (auto &fetch : args.glsl_ext_framebuffer_fetch) compiler->remap_ext_framebuffer_fetch(fetch.first, fetch.second, !args.glsl_ext_framebuffer_fetch_noncoherent); @@ -1704,6 +1711,11 @@ static int main_inner(int argc, char *argv[]) cbs.add("--glsl-ext-framebuffer-fetch-noncoherent", [&args](CLIParser &) { args.glsl_ext_framebuffer_fetch_noncoherent = true; }); + cbs.add("--glsl-descriptor-heap-set-binding", [&args](CLIParser &parser) + { + args.glsl_descriptor_heap_set = parser.next_uint(); + args.glsl_descriptor_heap_binding = parser.next_uint(); + }); cbs.add("--vulkan-glsl-disable-ext-samplerless-texture-functions", [&args](CLIParser &) { args.vulkan_glsl_disable_ext_samplerless_texture_functions = true; }); cbs.add("--disable-storage-image-qualifier-deduction", diff --git a/3rdparty/spirv-cross/spirv.h b/3rdparty/spirv-cross/spirv.h index 26bc6d9d8..d00cf0f4e 100644 --- a/3rdparty/spirv-cross/spirv.h +++ b/3rdparty/spirv-cross/spirv.h @@ -179,6 +179,7 @@ typedef enum SpvExecutionMode_ { SpvExecutionModeQuadDerivativesKHR = 5088, SpvExecutionModeRequireFullQuadsKHR = 5089, SpvExecutionModeSharesInputWithAMDX = 5102, + SpvExecutionModeArithmeticPoisonKHR = 5157, SpvExecutionModeOutputLinesEXT = 5269, SpvExecutionModeOutputLinesNV = 5269, SpvExecutionModeOutputPrimitivesEXT = 5270, @@ -195,6 +196,7 @@ typedef enum SpvExecutionMode_ { SpvExecutionModeSampleInterlockUnorderedEXT = 5369, SpvExecutionModeShadingRateInterlockOrderedEXT = 5370, SpvExecutionModeShadingRateInterlockUnorderedEXT = 5371, + SpvExecutionModeShader64BitIndexingEXT = 5427, SpvExecutionModeSharedLocalMemorySizeINTEL = 5618, SpvExecutionModeRoundingModeRTPINTEL = 5620, SpvExecutionModeRoundingModeRTNINTEL = 5621, @@ -249,8 +251,11 @@ typedef enum SpvStorageClass_ { SpvStorageClassPhysicalStorageBufferEXT = 5349, SpvStorageClassHitObjectAttributeNV = 5385, SpvStorageClassTaskPayloadWorkgroupEXT = 5402, + SpvStorageClassHitObjectAttributeEXT = 5411, SpvStorageClassCodeSectionINTEL = 5605, + SpvStorageClassDeviceOnlyALTERA = 5936, SpvStorageClassDeviceOnlyINTEL = 5936, + SpvStorageClassHostOnlyALTERA = 5937, SpvStorageClassHostOnlyINTEL = 5937, SpvStorageClassMax = 0x7fffffff, } SpvStorageClass; @@ -489,6 +494,7 @@ typedef enum SpvFunctionParameterAttribute_ { SpvFunctionParameterAttributeNoCapture = 5, SpvFunctionParameterAttributeNoWrite = 6, SpvFunctionParameterAttributeNoReadWrite = 7, + SpvFunctionParameterAttributeRuntimeAlignedALTERA = 5940, SpvFunctionParameterAttributeRuntimeAlignedINTEL = 5940, SpvFunctionParameterAttributeMax = 0x7fffffff, } SpvFunctionParameterAttribute; @@ -556,6 +562,9 @@ typedef enum SpvDecoration_ { SpvDecorationPayloadNodeSparseArrayAMDX = 5099, SpvDecorationPayloadNodeArraySizeAMDX = 5100, SpvDecorationPayloadDispatchIndirectAMDX = 5105, + SpvDecorationArrayStrideIdEXT = 5124, + SpvDecorationOffsetIdEXT = 5125, + SpvDecorationUTFEncodedKHR = 5145, SpvDecorationOverrideCoverageNV = 5248, SpvDecorationPassthroughNV = 5250, SpvDecorationViewportRelativeNV = 5252, @@ -572,7 +581,10 @@ typedef enum SpvDecoration_ { SpvDecorationRestrictPointerEXT = 5355, SpvDecorationAliasedPointer = 5356, SpvDecorationAliasedPointerEXT = 5356, + SpvDecorationMemberOffsetNV = 5358, SpvDecorationHitObjectShaderRecordBufferNV = 5386, + SpvDecorationHitObjectShaderRecordBufferEXT = 5389, + SpvDecorationBankNV = 5397, SpvDecorationBindlessSamplerNV = 5398, SpvDecorationBindlessImageNV = 5399, SpvDecorationBoundSamplerNV = 5400, @@ -593,54 +605,95 @@ typedef enum SpvDecoration_ { SpvDecorationUserTypeGOOGLE = 5636, SpvDecorationFunctionRoundingModeINTEL = 5822, SpvDecorationFunctionDenormModeINTEL = 5823, + SpvDecorationRegisterALTERA = 5825, SpvDecorationRegisterINTEL = 5825, + SpvDecorationMemoryALTERA = 5826, SpvDecorationMemoryINTEL = 5826, + SpvDecorationNumbanksALTERA = 5827, SpvDecorationNumbanksINTEL = 5827, + SpvDecorationBankwidthALTERA = 5828, SpvDecorationBankwidthINTEL = 5828, + SpvDecorationMaxPrivateCopiesALTERA = 5829, SpvDecorationMaxPrivateCopiesINTEL = 5829, + SpvDecorationSinglepumpALTERA = 5830, SpvDecorationSinglepumpINTEL = 5830, + SpvDecorationDoublepumpALTERA = 5831, SpvDecorationDoublepumpINTEL = 5831, + SpvDecorationMaxReplicatesALTERA = 5832, SpvDecorationMaxReplicatesINTEL = 5832, + SpvDecorationSimpleDualPortALTERA = 5833, SpvDecorationSimpleDualPortINTEL = 5833, + SpvDecorationMergeALTERA = 5834, SpvDecorationMergeINTEL = 5834, + SpvDecorationBankBitsALTERA = 5835, SpvDecorationBankBitsINTEL = 5835, + SpvDecorationForcePow2DepthALTERA = 5836, SpvDecorationForcePow2DepthINTEL = 5836, + SpvDecorationStridesizeALTERA = 5883, SpvDecorationStridesizeINTEL = 5883, + SpvDecorationWordsizeALTERA = 5884, SpvDecorationWordsizeINTEL = 5884, + SpvDecorationTrueDualPortALTERA = 5885, SpvDecorationTrueDualPortINTEL = 5885, + SpvDecorationBurstCoalesceALTERA = 5899, SpvDecorationBurstCoalesceINTEL = 5899, + SpvDecorationCacheSizeALTERA = 5900, SpvDecorationCacheSizeINTEL = 5900, + SpvDecorationDontStaticallyCoalesceALTERA = 5901, SpvDecorationDontStaticallyCoalesceINTEL = 5901, + SpvDecorationPrefetchALTERA = 5902, SpvDecorationPrefetchINTEL = 5902, + SpvDecorationStallEnableALTERA = 5905, SpvDecorationStallEnableINTEL = 5905, + SpvDecorationFuseLoopsInFunctionALTERA = 5907, SpvDecorationFuseLoopsInFunctionINTEL = 5907, + SpvDecorationMathOpDSPModeALTERA = 5909, SpvDecorationMathOpDSPModeINTEL = 5909, SpvDecorationAliasScopeINTEL = 5914, SpvDecorationNoAliasINTEL = 5915, + SpvDecorationInitiationIntervalALTERA = 5917, SpvDecorationInitiationIntervalINTEL = 5917, + SpvDecorationMaxConcurrencyALTERA = 5918, SpvDecorationMaxConcurrencyINTEL = 5918, + SpvDecorationPipelineEnableALTERA = 5919, SpvDecorationPipelineEnableINTEL = 5919, + SpvDecorationBufferLocationALTERA = 5921, SpvDecorationBufferLocationINTEL = 5921, + SpvDecorationIOPipeStorageALTERA = 5944, SpvDecorationIOPipeStorageINTEL = 5944, SpvDecorationFunctionFloatingPointModeINTEL = 6080, SpvDecorationSingleElementVectorINTEL = 6085, SpvDecorationVectorComputeCallableFunctionINTEL = 6087, SpvDecorationMediaBlockIOINTEL = 6140, + SpvDecorationStallFreeALTERA = 6151, SpvDecorationStallFreeINTEL = 6151, SpvDecorationFPMaxErrorDecorationINTEL = 6170, + SpvDecorationLatencyControlLabelALTERA = 6172, SpvDecorationLatencyControlLabelINTEL = 6172, + SpvDecorationLatencyControlConstraintALTERA = 6173, SpvDecorationLatencyControlConstraintINTEL = 6173, + SpvDecorationConduitKernelArgumentALTERA = 6175, SpvDecorationConduitKernelArgumentINTEL = 6175, + SpvDecorationRegisterMapKernelArgumentALTERA = 6176, SpvDecorationRegisterMapKernelArgumentINTEL = 6176, + SpvDecorationMMHostInterfaceAddressWidthALTERA = 6177, SpvDecorationMMHostInterfaceAddressWidthINTEL = 6177, + SpvDecorationMMHostInterfaceDataWidthALTERA = 6178, SpvDecorationMMHostInterfaceDataWidthINTEL = 6178, + SpvDecorationMMHostInterfaceLatencyALTERA = 6179, SpvDecorationMMHostInterfaceLatencyINTEL = 6179, + SpvDecorationMMHostInterfaceReadWriteModeALTERA = 6180, SpvDecorationMMHostInterfaceReadWriteModeINTEL = 6180, + SpvDecorationMMHostInterfaceMaxBurstALTERA = 6181, SpvDecorationMMHostInterfaceMaxBurstINTEL = 6181, + SpvDecorationMMHostInterfaceWaitRequestALTERA = 6182, SpvDecorationMMHostInterfaceWaitRequestINTEL = 6182, + SpvDecorationStableKernelArgumentALTERA = 6183, SpvDecorationStableKernelArgumentINTEL = 6183, SpvDecorationHostAccessINTEL = 6188, + SpvDecorationInitModeALTERA = 6190, SpvDecorationInitModeINTEL = 6190, + SpvDecorationImplementInRegisterMapALTERA = 6191, SpvDecorationImplementInRegisterMapINTEL = 6191, SpvDecorationConditionalINTEL = 6247, SpvDecorationCacheControlLoadINTEL = 6442, @@ -725,6 +778,8 @@ typedef enum SpvBuiltIn_ { SpvBuiltInFragStencilRefEXT = 5014, SpvBuiltInRemainingRecursionLevelsAMDX = 5021, SpvBuiltInShaderIndexAMDX = 5073, + SpvBuiltInSamplerHeapEXT = 5122, + SpvBuiltInResourceHeapEXT = 5123, SpvBuiltInViewportMaskNV = 5253, SpvBuiltInSecondaryPositionNV = 5257, SpvBuiltInSecondaryViewportMaskNV = 5258, @@ -822,15 +877,25 @@ typedef enum SpvLoopControlShift_ { SpvLoopControlIterationMultipleShift = 6, SpvLoopControlPeelCountShift = 7, SpvLoopControlPartialCountShift = 8, + SpvLoopControlInitiationIntervalALTERAShift = 16, SpvLoopControlInitiationIntervalINTELShift = 16, + SpvLoopControlMaxConcurrencyALTERAShift = 17, SpvLoopControlMaxConcurrencyINTELShift = 17, + SpvLoopControlDependencyArrayALTERAShift = 18, SpvLoopControlDependencyArrayINTELShift = 18, + SpvLoopControlPipelineEnableALTERAShift = 19, SpvLoopControlPipelineEnableINTELShift = 19, + SpvLoopControlLoopCoalesceALTERAShift = 20, SpvLoopControlLoopCoalesceINTELShift = 20, + SpvLoopControlMaxInterleavingALTERAShift = 21, SpvLoopControlMaxInterleavingINTELShift = 21, + SpvLoopControlSpeculatedIterationsALTERAShift = 22, SpvLoopControlSpeculatedIterationsINTELShift = 22, + SpvLoopControlNoFusionALTERAShift = 23, SpvLoopControlNoFusionINTELShift = 23, + SpvLoopControlLoopCountALTERAShift = 24, SpvLoopControlLoopCountINTELShift = 24, + SpvLoopControlMaxReinvocationDelayALTERAShift = 25, SpvLoopControlMaxReinvocationDelayINTELShift = 25, SpvLoopControlMax = 0x7fffffff, } SpvLoopControlShift; @@ -846,15 +911,25 @@ typedef enum SpvLoopControlMask_ { SpvLoopControlIterationMultipleMask = 0x00000040, SpvLoopControlPeelCountMask = 0x00000080, SpvLoopControlPartialCountMask = 0x00000100, + SpvLoopControlInitiationIntervalALTERAMask = 0x00010000, SpvLoopControlInitiationIntervalINTELMask = 0x00010000, + SpvLoopControlMaxConcurrencyALTERAMask = 0x00020000, SpvLoopControlMaxConcurrencyINTELMask = 0x00020000, + SpvLoopControlDependencyArrayALTERAMask = 0x00040000, SpvLoopControlDependencyArrayINTELMask = 0x00040000, + SpvLoopControlPipelineEnableALTERAMask = 0x00080000, SpvLoopControlPipelineEnableINTELMask = 0x00080000, + SpvLoopControlLoopCoalesceALTERAMask = 0x00100000, SpvLoopControlLoopCoalesceINTELMask = 0x00100000, + SpvLoopControlMaxInterleavingALTERAMask = 0x00200000, SpvLoopControlMaxInterleavingINTELMask = 0x00200000, + SpvLoopControlSpeculatedIterationsALTERAMask = 0x00400000, SpvLoopControlSpeculatedIterationsINTELMask = 0x00400000, + SpvLoopControlNoFusionALTERAMask = 0x00800000, SpvLoopControlNoFusionINTELMask = 0x00800000, + SpvLoopControlLoopCountALTERAMask = 0x01000000, SpvLoopControlLoopCountINTELMask = 0x01000000, + SpvLoopControlMaxReinvocationDelayALTERAMask = 0x02000000, SpvLoopControlMaxReinvocationDelayINTELMask = 0x02000000, } SpvLoopControlMask; @@ -967,8 +1042,11 @@ typedef enum SpvGroupOperation_ { SpvGroupOperationInclusiveScan = 1, SpvGroupOperationExclusiveScan = 2, SpvGroupOperationClusteredReduce = 3, + SpvGroupOperationPartitionedReduceEXT = 6, SpvGroupOperationPartitionedReduceNV = 6, + SpvGroupOperationPartitionedInclusiveScanEXT = 7, SpvGroupOperationPartitionedInclusiveScanNV = 7, + SpvGroupOperationPartitionedExclusiveScanEXT = 8, SpvGroupOperationPartitionedExclusiveScanNV = 8, SpvGroupOperationMax = 0x7fffffff, } SpvGroupOperation; @@ -1124,6 +1202,10 @@ typedef enum SpvCapability_ { SpvCapabilityBFloat16TypeKHR = 5116, SpvCapabilityBFloat16DotProductKHR = 5117, SpvCapabilityBFloat16CooperativeMatrixKHR = 5118, + SpvCapabilityAbortKHR = 5120, + SpvCapabilityDescriptorHeapEXT = 5128, + SpvCapabilityConstantDataKHR = 5146, + SpvCapabilityPoisonFreezeKHR = 5156, SpvCapabilitySampleMaskOverrideCoverageNV = 5249, SpvCapabilityGeometryShaderPassthroughNV = 5251, SpvCapabilityShaderViewportIndexLayerEXT = 5254, @@ -1141,6 +1223,7 @@ typedef enum SpvCapability_ { SpvCapabilityComputeDerivativeGroupQuadsNV = 5288, SpvCapabilityFragmentDensityEXT = 5291, SpvCapabilityShadingRateNV = 5291, + SpvCapabilityGroupNonUniformPartitionedEXT = 5297, SpvCapabilityGroupNonUniformPartitionedNV = 5297, SpvCapabilityShaderNonUniform = 5301, SpvCapabilityShaderNonUniformEXT = 5301, @@ -1188,6 +1271,7 @@ typedef enum SpvCapability_ { SpvCapabilityDisplacementMicromapNV = 5380, SpvCapabilityRayTracingOpacityMicromapEXT = 5381, SpvCapabilityShaderInvocationReorderNV = 5383, + SpvCapabilityShaderInvocationReorderEXT = 5388, SpvCapabilityBindlessTextureNV = 5390, SpvCapabilityRayQueryPositionFetchKHR = 5391, SpvCapabilityCooperativeVectorNV = 5394, @@ -1196,6 +1280,9 @@ typedef enum SpvCapability_ { SpvCapabilityRawAccessChainsNV = 5414, SpvCapabilityRayTracingSpheresGeometryNV = 5418, SpvCapabilityRayTracingLinearSweptSpheresGeometryNV = 5419, + SpvCapabilityPushConstantBanksNV = 5423, + SpvCapabilityLongVectorEXT = 5425, + SpvCapabilityShader64BitIndexingEXT = 5426, SpvCapabilityCooperativeMatrixReductionsNV = 5430, SpvCapabilityCooperativeMatrixConversionsNV = 5431, SpvCapabilityCooperativeMatrixPerElementOperationsNV = 5432, @@ -1225,26 +1312,42 @@ typedef enum SpvCapability_ { SpvCapabilitySubgroupAvcMotionEstimationChromaINTEL = 5698, SpvCapabilityVariableLengthArrayINTEL = 5817, SpvCapabilityFunctionFloatControlINTEL = 5821, + SpvCapabilityFPGAMemoryAttributesALTERA = 5824, SpvCapabilityFPGAMemoryAttributesINTEL = 5824, SpvCapabilityFPFastMathModeINTEL = 5837, + SpvCapabilityArbitraryPrecisionIntegersALTERA = 5844, SpvCapabilityArbitraryPrecisionIntegersINTEL = 5844, + SpvCapabilityArbitraryPrecisionFloatingPointALTERA = 5845, SpvCapabilityArbitraryPrecisionFloatingPointINTEL = 5845, SpvCapabilityUnstructuredLoopControlsINTEL = 5886, + SpvCapabilityFPGALoopControlsALTERA = 5888, SpvCapabilityFPGALoopControlsINTEL = 5888, SpvCapabilityKernelAttributesINTEL = 5892, SpvCapabilityFPGAKernelAttributesINTEL = 5897, + SpvCapabilityFPGAMemoryAccessesALTERA = 5898, SpvCapabilityFPGAMemoryAccessesINTEL = 5898, + SpvCapabilityFPGAClusterAttributesALTERA = 5904, SpvCapabilityFPGAClusterAttributesINTEL = 5904, + SpvCapabilityLoopFuseALTERA = 5906, SpvCapabilityLoopFuseINTEL = 5906, + SpvCapabilityFPGADSPControlALTERA = 5908, SpvCapabilityFPGADSPControlINTEL = 5908, SpvCapabilityMemoryAccessAliasingINTEL = 5910, + SpvCapabilityFPGAInvocationPipeliningAttributesALTERA = 5916, SpvCapabilityFPGAInvocationPipeliningAttributesINTEL = 5916, + SpvCapabilityFPGABufferLocationALTERA = 5920, SpvCapabilityFPGABufferLocationINTEL = 5920, + SpvCapabilityArbitraryPrecisionFixedPointALTERA = 5922, SpvCapabilityArbitraryPrecisionFixedPointINTEL = 5922, + SpvCapabilityUSMStorageClassesALTERA = 5935, SpvCapabilityUSMStorageClassesINTEL = 5935, + SpvCapabilityRuntimeAlignedAttributeALTERA = 5939, SpvCapabilityRuntimeAlignedAttributeINTEL = 5939, + SpvCapabilityIOPipesALTERA = 5943, SpvCapabilityIOPipesINTEL = 5943, + SpvCapabilityBlockingPipesALTERA = 5945, SpvCapabilityBlockingPipesINTEL = 5945, + SpvCapabilityFPGARegALTERA = 5948, SpvCapabilityFPGARegINTEL = 5948, SpvCapabilityDotProductInputAll = 6016, SpvCapabilityDotProductInputAllKHR = 6016, @@ -1271,13 +1374,18 @@ typedef enum SpvCapability_ { SpvCapabilityBFloat16ConversionINTEL = 6115, SpvCapabilitySplitBarrierINTEL = 6141, SpvCapabilityArithmeticFenceEXT = 6144, + SpvCapabilityFPGAClusterAttributesV2ALTERA = 6150, SpvCapabilityFPGAClusterAttributesV2INTEL = 6150, SpvCapabilityFPGAKernelAttributesv2INTEL = 6161, + SpvCapabilityTaskSequenceALTERA = 6162, SpvCapabilityTaskSequenceINTEL = 6162, SpvCapabilityFPMaxErrorINTEL = 6169, + SpvCapabilityFPGALatencyControlALTERA = 6171, SpvCapabilityFPGALatencyControlINTEL = 6171, + SpvCapabilityFPGAArgumentInterfacesALTERA = 6174, SpvCapabilityFPGAArgumentInterfacesINTEL = 6174, SpvCapabilityGlobalVariableHostAccessINTEL = 6187, + SpvCapabilityGlobalVariableFPGADecorationsALTERA = 6189, SpvCapabilityGlobalVariableFPGADecorationsINTEL = 6189, SpvCapabilitySubgroupBufferPrefetchINTEL = 6220, SpvCapabilitySubgroup2DBlockIOINTEL = 6228, @@ -1294,6 +1402,10 @@ typedef enum SpvCapability_ { SpvCapabilityCacheControlsINTEL = 6441, SpvCapabilityRegisterLimitsINTEL = 6460, SpvCapabilityBindlessImagesINTEL = 6528, + SpvCapabilityDotProductFloat16AccFloat32VALVE = 6912, + SpvCapabilityDotProductFloat16AccFloat16VALVE = 6913, + SpvCapabilityDotProductBFloat16AccVALVE = 6914, + SpvCapabilityDotProductFloat8AccFloat32VALVE = 6915, SpvCapabilityMax = 0x7fffffff, } SpvCapability; @@ -1489,7 +1601,9 @@ typedef enum SpvTensorOperandsMask_ { } SpvTensorOperandsMask; typedef enum SpvInitializationModeQualifier_ { + SpvInitializationModeQualifierInitOnDeviceReprogramALTERA = 0, SpvInitializationModeQualifierInitOnDeviceReprogramINTEL = 0, + SpvInitializationModeQualifierInitOnDeviceResetALTERA = 1, SpvInitializationModeQualifierInitOnDeviceResetINTEL = 1, SpvInitializationModeQualifierMax = 0x7fffffff, } SpvInitializationModeQualifier; @@ -2049,6 +2163,16 @@ typedef enum SpvOp_ { SpvOpSpecConstantStringAMDX = 5104, SpvOpGroupNonUniformQuadAllKHR = 5110, SpvOpGroupNonUniformQuadAnyKHR = 5111, + SpvOpTypeBufferEXT = 5115, + SpvOpBufferPointerEXT = 5119, + SpvOpAbortKHR = 5121, + SpvOpUntypedImageTexelPointerEXT = 5126, + SpvOpMemberDecorateIdEXT = 5127, + SpvOpConstantSizeOfEXT = 5129, + SpvOpConstantDataKHR = 5147, + SpvOpSpecConstantDataKHR = 5148, + SpvOpPoisonKHR = 5158, + SpvOpFreezeKHR = 5159, SpvOpHitObjectRecordHitMotionNV = 5249, SpvOpHitObjectRecordHitWithIndexMotionNV = 5250, SpvOpHitObjectRecordMissMotionNV = 5251, @@ -2084,6 +2208,7 @@ typedef enum SpvOp_ { SpvOpTypeHitObjectNV = 5281, SpvOpImageSampleFootprintNV = 5283, SpvOpTypeCooperativeVectorNV = 5288, + SpvOpTypeVectorIdEXT = 5288, SpvOpCooperativeVectorMatrixMulNV = 5289, SpvOpCooperativeVectorOuterProductAccumulateNV = 5290, SpvOpCooperativeVectorReduceSumAccumulateNV = 5291, @@ -2091,12 +2216,43 @@ typedef enum SpvOp_ { SpvOpCooperativeMatrixConvertNV = 5293, SpvOpEmitMeshTasksEXT = 5294, SpvOpSetMeshOutputsEXT = 5295, + SpvOpGroupNonUniformPartitionEXT = 5296, SpvOpGroupNonUniformPartitionNV = 5296, SpvOpWritePackedPrimitiveIndices4x8NV = 5299, SpvOpFetchMicroTriangleVertexPositionNV = 5300, SpvOpFetchMicroTriangleVertexBarycentricNV = 5301, SpvOpCooperativeVectorLoadNV = 5302, SpvOpCooperativeVectorStoreNV = 5303, + SpvOpHitObjectRecordFromQueryEXT = 5304, + SpvOpHitObjectRecordMissEXT = 5305, + SpvOpHitObjectRecordMissMotionEXT = 5306, + SpvOpHitObjectGetIntersectionTriangleVertexPositionsEXT = 5307, + SpvOpHitObjectGetRayFlagsEXT = 5308, + SpvOpHitObjectSetShaderBindingTableRecordIndexEXT = 5309, + SpvOpHitObjectReorderExecuteShaderEXT = 5310, + SpvOpHitObjectTraceReorderExecuteEXT = 5311, + SpvOpHitObjectTraceMotionReorderExecuteEXT = 5312, + SpvOpTypeHitObjectEXT = 5313, + SpvOpReorderThreadWithHintEXT = 5314, + SpvOpReorderThreadWithHitObjectEXT = 5315, + SpvOpHitObjectTraceRayEXT = 5316, + SpvOpHitObjectTraceRayMotionEXT = 5317, + SpvOpHitObjectRecordEmptyEXT = 5318, + SpvOpHitObjectExecuteShaderEXT = 5319, + SpvOpHitObjectGetCurrentTimeEXT = 5320, + SpvOpHitObjectGetAttributesEXT = 5321, + SpvOpHitObjectGetHitKindEXT = 5322, + SpvOpHitObjectGetPrimitiveIndexEXT = 5323, + SpvOpHitObjectGetGeometryIndexEXT = 5324, + SpvOpHitObjectGetInstanceIdEXT = 5325, + SpvOpHitObjectGetInstanceCustomIndexEXT = 5326, + SpvOpHitObjectGetObjectRayOriginEXT = 5327, + SpvOpHitObjectGetObjectRayDirectionEXT = 5328, + SpvOpHitObjectGetWorldRayDirectionEXT = 5329, + SpvOpHitObjectGetWorldRayOriginEXT = 5330, + SpvOpHitObjectGetObjectToWorldEXT = 5331, + SpvOpHitObjectGetWorldToObjectEXT = 5332, + SpvOpHitObjectGetRayTMaxEXT = 5333, SpvOpReportIntersectionKHR = 5334, SpvOpReportIntersectionNV = 5334, SpvOpIgnoreIntersectionNV = 5335, @@ -2111,6 +2267,12 @@ typedef enum SpvOp_ { SpvOpRayQueryGetClusterIdNV = 5345, SpvOpRayQueryGetIntersectionClusterIdNV = 5345, SpvOpHitObjectGetClusterIdNV = 5346, + SpvOpHitObjectGetRayTMinEXT = 5347, + SpvOpHitObjectGetShaderBindingTableRecordIndexEXT = 5348, + SpvOpHitObjectGetShaderRecordBufferHandleEXT = 5349, + SpvOpHitObjectIsEmptyEXT = 5350, + SpvOpHitObjectIsHitEXT = 5351, + SpvOpHitObjectIsMissEXT = 5352, SpvOpTypeCooperativeMatrixNV = 5358, SpvOpCooperativeMatrixLoadNV = 5359, SpvOpCooperativeMatrixStoreNV = 5360, @@ -2317,23 +2479,41 @@ typedef enum SpvOp_ { SpvOpVariableLengthArrayINTEL = 5818, SpvOpSaveMemoryINTEL = 5819, SpvOpRestoreMemoryINTEL = 5820, + SpvOpArbitraryFloatSinCosPiALTERA = 5840, SpvOpArbitraryFloatSinCosPiINTEL = 5840, + SpvOpArbitraryFloatCastALTERA = 5841, SpvOpArbitraryFloatCastINTEL = 5841, + SpvOpArbitraryFloatCastFromIntALTERA = 5842, SpvOpArbitraryFloatCastFromIntINTEL = 5842, + SpvOpArbitraryFloatCastToIntALTERA = 5843, SpvOpArbitraryFloatCastToIntINTEL = 5843, + SpvOpArbitraryFloatAddALTERA = 5846, SpvOpArbitraryFloatAddINTEL = 5846, + SpvOpArbitraryFloatSubALTERA = 5847, SpvOpArbitraryFloatSubINTEL = 5847, + SpvOpArbitraryFloatMulALTERA = 5848, SpvOpArbitraryFloatMulINTEL = 5848, + SpvOpArbitraryFloatDivALTERA = 5849, SpvOpArbitraryFloatDivINTEL = 5849, + SpvOpArbitraryFloatGTALTERA = 5850, SpvOpArbitraryFloatGTINTEL = 5850, + SpvOpArbitraryFloatGEALTERA = 5851, SpvOpArbitraryFloatGEINTEL = 5851, + SpvOpArbitraryFloatLTALTERA = 5852, SpvOpArbitraryFloatLTINTEL = 5852, + SpvOpArbitraryFloatLEALTERA = 5853, SpvOpArbitraryFloatLEINTEL = 5853, + SpvOpArbitraryFloatEQALTERA = 5854, SpvOpArbitraryFloatEQINTEL = 5854, + SpvOpArbitraryFloatRecipALTERA = 5855, SpvOpArbitraryFloatRecipINTEL = 5855, + SpvOpArbitraryFloatRSqrtALTERA = 5856, SpvOpArbitraryFloatRSqrtINTEL = 5856, + SpvOpArbitraryFloatCbrtALTERA = 5857, SpvOpArbitraryFloatCbrtINTEL = 5857, + SpvOpArbitraryFloatHypotALTERA = 5858, SpvOpArbitraryFloatHypotINTEL = 5858, + SpvOpArbitraryFloatSqrtALTERA = 5859, SpvOpArbitraryFloatSqrtINTEL = 5859, SpvOpArbitraryFloatLogINTEL = 5860, SpvOpArbitraryFloatLog2INTEL = 5861, @@ -2362,21 +2542,37 @@ typedef enum SpvOp_ { SpvOpAliasDomainDeclINTEL = 5911, SpvOpAliasScopeDeclINTEL = 5912, SpvOpAliasScopeListDeclINTEL = 5913, + SpvOpFixedSqrtALTERA = 5923, SpvOpFixedSqrtINTEL = 5923, + SpvOpFixedRecipALTERA = 5924, SpvOpFixedRecipINTEL = 5924, + SpvOpFixedRsqrtALTERA = 5925, SpvOpFixedRsqrtINTEL = 5925, + SpvOpFixedSinALTERA = 5926, SpvOpFixedSinINTEL = 5926, + SpvOpFixedCosALTERA = 5927, SpvOpFixedCosINTEL = 5927, + SpvOpFixedSinCosALTERA = 5928, SpvOpFixedSinCosINTEL = 5928, + SpvOpFixedSinPiALTERA = 5929, SpvOpFixedSinPiINTEL = 5929, + SpvOpFixedCosPiALTERA = 5930, SpvOpFixedCosPiINTEL = 5930, + SpvOpFixedSinCosPiALTERA = 5931, SpvOpFixedSinCosPiINTEL = 5931, + SpvOpFixedLogALTERA = 5932, SpvOpFixedLogINTEL = 5932, + SpvOpFixedExpALTERA = 5933, SpvOpFixedExpINTEL = 5933, + SpvOpPtrCastToCrossWorkgroupALTERA = 5934, SpvOpPtrCastToCrossWorkgroupINTEL = 5934, + SpvOpCrossWorkgroupCastToPtrALTERA = 5938, SpvOpCrossWorkgroupCastToPtrINTEL = 5938, + SpvOpReadPipeBlockingALTERA = 5946, SpvOpReadPipeBlockingINTEL = 5946, + SpvOpWritePipeBlockingALTERA = 5947, SpvOpWritePipeBlockingINTEL = 5947, + SpvOpFPGARegALTERA = 5949, SpvOpFPGARegINTEL = 5949, SpvOpRayQueryGetRayTMinKHR = 6016, SpvOpRayQueryGetRayFlagsKHR = 6017, @@ -2406,10 +2602,15 @@ typedef enum SpvOp_ { SpvOpControlBarrierArriveINTEL = 6142, SpvOpControlBarrierWaitINTEL = 6143, SpvOpArithmeticFenceEXT = 6145, + SpvOpTaskSequenceCreateALTERA = 6163, SpvOpTaskSequenceCreateINTEL = 6163, + SpvOpTaskSequenceAsyncALTERA = 6164, SpvOpTaskSequenceAsyncINTEL = 6164, + SpvOpTaskSequenceGetALTERA = 6165, SpvOpTaskSequenceGetINTEL = 6165, + SpvOpTaskSequenceReleaseALTERA = 6166, SpvOpTaskSequenceReleaseINTEL = 6166, + SpvOpTypeTaskSequenceALTERA = 6199, SpvOpTypeTaskSequenceINTEL = 6199, SpvOpSubgroupBlockPrefetchINTEL = 6221, SpvOpSubgroup2DBlockLoadINTEL = 6231, @@ -2441,6 +2642,9 @@ typedef enum SpvOp_ { SpvOpConvertHandleToImageINTEL = 6529, SpvOpConvertHandleToSamplerINTEL = 6530, SpvOpConvertHandleToSampledImageINTEL = 6531, + SpvOpFDot2MixAcc32VALVE = 6916, + SpvOpFDot2MixAcc16VALVE = 6917, + SpvOpFDot4MixAcc32VALVE = 6918, SpvOpMax = 0x7fffffff, } SpvOp; @@ -2888,6 +3092,16 @@ inline void SpvHasResultAndType(SpvOp opcode, bool *hasResult, bool *hasResultTy case SpvOpSpecConstantStringAMDX: *hasResult = true; *hasResultType = false; break; case SpvOpGroupNonUniformQuadAllKHR: *hasResult = true; *hasResultType = true; break; case SpvOpGroupNonUniformQuadAnyKHR: *hasResult = true; *hasResultType = true; break; + case SpvOpTypeBufferEXT: *hasResult = true; *hasResultType = false; break; + case SpvOpBufferPointerEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpAbortKHR: *hasResult = false; *hasResultType = false; break; + case SpvOpUntypedImageTexelPointerEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpMemberDecorateIdEXT: *hasResult = false; *hasResultType = false; break; + case SpvOpConstantSizeOfEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpConstantDataKHR: *hasResult = true; *hasResultType = true; break; + case SpvOpSpecConstantDataKHR: *hasResult = true; *hasResultType = true; break; + case SpvOpPoisonKHR: *hasResult = true; *hasResultType = true; break; + case SpvOpFreezeKHR: *hasResult = true; *hasResultType = true; break; case SpvOpHitObjectRecordHitMotionNV: *hasResult = false; *hasResultType = false; break; case SpvOpHitObjectRecordHitWithIndexMotionNV: *hasResult = false; *hasResultType = false; break; case SpvOpHitObjectRecordMissMotionNV: *hasResult = false; *hasResultType = false; break; @@ -2922,7 +3136,7 @@ inline void SpvHasResultAndType(SpvOp opcode, bool *hasResult, bool *hasResultTy case SpvOpReorderThreadWithHintNV: *hasResult = false; *hasResultType = false; break; case SpvOpTypeHitObjectNV: *hasResult = true; *hasResultType = false; break; case SpvOpImageSampleFootprintNV: *hasResult = true; *hasResultType = true; break; - case SpvOpTypeCooperativeVectorNV: *hasResult = true; *hasResultType = false; break; + case SpvOpTypeVectorIdEXT: *hasResult = true; *hasResultType = false; break; case SpvOpCooperativeVectorMatrixMulNV: *hasResult = true; *hasResultType = true; break; case SpvOpCooperativeVectorOuterProductAccumulateNV: *hasResult = false; *hasResultType = false; break; case SpvOpCooperativeVectorReduceSumAccumulateNV: *hasResult = false; *hasResultType = false; break; @@ -2930,12 +3144,42 @@ inline void SpvHasResultAndType(SpvOp opcode, bool *hasResult, bool *hasResultTy case SpvOpCooperativeMatrixConvertNV: *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 SpvOpGroupNonUniformPartitionEXT: *hasResult = true; *hasResultType = true; break; case SpvOpWritePackedPrimitiveIndices4x8NV: *hasResult = false; *hasResultType = false; break; case SpvOpFetchMicroTriangleVertexPositionNV: *hasResult = true; *hasResultType = true; break; case SpvOpFetchMicroTriangleVertexBarycentricNV: *hasResult = true; *hasResultType = true; break; case SpvOpCooperativeVectorLoadNV: *hasResult = true; *hasResultType = true; break; case SpvOpCooperativeVectorStoreNV: *hasResult = false; *hasResultType = false; break; + case SpvOpHitObjectRecordFromQueryEXT: *hasResult = false; *hasResultType = false; break; + case SpvOpHitObjectRecordMissEXT: *hasResult = false; *hasResultType = false; break; + case SpvOpHitObjectRecordMissMotionEXT: *hasResult = false; *hasResultType = false; break; + case SpvOpHitObjectGetIntersectionTriangleVertexPositionsEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectGetRayFlagsEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectSetShaderBindingTableRecordIndexEXT: *hasResult = false; *hasResultType = false; break; + case SpvOpHitObjectReorderExecuteShaderEXT: *hasResult = false; *hasResultType = false; break; + case SpvOpHitObjectTraceReorderExecuteEXT: *hasResult = false; *hasResultType = false; break; + case SpvOpHitObjectTraceMotionReorderExecuteEXT: *hasResult = false; *hasResultType = false; break; + case SpvOpTypeHitObjectEXT: *hasResult = true; *hasResultType = false; break; + case SpvOpReorderThreadWithHintEXT: *hasResult = false; *hasResultType = false; break; + case SpvOpReorderThreadWithHitObjectEXT: *hasResult = false; *hasResultType = false; break; + case SpvOpHitObjectTraceRayEXT: *hasResult = false; *hasResultType = false; break; + case SpvOpHitObjectTraceRayMotionEXT: *hasResult = false; *hasResultType = false; break; + case SpvOpHitObjectRecordEmptyEXT: *hasResult = false; *hasResultType = false; break; + case SpvOpHitObjectExecuteShaderEXT: *hasResult = false; *hasResultType = false; break; + case SpvOpHitObjectGetCurrentTimeEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectGetAttributesEXT: *hasResult = false; *hasResultType = false; break; + case SpvOpHitObjectGetHitKindEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectGetPrimitiveIndexEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectGetGeometryIndexEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectGetInstanceIdEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectGetInstanceCustomIndexEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectGetObjectRayOriginEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectGetObjectRayDirectionEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectGetWorldRayDirectionEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectGetWorldRayOriginEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectGetObjectToWorldEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectGetWorldToObjectEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectGetRayTMaxEXT: *hasResult = true; *hasResultType = true; break; case SpvOpReportIntersectionKHR: *hasResult = true; *hasResultType = true; break; case SpvOpIgnoreIntersectionNV: *hasResult = false; *hasResultType = false; break; case SpvOpTerminateRayNV: *hasResult = false; *hasResultType = false; break; @@ -2947,6 +3191,12 @@ inline void SpvHasResultAndType(SpvOp opcode, bool *hasResult, bool *hasResultTy case SpvOpExecuteCallableNV: *hasResult = false; *hasResultType = false; break; case SpvOpRayQueryGetIntersectionClusterIdNV: *hasResult = true; *hasResultType = true; break; case SpvOpHitObjectGetClusterIdNV: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectGetRayTMinEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectGetShaderBindingTableRecordIndexEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectGetShaderRecordBufferHandleEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectIsEmptyEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectIsHitEXT: *hasResult = true; *hasResultType = true; break; + case SpvOpHitObjectIsMissEXT: *hasResult = true; *hasResultType = true; break; case SpvOpTypeCooperativeMatrixNV: *hasResult = true; *hasResultType = false; break; case SpvOpCooperativeMatrixLoadNV: *hasResult = true; *hasResultType = true; break; case SpvOpCooperativeMatrixStoreNV: *hasResult = false; *hasResultType = false; break; @@ -3150,24 +3400,24 @@ inline void SpvHasResultAndType(SpvOp opcode, bool *hasResult, bool *hasResultTy 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 SpvOpArbitraryFloatSinCosPiALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpArbitraryFloatCastALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpArbitraryFloatCastFromIntALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpArbitraryFloatCastToIntALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpArbitraryFloatAddALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpArbitraryFloatSubALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpArbitraryFloatMulALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpArbitraryFloatDivALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpArbitraryFloatGTALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpArbitraryFloatGEALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpArbitraryFloatLTALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpArbitraryFloatLEALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpArbitraryFloatEQALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpArbitraryFloatRecipALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpArbitraryFloatRSqrtALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpArbitraryFloatCbrtALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpArbitraryFloatHypotALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpArbitraryFloatSqrtALTERA: *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; @@ -3195,22 +3445,22 @@ inline void SpvHasResultAndType(SpvOp opcode, bool *hasResult, bool *hasResultTy 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; + case SpvOpFixedSqrtALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpFixedRecipALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpFixedRsqrtALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpFixedSinALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpFixedCosALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpFixedSinCosALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpFixedSinPiALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpFixedCosPiALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpFixedSinCosPiALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpFixedLogALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpFixedExpALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpPtrCastToCrossWorkgroupALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpCrossWorkgroupCastToPtrALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpReadPipeBlockingALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpWritePipeBlockingALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpFPGARegALTERA: *hasResult = true; *hasResultType = true; break; case SpvOpRayQueryGetRayTMinKHR: *hasResult = true; *hasResultType = true; break; case SpvOpRayQueryGetRayFlagsKHR: *hasResult = true; *hasResultType = true; break; case SpvOpRayQueryGetIntersectionTKHR: *hasResult = true; *hasResultType = true; break; @@ -3239,11 +3489,11 @@ inline void SpvHasResultAndType(SpvOp opcode, bool *hasResult, bool *hasResultTy case SpvOpControlBarrierArriveINTEL: *hasResult = false; *hasResultType = false; break; case SpvOpControlBarrierWaitINTEL: *hasResult = false; *hasResultType = false; break; case SpvOpArithmeticFenceEXT: *hasResult = true; *hasResultType = true; break; - case SpvOpTaskSequenceCreateINTEL: *hasResult = true; *hasResultType = true; break; - case SpvOpTaskSequenceAsyncINTEL: *hasResult = false; *hasResultType = false; break; - case SpvOpTaskSequenceGetINTEL: *hasResult = true; *hasResultType = true; break; - case SpvOpTaskSequenceReleaseINTEL: *hasResult = false; *hasResultType = false; break; - case SpvOpTypeTaskSequenceINTEL: *hasResult = true; *hasResultType = false; break; + case SpvOpTaskSequenceCreateALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpTaskSequenceAsyncALTERA: *hasResult = false; *hasResultType = false; break; + case SpvOpTaskSequenceGetALTERA: *hasResult = true; *hasResultType = true; break; + case SpvOpTaskSequenceReleaseALTERA: *hasResult = false; *hasResultType = false; break; + case SpvOpTypeTaskSequenceALTERA: *hasResult = true; *hasResultType = false; break; case SpvOpSubgroupBlockPrefetchINTEL: *hasResult = false; *hasResultType = false; break; case SpvOpSubgroup2DBlockLoadINTEL: *hasResult = false; *hasResultType = false; break; case SpvOpSubgroup2DBlockLoadTransformINTEL: *hasResult = false; *hasResultType = false; break; @@ -3274,6 +3524,9 @@ inline void SpvHasResultAndType(SpvOp opcode, bool *hasResult, bool *hasResultTy case SpvOpConvertHandleToImageINTEL: *hasResult = true; *hasResultType = true; break; case SpvOpConvertHandleToSamplerINTEL: *hasResult = true; *hasResultType = true; break; case SpvOpConvertHandleToSampledImageINTEL: *hasResult = true; *hasResultType = true; break; + case SpvOpFDot2MixAcc32VALVE: *hasResult = true; *hasResultType = true; break; + case SpvOpFDot2MixAcc16VALVE: *hasResult = true; *hasResultType = true; break; + case SpvOpFDot4MixAcc32VALVE: *hasResult = true; *hasResultType = true; break; } } inline const char* SpvSourceLanguageToString(SpvSourceLanguage value) { @@ -3408,6 +3661,7 @@ inline const char* SpvExecutionModeToString(SpvExecutionMode value) { case SpvExecutionModeQuadDerivativesKHR: return "QuadDerivativesKHR"; case SpvExecutionModeRequireFullQuadsKHR: return "RequireFullQuadsKHR"; case SpvExecutionModeSharesInputWithAMDX: return "SharesInputWithAMDX"; + case SpvExecutionModeArithmeticPoisonKHR: return "ArithmeticPoisonKHR"; case SpvExecutionModeOutputLinesEXT: return "OutputLinesEXT"; case SpvExecutionModeOutputPrimitivesEXT: return "OutputPrimitivesEXT"; case SpvExecutionModeDerivativeGroupQuadsKHR: return "DerivativeGroupQuadsKHR"; @@ -3419,6 +3673,7 @@ inline const char* SpvExecutionModeToString(SpvExecutionMode value) { case SpvExecutionModeSampleInterlockUnorderedEXT: return "SampleInterlockUnorderedEXT"; case SpvExecutionModeShadingRateInterlockOrderedEXT: return "ShadingRateInterlockOrderedEXT"; case SpvExecutionModeShadingRateInterlockUnorderedEXT: return "ShadingRateInterlockUnorderedEXT"; + case SpvExecutionModeShader64BitIndexingEXT: return "Shader64BitIndexingEXT"; case SpvExecutionModeSharedLocalMemorySizeINTEL: return "SharedLocalMemorySizeINTEL"; case SpvExecutionModeRoundingModeRTPINTEL: return "RoundingModeRTPINTEL"; case SpvExecutionModeRoundingModeRTNINTEL: return "RoundingModeRTNINTEL"; @@ -3468,9 +3723,10 @@ inline const char* SpvStorageClassToString(SpvStorageClass value) { case SpvStorageClassPhysicalStorageBuffer: return "PhysicalStorageBuffer"; case SpvStorageClassHitObjectAttributeNV: return "HitObjectAttributeNV"; case SpvStorageClassTaskPayloadWorkgroupEXT: return "TaskPayloadWorkgroupEXT"; + case SpvStorageClassHitObjectAttributeEXT: return "HitObjectAttributeEXT"; case SpvStorageClassCodeSectionINTEL: return "CodeSectionINTEL"; - case SpvStorageClassDeviceOnlyINTEL: return "DeviceOnlyINTEL"; - case SpvStorageClassHostOnlyINTEL: return "HostOnlyINTEL"; + case SpvStorageClassDeviceOnlyALTERA: return "DeviceOnlyALTERA"; + case SpvStorageClassHostOnlyALTERA: return "HostOnlyALTERA"; default: return "Unknown"; } } @@ -3652,7 +3908,7 @@ inline const char* SpvFunctionParameterAttributeToString(SpvFunctionParameterAtt case SpvFunctionParameterAttributeNoCapture: return "NoCapture"; case SpvFunctionParameterAttributeNoWrite: return "NoWrite"; case SpvFunctionParameterAttributeNoReadWrite: return "NoReadWrite"; - case SpvFunctionParameterAttributeRuntimeAlignedINTEL: return "RuntimeAlignedINTEL"; + case SpvFunctionParameterAttributeRuntimeAlignedALTERA: return "RuntimeAlignedALTERA"; default: return "Unknown"; } } @@ -3721,6 +3977,9 @@ inline const char* SpvDecorationToString(SpvDecoration value) { case SpvDecorationPayloadNodeSparseArrayAMDX: return "PayloadNodeSparseArrayAMDX"; case SpvDecorationPayloadNodeArraySizeAMDX: return "PayloadNodeArraySizeAMDX"; case SpvDecorationPayloadDispatchIndirectAMDX: return "PayloadDispatchIndirectAMDX"; + case SpvDecorationArrayStrideIdEXT: return "ArrayStrideIdEXT"; + case SpvDecorationOffsetIdEXT: return "OffsetIdEXT"; + case SpvDecorationUTFEncodedKHR: return "UTFEncodedKHR"; case SpvDecorationOverrideCoverageNV: return "OverrideCoverageNV"; case SpvDecorationPassthroughNV: return "PassthroughNV"; case SpvDecorationViewportRelativeNV: return "ViewportRelativeNV"; @@ -3732,7 +3991,10 @@ inline const char* SpvDecorationToString(SpvDecoration value) { case SpvDecorationNonUniform: return "NonUniform"; case SpvDecorationRestrictPointer: return "RestrictPointer"; case SpvDecorationAliasedPointer: return "AliasedPointer"; + case SpvDecorationMemberOffsetNV: return "MemberOffsetNV"; case SpvDecorationHitObjectShaderRecordBufferNV: return "HitObjectShaderRecordBufferNV"; + case SpvDecorationHitObjectShaderRecordBufferEXT: return "HitObjectShaderRecordBufferEXT"; + case SpvDecorationBankNV: return "BankNV"; case SpvDecorationBindlessSamplerNV: return "BindlessSamplerNV"; case SpvDecorationBindlessImageNV: return "BindlessImageNV"; case SpvDecorationBoundSamplerNV: return "BoundSamplerNV"; @@ -3751,55 +4013,55 @@ inline const char* SpvDecorationToString(SpvDecoration value) { case SpvDecorationUserTypeGOOGLE: return "UserTypeGOOGLE"; case SpvDecorationFunctionRoundingModeINTEL: return "FunctionRoundingModeINTEL"; case SpvDecorationFunctionDenormModeINTEL: return "FunctionDenormModeINTEL"; - case SpvDecorationRegisterINTEL: return "RegisterINTEL"; - case SpvDecorationMemoryINTEL: return "MemoryINTEL"; - case SpvDecorationNumbanksINTEL: return "NumbanksINTEL"; - case SpvDecorationBankwidthINTEL: return "BankwidthINTEL"; - case SpvDecorationMaxPrivateCopiesINTEL: return "MaxPrivateCopiesINTEL"; - case SpvDecorationSinglepumpINTEL: return "SinglepumpINTEL"; - case SpvDecorationDoublepumpINTEL: return "DoublepumpINTEL"; - case SpvDecorationMaxReplicatesINTEL: return "MaxReplicatesINTEL"; - case SpvDecorationSimpleDualPortINTEL: return "SimpleDualPortINTEL"; - case SpvDecorationMergeINTEL: return "MergeINTEL"; - case SpvDecorationBankBitsINTEL: return "BankBitsINTEL"; - case SpvDecorationForcePow2DepthINTEL: return "ForcePow2DepthINTEL"; - case SpvDecorationStridesizeINTEL: return "StridesizeINTEL"; - case SpvDecorationWordsizeINTEL: return "WordsizeINTEL"; - case SpvDecorationTrueDualPortINTEL: return "TrueDualPortINTEL"; - case SpvDecorationBurstCoalesceINTEL: return "BurstCoalesceINTEL"; - case SpvDecorationCacheSizeINTEL: return "CacheSizeINTEL"; - case SpvDecorationDontStaticallyCoalesceINTEL: return "DontStaticallyCoalesceINTEL"; - case SpvDecorationPrefetchINTEL: return "PrefetchINTEL"; - case SpvDecorationStallEnableINTEL: return "StallEnableINTEL"; - case SpvDecorationFuseLoopsInFunctionINTEL: return "FuseLoopsInFunctionINTEL"; - case SpvDecorationMathOpDSPModeINTEL: return "MathOpDSPModeINTEL"; + case SpvDecorationRegisterALTERA: return "RegisterALTERA"; + case SpvDecorationMemoryALTERA: return "MemoryALTERA"; + case SpvDecorationNumbanksALTERA: return "NumbanksALTERA"; + case SpvDecorationBankwidthALTERA: return "BankwidthALTERA"; + case SpvDecorationMaxPrivateCopiesALTERA: return "MaxPrivateCopiesALTERA"; + case SpvDecorationSinglepumpALTERA: return "SinglepumpALTERA"; + case SpvDecorationDoublepumpALTERA: return "DoublepumpALTERA"; + case SpvDecorationMaxReplicatesALTERA: return "MaxReplicatesALTERA"; + case SpvDecorationSimpleDualPortALTERA: return "SimpleDualPortALTERA"; + case SpvDecorationMergeALTERA: return "MergeALTERA"; + case SpvDecorationBankBitsALTERA: return "BankBitsALTERA"; + case SpvDecorationForcePow2DepthALTERA: return "ForcePow2DepthALTERA"; + case SpvDecorationStridesizeALTERA: return "StridesizeALTERA"; + case SpvDecorationWordsizeALTERA: return "WordsizeALTERA"; + case SpvDecorationTrueDualPortALTERA: return "TrueDualPortALTERA"; + case SpvDecorationBurstCoalesceALTERA: return "BurstCoalesceALTERA"; + case SpvDecorationCacheSizeALTERA: return "CacheSizeALTERA"; + case SpvDecorationDontStaticallyCoalesceALTERA: return "DontStaticallyCoalesceALTERA"; + case SpvDecorationPrefetchALTERA: return "PrefetchALTERA"; + case SpvDecorationStallEnableALTERA: return "StallEnableALTERA"; + case SpvDecorationFuseLoopsInFunctionALTERA: return "FuseLoopsInFunctionALTERA"; + case SpvDecorationMathOpDSPModeALTERA: return "MathOpDSPModeALTERA"; case SpvDecorationAliasScopeINTEL: return "AliasScopeINTEL"; case SpvDecorationNoAliasINTEL: return "NoAliasINTEL"; - case SpvDecorationInitiationIntervalINTEL: return "InitiationIntervalINTEL"; - case SpvDecorationMaxConcurrencyINTEL: return "MaxConcurrencyINTEL"; - case SpvDecorationPipelineEnableINTEL: return "PipelineEnableINTEL"; - case SpvDecorationBufferLocationINTEL: return "BufferLocationINTEL"; - case SpvDecorationIOPipeStorageINTEL: return "IOPipeStorageINTEL"; + case SpvDecorationInitiationIntervalALTERA: return "InitiationIntervalALTERA"; + case SpvDecorationMaxConcurrencyALTERA: return "MaxConcurrencyALTERA"; + case SpvDecorationPipelineEnableALTERA: return "PipelineEnableALTERA"; + case SpvDecorationBufferLocationALTERA: return "BufferLocationALTERA"; + case SpvDecorationIOPipeStorageALTERA: return "IOPipeStorageALTERA"; case SpvDecorationFunctionFloatingPointModeINTEL: return "FunctionFloatingPointModeINTEL"; case SpvDecorationSingleElementVectorINTEL: return "SingleElementVectorINTEL"; case SpvDecorationVectorComputeCallableFunctionINTEL: return "VectorComputeCallableFunctionINTEL"; case SpvDecorationMediaBlockIOINTEL: return "MediaBlockIOINTEL"; - case SpvDecorationStallFreeINTEL: return "StallFreeINTEL"; + case SpvDecorationStallFreeALTERA: return "StallFreeALTERA"; case SpvDecorationFPMaxErrorDecorationINTEL: return "FPMaxErrorDecorationINTEL"; - case SpvDecorationLatencyControlLabelINTEL: return "LatencyControlLabelINTEL"; - case SpvDecorationLatencyControlConstraintINTEL: return "LatencyControlConstraintINTEL"; - case SpvDecorationConduitKernelArgumentINTEL: return "ConduitKernelArgumentINTEL"; - case SpvDecorationRegisterMapKernelArgumentINTEL: return "RegisterMapKernelArgumentINTEL"; - case SpvDecorationMMHostInterfaceAddressWidthINTEL: return "MMHostInterfaceAddressWidthINTEL"; - case SpvDecorationMMHostInterfaceDataWidthINTEL: return "MMHostInterfaceDataWidthINTEL"; - case SpvDecorationMMHostInterfaceLatencyINTEL: return "MMHostInterfaceLatencyINTEL"; - case SpvDecorationMMHostInterfaceReadWriteModeINTEL: return "MMHostInterfaceReadWriteModeINTEL"; - case SpvDecorationMMHostInterfaceMaxBurstINTEL: return "MMHostInterfaceMaxBurstINTEL"; - case SpvDecorationMMHostInterfaceWaitRequestINTEL: return "MMHostInterfaceWaitRequestINTEL"; - case SpvDecorationStableKernelArgumentINTEL: return "StableKernelArgumentINTEL"; + case SpvDecorationLatencyControlLabelALTERA: return "LatencyControlLabelALTERA"; + case SpvDecorationLatencyControlConstraintALTERA: return "LatencyControlConstraintALTERA"; + case SpvDecorationConduitKernelArgumentALTERA: return "ConduitKernelArgumentALTERA"; + case SpvDecorationRegisterMapKernelArgumentALTERA: return "RegisterMapKernelArgumentALTERA"; + case SpvDecorationMMHostInterfaceAddressWidthALTERA: return "MMHostInterfaceAddressWidthALTERA"; + case SpvDecorationMMHostInterfaceDataWidthALTERA: return "MMHostInterfaceDataWidthALTERA"; + case SpvDecorationMMHostInterfaceLatencyALTERA: return "MMHostInterfaceLatencyALTERA"; + case SpvDecorationMMHostInterfaceReadWriteModeALTERA: return "MMHostInterfaceReadWriteModeALTERA"; + case SpvDecorationMMHostInterfaceMaxBurstALTERA: return "MMHostInterfaceMaxBurstALTERA"; + case SpvDecorationMMHostInterfaceWaitRequestALTERA: return "MMHostInterfaceWaitRequestALTERA"; + case SpvDecorationStableKernelArgumentALTERA: return "StableKernelArgumentALTERA"; case SpvDecorationHostAccessINTEL: return "HostAccessINTEL"; - case SpvDecorationInitModeINTEL: return "InitModeINTEL"; - case SpvDecorationImplementInRegisterMapINTEL: return "ImplementInRegisterMapINTEL"; + case SpvDecorationInitModeALTERA: return "InitModeALTERA"; + case SpvDecorationImplementInRegisterMapALTERA: return "ImplementInRegisterMapALTERA"; case SpvDecorationConditionalINTEL: return "ConditionalINTEL"; case SpvDecorationCacheControlLoadINTEL: return "CacheControlLoadINTEL"; case SpvDecorationCacheControlStoreINTEL: return "CacheControlStoreINTEL"; @@ -3880,6 +4142,8 @@ inline const char* SpvBuiltInToString(SpvBuiltIn value) { case SpvBuiltInFragStencilRefEXT: return "FragStencilRefEXT"; case SpvBuiltInRemainingRecursionLevelsAMDX: return "RemainingRecursionLevelsAMDX"; case SpvBuiltInShaderIndexAMDX: return "ShaderIndexAMDX"; + case SpvBuiltInSamplerHeapEXT: return "SamplerHeapEXT"; + case SpvBuiltInResourceHeapEXT: return "ResourceHeapEXT"; case SpvBuiltInViewportMaskNV: return "ViewportMaskNV"; case SpvBuiltInSecondaryPositionNV: return "SecondaryPositionNV"; case SpvBuiltInSecondaryViewportMaskNV: return "SecondaryViewportMaskNV"; @@ -3958,9 +4222,9 @@ inline const char* SpvGroupOperationToString(SpvGroupOperation value) { case SpvGroupOperationInclusiveScan: return "InclusiveScan"; case SpvGroupOperationExclusiveScan: return "ExclusiveScan"; case SpvGroupOperationClusteredReduce: return "ClusteredReduce"; - case SpvGroupOperationPartitionedReduceNV: return "PartitionedReduceNV"; - case SpvGroupOperationPartitionedInclusiveScanNV: return "PartitionedInclusiveScanNV"; - case SpvGroupOperationPartitionedExclusiveScanNV: return "PartitionedExclusiveScanNV"; + case SpvGroupOperationPartitionedReduceEXT: return "PartitionedReduceEXT"; + case SpvGroupOperationPartitionedInclusiveScanEXT: return "PartitionedInclusiveScanEXT"; + case SpvGroupOperationPartitionedExclusiveScanEXT: return "PartitionedExclusiveScanEXT"; default: return "Unknown"; } } @@ -4107,6 +4371,10 @@ inline const char* SpvCapabilityToString(SpvCapability value) { case SpvCapabilityBFloat16TypeKHR: return "BFloat16TypeKHR"; case SpvCapabilityBFloat16DotProductKHR: return "BFloat16DotProductKHR"; case SpvCapabilityBFloat16CooperativeMatrixKHR: return "BFloat16CooperativeMatrixKHR"; + case SpvCapabilityAbortKHR: return "AbortKHR"; + case SpvCapabilityDescriptorHeapEXT: return "DescriptorHeapEXT"; + case SpvCapabilityConstantDataKHR: return "ConstantDataKHR"; + case SpvCapabilityPoisonFreezeKHR: return "PoisonFreezeKHR"; case SpvCapabilitySampleMaskOverrideCoverageNV: return "SampleMaskOverrideCoverageNV"; case SpvCapabilityGeometryShaderPassthroughNV: return "GeometryShaderPassthroughNV"; case SpvCapabilityShaderViewportIndexLayerEXT: return "ShaderViewportIndexLayerEXT"; @@ -4120,7 +4388,7 @@ inline const char* SpvCapabilityToString(SpvCapability value) { case SpvCapabilityFragmentBarycentricKHR: return "FragmentBarycentricKHR"; case SpvCapabilityComputeDerivativeGroupQuadsKHR: return "ComputeDerivativeGroupQuadsKHR"; case SpvCapabilityFragmentDensityEXT: return "FragmentDensityEXT"; - case SpvCapabilityGroupNonUniformPartitionedNV: return "GroupNonUniformPartitionedNV"; + case SpvCapabilityGroupNonUniformPartitionedEXT: return "GroupNonUniformPartitionedEXT"; case SpvCapabilityShaderNonUniform: return "ShaderNonUniform"; case SpvCapabilityRuntimeDescriptorArray: return "RuntimeDescriptorArray"; case SpvCapabilityInputAttachmentArrayDynamicIndexing: return "InputAttachmentArrayDynamicIndexing"; @@ -4150,6 +4418,7 @@ inline const char* SpvCapabilityToString(SpvCapability value) { case SpvCapabilityDisplacementMicromapNV: return "DisplacementMicromapNV"; case SpvCapabilityRayTracingOpacityMicromapEXT: return "RayTracingOpacityMicromapEXT"; case SpvCapabilityShaderInvocationReorderNV: return "ShaderInvocationReorderNV"; + case SpvCapabilityShaderInvocationReorderEXT: return "ShaderInvocationReorderEXT"; case SpvCapabilityBindlessTextureNV: return "BindlessTextureNV"; case SpvCapabilityRayQueryPositionFetchKHR: return "RayQueryPositionFetchKHR"; case SpvCapabilityCooperativeVectorNV: return "CooperativeVectorNV"; @@ -4158,6 +4427,9 @@ inline const char* SpvCapabilityToString(SpvCapability value) { case SpvCapabilityRawAccessChainsNV: return "RawAccessChainsNV"; case SpvCapabilityRayTracingSpheresGeometryNV: return "RayTracingSpheresGeometryNV"; case SpvCapabilityRayTracingLinearSweptSpheresGeometryNV: return "RayTracingLinearSweptSpheresGeometryNV"; + case SpvCapabilityPushConstantBanksNV: return "PushConstantBanksNV"; + case SpvCapabilityLongVectorEXT: return "LongVectorEXT"; + case SpvCapabilityShader64BitIndexingEXT: return "Shader64BitIndexingEXT"; case SpvCapabilityCooperativeMatrixReductionsNV: return "CooperativeMatrixReductionsNV"; case SpvCapabilityCooperativeMatrixConversionsNV: return "CooperativeMatrixConversionsNV"; case SpvCapabilityCooperativeMatrixPerElementOperationsNV: return "CooperativeMatrixPerElementOperationsNV"; @@ -4187,27 +4459,27 @@ inline const char* SpvCapabilityToString(SpvCapability value) { case SpvCapabilitySubgroupAvcMotionEstimationChromaINTEL: return "SubgroupAvcMotionEstimationChromaINTEL"; case SpvCapabilityVariableLengthArrayINTEL: return "VariableLengthArrayINTEL"; case SpvCapabilityFunctionFloatControlINTEL: return "FunctionFloatControlINTEL"; - case SpvCapabilityFPGAMemoryAttributesINTEL: return "FPGAMemoryAttributesINTEL"; + case SpvCapabilityFPGAMemoryAttributesALTERA: return "FPGAMemoryAttributesALTERA"; case SpvCapabilityFPFastMathModeINTEL: return "FPFastMathModeINTEL"; - case SpvCapabilityArbitraryPrecisionIntegersINTEL: return "ArbitraryPrecisionIntegersINTEL"; - case SpvCapabilityArbitraryPrecisionFloatingPointINTEL: return "ArbitraryPrecisionFloatingPointINTEL"; + case SpvCapabilityArbitraryPrecisionIntegersALTERA: return "ArbitraryPrecisionIntegersALTERA"; + case SpvCapabilityArbitraryPrecisionFloatingPointALTERA: return "ArbitraryPrecisionFloatingPointALTERA"; case SpvCapabilityUnstructuredLoopControlsINTEL: return "UnstructuredLoopControlsINTEL"; - case SpvCapabilityFPGALoopControlsINTEL: return "FPGALoopControlsINTEL"; + case SpvCapabilityFPGALoopControlsALTERA: return "FPGALoopControlsALTERA"; case SpvCapabilityKernelAttributesINTEL: return "KernelAttributesINTEL"; case SpvCapabilityFPGAKernelAttributesINTEL: return "FPGAKernelAttributesINTEL"; - case SpvCapabilityFPGAMemoryAccessesINTEL: return "FPGAMemoryAccessesINTEL"; - case SpvCapabilityFPGAClusterAttributesINTEL: return "FPGAClusterAttributesINTEL"; - case SpvCapabilityLoopFuseINTEL: return "LoopFuseINTEL"; - case SpvCapabilityFPGADSPControlINTEL: return "FPGADSPControlINTEL"; + case SpvCapabilityFPGAMemoryAccessesALTERA: return "FPGAMemoryAccessesALTERA"; + case SpvCapabilityFPGAClusterAttributesALTERA: return "FPGAClusterAttributesALTERA"; + case SpvCapabilityLoopFuseALTERA: return "LoopFuseALTERA"; + case SpvCapabilityFPGADSPControlALTERA: return "FPGADSPControlALTERA"; case SpvCapabilityMemoryAccessAliasingINTEL: return "MemoryAccessAliasingINTEL"; - case SpvCapabilityFPGAInvocationPipeliningAttributesINTEL: return "FPGAInvocationPipeliningAttributesINTEL"; - case SpvCapabilityFPGABufferLocationINTEL: return "FPGABufferLocationINTEL"; - case SpvCapabilityArbitraryPrecisionFixedPointINTEL: return "ArbitraryPrecisionFixedPointINTEL"; - case SpvCapabilityUSMStorageClassesINTEL: return "USMStorageClassesINTEL"; - case SpvCapabilityRuntimeAlignedAttributeINTEL: return "RuntimeAlignedAttributeINTEL"; - case SpvCapabilityIOPipesINTEL: return "IOPipesINTEL"; - case SpvCapabilityBlockingPipesINTEL: return "BlockingPipesINTEL"; - case SpvCapabilityFPGARegINTEL: return "FPGARegINTEL"; + case SpvCapabilityFPGAInvocationPipeliningAttributesALTERA: return "FPGAInvocationPipeliningAttributesALTERA"; + case SpvCapabilityFPGABufferLocationALTERA: return "FPGABufferLocationALTERA"; + case SpvCapabilityArbitraryPrecisionFixedPointALTERA: return "ArbitraryPrecisionFixedPointALTERA"; + case SpvCapabilityUSMStorageClassesALTERA: return "USMStorageClassesALTERA"; + case SpvCapabilityRuntimeAlignedAttributeALTERA: return "RuntimeAlignedAttributeALTERA"; + case SpvCapabilityIOPipesALTERA: return "IOPipesALTERA"; + case SpvCapabilityBlockingPipesALTERA: return "BlockingPipesALTERA"; + case SpvCapabilityFPGARegALTERA: return "FPGARegALTERA"; case SpvCapabilityDotProductInputAll: return "DotProductInputAll"; case SpvCapabilityDotProductInput4x8Bit: return "DotProductInput4x8Bit"; case SpvCapabilityDotProductInput4x8BitPacked: return "DotProductInput4x8BitPacked"; @@ -4228,14 +4500,14 @@ inline const char* SpvCapabilityToString(SpvCapability value) { case SpvCapabilityBFloat16ConversionINTEL: return "BFloat16ConversionINTEL"; case SpvCapabilitySplitBarrierINTEL: return "SplitBarrierINTEL"; case SpvCapabilityArithmeticFenceEXT: return "ArithmeticFenceEXT"; - case SpvCapabilityFPGAClusterAttributesV2INTEL: return "FPGAClusterAttributesV2INTEL"; + case SpvCapabilityFPGAClusterAttributesV2ALTERA: return "FPGAClusterAttributesV2ALTERA"; case SpvCapabilityFPGAKernelAttributesv2INTEL: return "FPGAKernelAttributesv2INTEL"; - case SpvCapabilityTaskSequenceINTEL: return "TaskSequenceINTEL"; + case SpvCapabilityTaskSequenceALTERA: return "TaskSequenceALTERA"; case SpvCapabilityFPMaxErrorINTEL: return "FPMaxErrorINTEL"; - case SpvCapabilityFPGALatencyControlINTEL: return "FPGALatencyControlINTEL"; - case SpvCapabilityFPGAArgumentInterfacesINTEL: return "FPGAArgumentInterfacesINTEL"; + case SpvCapabilityFPGALatencyControlALTERA: return "FPGALatencyControlALTERA"; + case SpvCapabilityFPGAArgumentInterfacesALTERA: return "FPGAArgumentInterfacesALTERA"; case SpvCapabilityGlobalVariableHostAccessINTEL: return "GlobalVariableHostAccessINTEL"; - case SpvCapabilityGlobalVariableFPGADecorationsINTEL: return "GlobalVariableFPGADecorationsINTEL"; + case SpvCapabilityGlobalVariableFPGADecorationsALTERA: return "GlobalVariableFPGADecorationsALTERA"; case SpvCapabilitySubgroupBufferPrefetchINTEL: return "SubgroupBufferPrefetchINTEL"; case SpvCapabilitySubgroup2DBlockIOINTEL: return "Subgroup2DBlockIOINTEL"; case SpvCapabilitySubgroup2DBlockTransformINTEL: return "Subgroup2DBlockTransformINTEL"; @@ -4251,6 +4523,10 @@ inline const char* SpvCapabilityToString(SpvCapability value) { case SpvCapabilityCacheControlsINTEL: return "CacheControlsINTEL"; case SpvCapabilityRegisterLimitsINTEL: return "RegisterLimitsINTEL"; case SpvCapabilityBindlessImagesINTEL: return "BindlessImagesINTEL"; + case SpvCapabilityDotProductFloat16AccFloat32VALVE: return "DotProductFloat16AccFloat32VALVE"; + case SpvCapabilityDotProductFloat16AccFloat16VALVE: return "DotProductFloat16AccFloat16VALVE"; + case SpvCapabilityDotProductBFloat16AccVALVE: return "DotProductBFloat16AccVALVE"; + case SpvCapabilityDotProductFloat8AccFloat32VALVE: return "DotProductFloat8AccFloat32VALVE"; default: return "Unknown"; } } @@ -4359,8 +4635,8 @@ inline const char* SpvTensorClampModeToString(SpvTensorClampMode value) { inline const char* SpvInitializationModeQualifierToString(SpvInitializationModeQualifier value) { switch (value) { - case SpvInitializationModeQualifierInitOnDeviceReprogramINTEL: return "InitOnDeviceReprogramINTEL"; - case SpvInitializationModeQualifierInitOnDeviceResetINTEL: return "InitOnDeviceResetINTEL"; + case SpvInitializationModeQualifierInitOnDeviceReprogramALTERA: return "InitOnDeviceReprogramALTERA"; + case SpvInitializationModeQualifierInitOnDeviceResetALTERA: return "InitOnDeviceResetALTERA"; default: return "Unknown"; } } @@ -4881,6 +5157,16 @@ inline const char* SpvOpToString(SpvOp value) { case SpvOpSpecConstantStringAMDX: return "OpSpecConstantStringAMDX"; case SpvOpGroupNonUniformQuadAllKHR: return "OpGroupNonUniformQuadAllKHR"; case SpvOpGroupNonUniformQuadAnyKHR: return "OpGroupNonUniformQuadAnyKHR"; + case SpvOpTypeBufferEXT: return "OpTypeBufferEXT"; + case SpvOpBufferPointerEXT: return "OpBufferPointerEXT"; + case SpvOpAbortKHR: return "OpAbortKHR"; + case SpvOpUntypedImageTexelPointerEXT: return "OpUntypedImageTexelPointerEXT"; + case SpvOpMemberDecorateIdEXT: return "OpMemberDecorateIdEXT"; + case SpvOpConstantSizeOfEXT: return "OpConstantSizeOfEXT"; + case SpvOpConstantDataKHR: return "OpConstantDataKHR"; + case SpvOpSpecConstantDataKHR: return "OpSpecConstantDataKHR"; + case SpvOpPoisonKHR: return "OpPoisonKHR"; + case SpvOpFreezeKHR: return "OpFreezeKHR"; case SpvOpHitObjectRecordHitMotionNV: return "OpHitObjectRecordHitMotionNV"; case SpvOpHitObjectRecordHitWithIndexMotionNV: return "OpHitObjectRecordHitWithIndexMotionNV"; case SpvOpHitObjectRecordMissMotionNV: return "OpHitObjectRecordMissMotionNV"; @@ -4923,12 +5209,42 @@ inline const char* SpvOpToString(SpvOp value) { case SpvOpCooperativeMatrixConvertNV: return "OpCooperativeMatrixConvertNV"; case SpvOpEmitMeshTasksEXT: return "OpEmitMeshTasksEXT"; case SpvOpSetMeshOutputsEXT: return "OpSetMeshOutputsEXT"; - case SpvOpGroupNonUniformPartitionNV: return "OpGroupNonUniformPartitionNV"; + case SpvOpGroupNonUniformPartitionEXT: return "OpGroupNonUniformPartitionEXT"; case SpvOpWritePackedPrimitiveIndices4x8NV: return "OpWritePackedPrimitiveIndices4x8NV"; case SpvOpFetchMicroTriangleVertexPositionNV: return "OpFetchMicroTriangleVertexPositionNV"; case SpvOpFetchMicroTriangleVertexBarycentricNV: return "OpFetchMicroTriangleVertexBarycentricNV"; case SpvOpCooperativeVectorLoadNV: return "OpCooperativeVectorLoadNV"; case SpvOpCooperativeVectorStoreNV: return "OpCooperativeVectorStoreNV"; + case SpvOpHitObjectRecordFromQueryEXT: return "OpHitObjectRecordFromQueryEXT"; + case SpvOpHitObjectRecordMissEXT: return "OpHitObjectRecordMissEXT"; + case SpvOpHitObjectRecordMissMotionEXT: return "OpHitObjectRecordMissMotionEXT"; + case SpvOpHitObjectGetIntersectionTriangleVertexPositionsEXT: return "OpHitObjectGetIntersectionTriangleVertexPositionsEXT"; + case SpvOpHitObjectGetRayFlagsEXT: return "OpHitObjectGetRayFlagsEXT"; + case SpvOpHitObjectSetShaderBindingTableRecordIndexEXT: return "OpHitObjectSetShaderBindingTableRecordIndexEXT"; + case SpvOpHitObjectReorderExecuteShaderEXT: return "OpHitObjectReorderExecuteShaderEXT"; + case SpvOpHitObjectTraceReorderExecuteEXT: return "OpHitObjectTraceReorderExecuteEXT"; + case SpvOpHitObjectTraceMotionReorderExecuteEXT: return "OpHitObjectTraceMotionReorderExecuteEXT"; + case SpvOpTypeHitObjectEXT: return "OpTypeHitObjectEXT"; + case SpvOpReorderThreadWithHintEXT: return "OpReorderThreadWithHintEXT"; + case SpvOpReorderThreadWithHitObjectEXT: return "OpReorderThreadWithHitObjectEXT"; + case SpvOpHitObjectTraceRayEXT: return "OpHitObjectTraceRayEXT"; + case SpvOpHitObjectTraceRayMotionEXT: return "OpHitObjectTraceRayMotionEXT"; + case SpvOpHitObjectRecordEmptyEXT: return "OpHitObjectRecordEmptyEXT"; + case SpvOpHitObjectExecuteShaderEXT: return "OpHitObjectExecuteShaderEXT"; + case SpvOpHitObjectGetCurrentTimeEXT: return "OpHitObjectGetCurrentTimeEXT"; + case SpvOpHitObjectGetAttributesEXT: return "OpHitObjectGetAttributesEXT"; + case SpvOpHitObjectGetHitKindEXT: return "OpHitObjectGetHitKindEXT"; + case SpvOpHitObjectGetPrimitiveIndexEXT: return "OpHitObjectGetPrimitiveIndexEXT"; + case SpvOpHitObjectGetGeometryIndexEXT: return "OpHitObjectGetGeometryIndexEXT"; + case SpvOpHitObjectGetInstanceIdEXT: return "OpHitObjectGetInstanceIdEXT"; + case SpvOpHitObjectGetInstanceCustomIndexEXT: return "OpHitObjectGetInstanceCustomIndexEXT"; + case SpvOpHitObjectGetObjectRayOriginEXT: return "OpHitObjectGetObjectRayOriginEXT"; + case SpvOpHitObjectGetObjectRayDirectionEXT: return "OpHitObjectGetObjectRayDirectionEXT"; + case SpvOpHitObjectGetWorldRayDirectionEXT: return "OpHitObjectGetWorldRayDirectionEXT"; + case SpvOpHitObjectGetWorldRayOriginEXT: return "OpHitObjectGetWorldRayOriginEXT"; + case SpvOpHitObjectGetObjectToWorldEXT: return "OpHitObjectGetObjectToWorldEXT"; + case SpvOpHitObjectGetWorldToObjectEXT: return "OpHitObjectGetWorldToObjectEXT"; + case SpvOpHitObjectGetRayTMaxEXT: return "OpHitObjectGetRayTMaxEXT"; case SpvOpReportIntersectionKHR: return "OpReportIntersectionKHR"; case SpvOpIgnoreIntersectionNV: return "OpIgnoreIntersectionNV"; case SpvOpTerminateRayNV: return "OpTerminateRayNV"; @@ -4940,6 +5256,12 @@ inline const char* SpvOpToString(SpvOp value) { case SpvOpExecuteCallableNV: return "OpExecuteCallableNV"; case SpvOpRayQueryGetClusterIdNV: return "OpRayQueryGetClusterIdNV"; case SpvOpHitObjectGetClusterIdNV: return "OpHitObjectGetClusterIdNV"; + case SpvOpHitObjectGetRayTMinEXT: return "OpHitObjectGetRayTMinEXT"; + case SpvOpHitObjectGetShaderBindingTableRecordIndexEXT: return "OpHitObjectGetShaderBindingTableRecordIndexEXT"; + case SpvOpHitObjectGetShaderRecordBufferHandleEXT: return "OpHitObjectGetShaderRecordBufferHandleEXT"; + case SpvOpHitObjectIsEmptyEXT: return "OpHitObjectIsEmptyEXT"; + case SpvOpHitObjectIsHitEXT: return "OpHitObjectIsHitEXT"; + case SpvOpHitObjectIsMissEXT: return "OpHitObjectIsMissEXT"; case SpvOpTypeCooperativeMatrixNV: return "OpTypeCooperativeMatrixNV"; case SpvOpCooperativeMatrixLoadNV: return "OpCooperativeMatrixLoadNV"; case SpvOpCooperativeMatrixStoreNV: return "OpCooperativeMatrixStoreNV"; @@ -5143,24 +5465,24 @@ inline const char* SpvOpToString(SpvOp value) { case SpvOpVariableLengthArrayINTEL: return "OpVariableLengthArrayINTEL"; case SpvOpSaveMemoryINTEL: return "OpSaveMemoryINTEL"; case SpvOpRestoreMemoryINTEL: return "OpRestoreMemoryINTEL"; - case SpvOpArbitraryFloatSinCosPiINTEL: return "OpArbitraryFloatSinCosPiINTEL"; - case SpvOpArbitraryFloatCastINTEL: return "OpArbitraryFloatCastINTEL"; - case SpvOpArbitraryFloatCastFromIntINTEL: return "OpArbitraryFloatCastFromIntINTEL"; - case SpvOpArbitraryFloatCastToIntINTEL: return "OpArbitraryFloatCastToIntINTEL"; - case SpvOpArbitraryFloatAddINTEL: return "OpArbitraryFloatAddINTEL"; - case SpvOpArbitraryFloatSubINTEL: return "OpArbitraryFloatSubINTEL"; - case SpvOpArbitraryFloatMulINTEL: return "OpArbitraryFloatMulINTEL"; - case SpvOpArbitraryFloatDivINTEL: return "OpArbitraryFloatDivINTEL"; - case SpvOpArbitraryFloatGTINTEL: return "OpArbitraryFloatGTINTEL"; - case SpvOpArbitraryFloatGEINTEL: return "OpArbitraryFloatGEINTEL"; - case SpvOpArbitraryFloatLTINTEL: return "OpArbitraryFloatLTINTEL"; - case SpvOpArbitraryFloatLEINTEL: return "OpArbitraryFloatLEINTEL"; - case SpvOpArbitraryFloatEQINTEL: return "OpArbitraryFloatEQINTEL"; - case SpvOpArbitraryFloatRecipINTEL: return "OpArbitraryFloatRecipINTEL"; - case SpvOpArbitraryFloatRSqrtINTEL: return "OpArbitraryFloatRSqrtINTEL"; - case SpvOpArbitraryFloatCbrtINTEL: return "OpArbitraryFloatCbrtINTEL"; - case SpvOpArbitraryFloatHypotINTEL: return "OpArbitraryFloatHypotINTEL"; - case SpvOpArbitraryFloatSqrtINTEL: return "OpArbitraryFloatSqrtINTEL"; + case SpvOpArbitraryFloatSinCosPiALTERA: return "OpArbitraryFloatSinCosPiALTERA"; + case SpvOpArbitraryFloatCastALTERA: return "OpArbitraryFloatCastALTERA"; + case SpvOpArbitraryFloatCastFromIntALTERA: return "OpArbitraryFloatCastFromIntALTERA"; + case SpvOpArbitraryFloatCastToIntALTERA: return "OpArbitraryFloatCastToIntALTERA"; + case SpvOpArbitraryFloatAddALTERA: return "OpArbitraryFloatAddALTERA"; + case SpvOpArbitraryFloatSubALTERA: return "OpArbitraryFloatSubALTERA"; + case SpvOpArbitraryFloatMulALTERA: return "OpArbitraryFloatMulALTERA"; + case SpvOpArbitraryFloatDivALTERA: return "OpArbitraryFloatDivALTERA"; + case SpvOpArbitraryFloatGTALTERA: return "OpArbitraryFloatGTALTERA"; + case SpvOpArbitraryFloatGEALTERA: return "OpArbitraryFloatGEALTERA"; + case SpvOpArbitraryFloatLTALTERA: return "OpArbitraryFloatLTALTERA"; + case SpvOpArbitraryFloatLEALTERA: return "OpArbitraryFloatLEALTERA"; + case SpvOpArbitraryFloatEQALTERA: return "OpArbitraryFloatEQALTERA"; + case SpvOpArbitraryFloatRecipALTERA: return "OpArbitraryFloatRecipALTERA"; + case SpvOpArbitraryFloatRSqrtALTERA: return "OpArbitraryFloatRSqrtALTERA"; + case SpvOpArbitraryFloatCbrtALTERA: return "OpArbitraryFloatCbrtALTERA"; + case SpvOpArbitraryFloatHypotALTERA: return "OpArbitraryFloatHypotALTERA"; + case SpvOpArbitraryFloatSqrtALTERA: return "OpArbitraryFloatSqrtALTERA"; case SpvOpArbitraryFloatLogINTEL: return "OpArbitraryFloatLogINTEL"; case SpvOpArbitraryFloatLog2INTEL: return "OpArbitraryFloatLog2INTEL"; case SpvOpArbitraryFloatLog10INTEL: return "OpArbitraryFloatLog10INTEL"; @@ -5188,22 +5510,22 @@ inline const char* SpvOpToString(SpvOp value) { case SpvOpAliasDomainDeclINTEL: return "OpAliasDomainDeclINTEL"; case SpvOpAliasScopeDeclINTEL: return "OpAliasScopeDeclINTEL"; case SpvOpAliasScopeListDeclINTEL: return "OpAliasScopeListDeclINTEL"; - case SpvOpFixedSqrtINTEL: return "OpFixedSqrtINTEL"; - case SpvOpFixedRecipINTEL: return "OpFixedRecipINTEL"; - case SpvOpFixedRsqrtINTEL: return "OpFixedRsqrtINTEL"; - case SpvOpFixedSinINTEL: return "OpFixedSinINTEL"; - case SpvOpFixedCosINTEL: return "OpFixedCosINTEL"; - case SpvOpFixedSinCosINTEL: return "OpFixedSinCosINTEL"; - case SpvOpFixedSinPiINTEL: return "OpFixedSinPiINTEL"; - case SpvOpFixedCosPiINTEL: return "OpFixedCosPiINTEL"; - case SpvOpFixedSinCosPiINTEL: return "OpFixedSinCosPiINTEL"; - case SpvOpFixedLogINTEL: return "OpFixedLogINTEL"; - case SpvOpFixedExpINTEL: return "OpFixedExpINTEL"; - case SpvOpPtrCastToCrossWorkgroupINTEL: return "OpPtrCastToCrossWorkgroupINTEL"; - case SpvOpCrossWorkgroupCastToPtrINTEL: return "OpCrossWorkgroupCastToPtrINTEL"; - case SpvOpReadPipeBlockingINTEL: return "OpReadPipeBlockingINTEL"; - case SpvOpWritePipeBlockingINTEL: return "OpWritePipeBlockingINTEL"; - case SpvOpFPGARegINTEL: return "OpFPGARegINTEL"; + case SpvOpFixedSqrtALTERA: return "OpFixedSqrtALTERA"; + case SpvOpFixedRecipALTERA: return "OpFixedRecipALTERA"; + case SpvOpFixedRsqrtALTERA: return "OpFixedRsqrtALTERA"; + case SpvOpFixedSinALTERA: return "OpFixedSinALTERA"; + case SpvOpFixedCosALTERA: return "OpFixedCosALTERA"; + case SpvOpFixedSinCosALTERA: return "OpFixedSinCosALTERA"; + case SpvOpFixedSinPiALTERA: return "OpFixedSinPiALTERA"; + case SpvOpFixedCosPiALTERA: return "OpFixedCosPiALTERA"; + case SpvOpFixedSinCosPiALTERA: return "OpFixedSinCosPiALTERA"; + case SpvOpFixedLogALTERA: return "OpFixedLogALTERA"; + case SpvOpFixedExpALTERA: return "OpFixedExpALTERA"; + case SpvOpPtrCastToCrossWorkgroupALTERA: return "OpPtrCastToCrossWorkgroupALTERA"; + case SpvOpCrossWorkgroupCastToPtrALTERA: return "OpCrossWorkgroupCastToPtrALTERA"; + case SpvOpReadPipeBlockingALTERA: return "OpReadPipeBlockingALTERA"; + case SpvOpWritePipeBlockingALTERA: return "OpWritePipeBlockingALTERA"; + case SpvOpFPGARegALTERA: return "OpFPGARegALTERA"; case SpvOpRayQueryGetRayTMinKHR: return "OpRayQueryGetRayTMinKHR"; case SpvOpRayQueryGetRayFlagsKHR: return "OpRayQueryGetRayFlagsKHR"; case SpvOpRayQueryGetIntersectionTKHR: return "OpRayQueryGetIntersectionTKHR"; @@ -5232,11 +5554,11 @@ inline const char* SpvOpToString(SpvOp value) { case SpvOpControlBarrierArriveINTEL: return "OpControlBarrierArriveINTEL"; case SpvOpControlBarrierWaitINTEL: return "OpControlBarrierWaitINTEL"; case SpvOpArithmeticFenceEXT: return "OpArithmeticFenceEXT"; - case SpvOpTaskSequenceCreateINTEL: return "OpTaskSequenceCreateINTEL"; - case SpvOpTaskSequenceAsyncINTEL: return "OpTaskSequenceAsyncINTEL"; - case SpvOpTaskSequenceGetINTEL: return "OpTaskSequenceGetINTEL"; - case SpvOpTaskSequenceReleaseINTEL: return "OpTaskSequenceReleaseINTEL"; - case SpvOpTypeTaskSequenceINTEL: return "OpTypeTaskSequenceINTEL"; + case SpvOpTaskSequenceCreateALTERA: return "OpTaskSequenceCreateALTERA"; + case SpvOpTaskSequenceAsyncALTERA: return "OpTaskSequenceAsyncALTERA"; + case SpvOpTaskSequenceGetALTERA: return "OpTaskSequenceGetALTERA"; + case SpvOpTaskSequenceReleaseALTERA: return "OpTaskSequenceReleaseALTERA"; + case SpvOpTypeTaskSequenceALTERA: return "OpTypeTaskSequenceALTERA"; case SpvOpSubgroupBlockPrefetchINTEL: return "OpSubgroupBlockPrefetchINTEL"; case SpvOpSubgroup2DBlockLoadINTEL: return "OpSubgroup2DBlockLoadINTEL"; case SpvOpSubgroup2DBlockLoadTransformINTEL: return "OpSubgroup2DBlockLoadTransformINTEL"; @@ -5267,6 +5589,9 @@ inline const char* SpvOpToString(SpvOp value) { case SpvOpConvertHandleToImageINTEL: return "OpConvertHandleToImageINTEL"; case SpvOpConvertHandleToSamplerINTEL: return "OpConvertHandleToSamplerINTEL"; case SpvOpConvertHandleToSampledImageINTEL: return "OpConvertHandleToSampledImageINTEL"; + case SpvOpFDot2MixAcc32VALVE: return "OpFDot2MixAcc32VALVE"; + case SpvOpFDot2MixAcc16VALVE: return "OpFDot2MixAcc16VALVE"; + case SpvOpFDot4MixAcc32VALVE: return "OpFDot4MixAcc32VALVE"; default: return "Unknown"; } } diff --git a/3rdparty/spirv-cross/spirv.hpp b/3rdparty/spirv-cross/spirv.hpp index 086fcc48c..462c3f6bf 100644 --- a/3rdparty/spirv-cross/spirv.hpp +++ b/3rdparty/spirv-cross/spirv.hpp @@ -175,6 +175,7 @@ enum ExecutionMode { ExecutionModeQuadDerivativesKHR = 5088, ExecutionModeRequireFullQuadsKHR = 5089, ExecutionModeSharesInputWithAMDX = 5102, + ExecutionModeArithmeticPoisonKHR = 5157, ExecutionModeOutputLinesEXT = 5269, ExecutionModeOutputLinesNV = 5269, ExecutionModeOutputPrimitivesEXT = 5270, @@ -191,6 +192,7 @@ enum ExecutionMode { ExecutionModeSampleInterlockUnorderedEXT = 5369, ExecutionModeShadingRateInterlockOrderedEXT = 5370, ExecutionModeShadingRateInterlockUnorderedEXT = 5371, + ExecutionModeShader64BitIndexingEXT = 5427, ExecutionModeSharedLocalMemorySizeINTEL = 5618, ExecutionModeRoundingModeRTPINTEL = 5620, ExecutionModeRoundingModeRTNINTEL = 5621, @@ -245,8 +247,11 @@ enum StorageClass { StorageClassPhysicalStorageBufferEXT = 5349, StorageClassHitObjectAttributeNV = 5385, StorageClassTaskPayloadWorkgroupEXT = 5402, + StorageClassHitObjectAttributeEXT = 5411, StorageClassCodeSectionINTEL = 5605, + StorageClassDeviceOnlyALTERA = 5936, StorageClassDeviceOnlyINTEL = 5936, + StorageClassHostOnlyALTERA = 5937, StorageClassHostOnlyINTEL = 5937, StorageClassMax = 0x7fffffff, }; @@ -485,6 +490,7 @@ enum FunctionParameterAttribute { FunctionParameterAttributeNoCapture = 5, FunctionParameterAttributeNoWrite = 6, FunctionParameterAttributeNoReadWrite = 7, + FunctionParameterAttributeRuntimeAlignedALTERA = 5940, FunctionParameterAttributeRuntimeAlignedINTEL = 5940, FunctionParameterAttributeMax = 0x7fffffff, }; @@ -552,6 +558,9 @@ enum Decoration { DecorationPayloadNodeSparseArrayAMDX = 5099, DecorationPayloadNodeArraySizeAMDX = 5100, DecorationPayloadDispatchIndirectAMDX = 5105, + DecorationArrayStrideIdEXT = 5124, + DecorationOffsetIdEXT = 5125, + DecorationUTFEncodedKHR = 5145, DecorationOverrideCoverageNV = 5248, DecorationPassthroughNV = 5250, DecorationViewportRelativeNV = 5252, @@ -568,7 +577,10 @@ enum Decoration { DecorationRestrictPointerEXT = 5355, DecorationAliasedPointer = 5356, DecorationAliasedPointerEXT = 5356, + DecorationMemberOffsetNV = 5358, DecorationHitObjectShaderRecordBufferNV = 5386, + DecorationHitObjectShaderRecordBufferEXT = 5389, + DecorationBankNV = 5397, DecorationBindlessSamplerNV = 5398, DecorationBindlessImageNV = 5399, DecorationBoundSamplerNV = 5400, @@ -589,54 +601,95 @@ enum Decoration { DecorationUserTypeGOOGLE = 5636, DecorationFunctionRoundingModeINTEL = 5822, DecorationFunctionDenormModeINTEL = 5823, + DecorationRegisterALTERA = 5825, DecorationRegisterINTEL = 5825, + DecorationMemoryALTERA = 5826, DecorationMemoryINTEL = 5826, + DecorationNumbanksALTERA = 5827, DecorationNumbanksINTEL = 5827, + DecorationBankwidthALTERA = 5828, DecorationBankwidthINTEL = 5828, + DecorationMaxPrivateCopiesALTERA = 5829, DecorationMaxPrivateCopiesINTEL = 5829, + DecorationSinglepumpALTERA = 5830, DecorationSinglepumpINTEL = 5830, + DecorationDoublepumpALTERA = 5831, DecorationDoublepumpINTEL = 5831, + DecorationMaxReplicatesALTERA = 5832, DecorationMaxReplicatesINTEL = 5832, + DecorationSimpleDualPortALTERA = 5833, DecorationSimpleDualPortINTEL = 5833, + DecorationMergeALTERA = 5834, DecorationMergeINTEL = 5834, + DecorationBankBitsALTERA = 5835, DecorationBankBitsINTEL = 5835, + DecorationForcePow2DepthALTERA = 5836, DecorationForcePow2DepthINTEL = 5836, + DecorationStridesizeALTERA = 5883, DecorationStridesizeINTEL = 5883, + DecorationWordsizeALTERA = 5884, DecorationWordsizeINTEL = 5884, + DecorationTrueDualPortALTERA = 5885, DecorationTrueDualPortINTEL = 5885, + DecorationBurstCoalesceALTERA = 5899, DecorationBurstCoalesceINTEL = 5899, + DecorationCacheSizeALTERA = 5900, DecorationCacheSizeINTEL = 5900, + DecorationDontStaticallyCoalesceALTERA = 5901, DecorationDontStaticallyCoalesceINTEL = 5901, + DecorationPrefetchALTERA = 5902, DecorationPrefetchINTEL = 5902, + DecorationStallEnableALTERA = 5905, DecorationStallEnableINTEL = 5905, + DecorationFuseLoopsInFunctionALTERA = 5907, DecorationFuseLoopsInFunctionINTEL = 5907, + DecorationMathOpDSPModeALTERA = 5909, DecorationMathOpDSPModeINTEL = 5909, DecorationAliasScopeINTEL = 5914, DecorationNoAliasINTEL = 5915, + DecorationInitiationIntervalALTERA = 5917, DecorationInitiationIntervalINTEL = 5917, + DecorationMaxConcurrencyALTERA = 5918, DecorationMaxConcurrencyINTEL = 5918, + DecorationPipelineEnableALTERA = 5919, DecorationPipelineEnableINTEL = 5919, + DecorationBufferLocationALTERA = 5921, DecorationBufferLocationINTEL = 5921, + DecorationIOPipeStorageALTERA = 5944, DecorationIOPipeStorageINTEL = 5944, DecorationFunctionFloatingPointModeINTEL = 6080, DecorationSingleElementVectorINTEL = 6085, DecorationVectorComputeCallableFunctionINTEL = 6087, DecorationMediaBlockIOINTEL = 6140, + DecorationStallFreeALTERA = 6151, DecorationStallFreeINTEL = 6151, DecorationFPMaxErrorDecorationINTEL = 6170, + DecorationLatencyControlLabelALTERA = 6172, DecorationLatencyControlLabelINTEL = 6172, + DecorationLatencyControlConstraintALTERA = 6173, DecorationLatencyControlConstraintINTEL = 6173, + DecorationConduitKernelArgumentALTERA = 6175, DecorationConduitKernelArgumentINTEL = 6175, + DecorationRegisterMapKernelArgumentALTERA = 6176, DecorationRegisterMapKernelArgumentINTEL = 6176, + DecorationMMHostInterfaceAddressWidthALTERA = 6177, DecorationMMHostInterfaceAddressWidthINTEL = 6177, + DecorationMMHostInterfaceDataWidthALTERA = 6178, DecorationMMHostInterfaceDataWidthINTEL = 6178, + DecorationMMHostInterfaceLatencyALTERA = 6179, DecorationMMHostInterfaceLatencyINTEL = 6179, + DecorationMMHostInterfaceReadWriteModeALTERA = 6180, DecorationMMHostInterfaceReadWriteModeINTEL = 6180, + DecorationMMHostInterfaceMaxBurstALTERA = 6181, DecorationMMHostInterfaceMaxBurstINTEL = 6181, + DecorationMMHostInterfaceWaitRequestALTERA = 6182, DecorationMMHostInterfaceWaitRequestINTEL = 6182, + DecorationStableKernelArgumentALTERA = 6183, DecorationStableKernelArgumentINTEL = 6183, DecorationHostAccessINTEL = 6188, + DecorationInitModeALTERA = 6190, DecorationInitModeINTEL = 6190, + DecorationImplementInRegisterMapALTERA = 6191, DecorationImplementInRegisterMapINTEL = 6191, DecorationConditionalINTEL = 6247, DecorationCacheControlLoadINTEL = 6442, @@ -721,6 +774,8 @@ enum BuiltIn { BuiltInFragStencilRefEXT = 5014, BuiltInRemainingRecursionLevelsAMDX = 5021, BuiltInShaderIndexAMDX = 5073, + BuiltInSamplerHeapEXT = 5122, + BuiltInResourceHeapEXT = 5123, BuiltInViewportMaskNV = 5253, BuiltInSecondaryPositionNV = 5257, BuiltInSecondaryViewportMaskNV = 5258, @@ -818,15 +873,25 @@ enum LoopControlShift { LoopControlIterationMultipleShift = 6, LoopControlPeelCountShift = 7, LoopControlPartialCountShift = 8, + LoopControlInitiationIntervalALTERAShift = 16, LoopControlInitiationIntervalINTELShift = 16, + LoopControlMaxConcurrencyALTERAShift = 17, LoopControlMaxConcurrencyINTELShift = 17, + LoopControlDependencyArrayALTERAShift = 18, LoopControlDependencyArrayINTELShift = 18, + LoopControlPipelineEnableALTERAShift = 19, LoopControlPipelineEnableINTELShift = 19, + LoopControlLoopCoalesceALTERAShift = 20, LoopControlLoopCoalesceINTELShift = 20, + LoopControlMaxInterleavingALTERAShift = 21, LoopControlMaxInterleavingINTELShift = 21, + LoopControlSpeculatedIterationsALTERAShift = 22, LoopControlSpeculatedIterationsINTELShift = 22, + LoopControlNoFusionALTERAShift = 23, LoopControlNoFusionINTELShift = 23, + LoopControlLoopCountALTERAShift = 24, LoopControlLoopCountINTELShift = 24, + LoopControlMaxReinvocationDelayALTERAShift = 25, LoopControlMaxReinvocationDelayINTELShift = 25, LoopControlMax = 0x7fffffff, }; @@ -842,15 +907,25 @@ enum LoopControlMask { LoopControlIterationMultipleMask = 0x00000040, LoopControlPeelCountMask = 0x00000080, LoopControlPartialCountMask = 0x00000100, + LoopControlInitiationIntervalALTERAMask = 0x00010000, LoopControlInitiationIntervalINTELMask = 0x00010000, + LoopControlMaxConcurrencyALTERAMask = 0x00020000, LoopControlMaxConcurrencyINTELMask = 0x00020000, + LoopControlDependencyArrayALTERAMask = 0x00040000, LoopControlDependencyArrayINTELMask = 0x00040000, + LoopControlPipelineEnableALTERAMask = 0x00080000, LoopControlPipelineEnableINTELMask = 0x00080000, + LoopControlLoopCoalesceALTERAMask = 0x00100000, LoopControlLoopCoalesceINTELMask = 0x00100000, + LoopControlMaxInterleavingALTERAMask = 0x00200000, LoopControlMaxInterleavingINTELMask = 0x00200000, + LoopControlSpeculatedIterationsALTERAMask = 0x00400000, LoopControlSpeculatedIterationsINTELMask = 0x00400000, + LoopControlNoFusionALTERAMask = 0x00800000, LoopControlNoFusionINTELMask = 0x00800000, + LoopControlLoopCountALTERAMask = 0x01000000, LoopControlLoopCountINTELMask = 0x01000000, + LoopControlMaxReinvocationDelayALTERAMask = 0x02000000, LoopControlMaxReinvocationDelayINTELMask = 0x02000000, }; @@ -963,8 +1038,11 @@ enum GroupOperation { GroupOperationInclusiveScan = 1, GroupOperationExclusiveScan = 2, GroupOperationClusteredReduce = 3, + GroupOperationPartitionedReduceEXT = 6, GroupOperationPartitionedReduceNV = 6, + GroupOperationPartitionedInclusiveScanEXT = 7, GroupOperationPartitionedInclusiveScanNV = 7, + GroupOperationPartitionedExclusiveScanEXT = 8, GroupOperationPartitionedExclusiveScanNV = 8, GroupOperationMax = 0x7fffffff, }; @@ -1120,6 +1198,10 @@ enum Capability { CapabilityBFloat16TypeKHR = 5116, CapabilityBFloat16DotProductKHR = 5117, CapabilityBFloat16CooperativeMatrixKHR = 5118, + CapabilityAbortKHR = 5120, + CapabilityDescriptorHeapEXT = 5128, + CapabilityConstantDataKHR = 5146, + CapabilityPoisonFreezeKHR = 5156, CapabilitySampleMaskOverrideCoverageNV = 5249, CapabilityGeometryShaderPassthroughNV = 5251, CapabilityShaderViewportIndexLayerEXT = 5254, @@ -1137,6 +1219,7 @@ enum Capability { CapabilityComputeDerivativeGroupQuadsNV = 5288, CapabilityFragmentDensityEXT = 5291, CapabilityShadingRateNV = 5291, + CapabilityGroupNonUniformPartitionedEXT = 5297, CapabilityGroupNonUniformPartitionedNV = 5297, CapabilityShaderNonUniform = 5301, CapabilityShaderNonUniformEXT = 5301, @@ -1184,6 +1267,7 @@ enum Capability { CapabilityDisplacementMicromapNV = 5380, CapabilityRayTracingOpacityMicromapEXT = 5381, CapabilityShaderInvocationReorderNV = 5383, + CapabilityShaderInvocationReorderEXT = 5388, CapabilityBindlessTextureNV = 5390, CapabilityRayQueryPositionFetchKHR = 5391, CapabilityCooperativeVectorNV = 5394, @@ -1192,6 +1276,9 @@ enum Capability { CapabilityRawAccessChainsNV = 5414, CapabilityRayTracingSpheresGeometryNV = 5418, CapabilityRayTracingLinearSweptSpheresGeometryNV = 5419, + CapabilityPushConstantBanksNV = 5423, + CapabilityLongVectorEXT = 5425, + CapabilityShader64BitIndexingEXT = 5426, CapabilityCooperativeMatrixReductionsNV = 5430, CapabilityCooperativeMatrixConversionsNV = 5431, CapabilityCooperativeMatrixPerElementOperationsNV = 5432, @@ -1221,26 +1308,42 @@ enum Capability { CapabilitySubgroupAvcMotionEstimationChromaINTEL = 5698, CapabilityVariableLengthArrayINTEL = 5817, CapabilityFunctionFloatControlINTEL = 5821, + CapabilityFPGAMemoryAttributesALTERA = 5824, CapabilityFPGAMemoryAttributesINTEL = 5824, CapabilityFPFastMathModeINTEL = 5837, + CapabilityArbitraryPrecisionIntegersALTERA = 5844, CapabilityArbitraryPrecisionIntegersINTEL = 5844, + CapabilityArbitraryPrecisionFloatingPointALTERA = 5845, CapabilityArbitraryPrecisionFloatingPointINTEL = 5845, CapabilityUnstructuredLoopControlsINTEL = 5886, + CapabilityFPGALoopControlsALTERA = 5888, CapabilityFPGALoopControlsINTEL = 5888, CapabilityKernelAttributesINTEL = 5892, CapabilityFPGAKernelAttributesINTEL = 5897, + CapabilityFPGAMemoryAccessesALTERA = 5898, CapabilityFPGAMemoryAccessesINTEL = 5898, + CapabilityFPGAClusterAttributesALTERA = 5904, CapabilityFPGAClusterAttributesINTEL = 5904, + CapabilityLoopFuseALTERA = 5906, CapabilityLoopFuseINTEL = 5906, + CapabilityFPGADSPControlALTERA = 5908, CapabilityFPGADSPControlINTEL = 5908, CapabilityMemoryAccessAliasingINTEL = 5910, + CapabilityFPGAInvocationPipeliningAttributesALTERA = 5916, CapabilityFPGAInvocationPipeliningAttributesINTEL = 5916, + CapabilityFPGABufferLocationALTERA = 5920, CapabilityFPGABufferLocationINTEL = 5920, + CapabilityArbitraryPrecisionFixedPointALTERA = 5922, CapabilityArbitraryPrecisionFixedPointINTEL = 5922, + CapabilityUSMStorageClassesALTERA = 5935, CapabilityUSMStorageClassesINTEL = 5935, + CapabilityRuntimeAlignedAttributeALTERA = 5939, CapabilityRuntimeAlignedAttributeINTEL = 5939, + CapabilityIOPipesALTERA = 5943, CapabilityIOPipesINTEL = 5943, + CapabilityBlockingPipesALTERA = 5945, CapabilityBlockingPipesINTEL = 5945, + CapabilityFPGARegALTERA = 5948, CapabilityFPGARegINTEL = 5948, CapabilityDotProductInputAll = 6016, CapabilityDotProductInputAllKHR = 6016, @@ -1267,13 +1370,18 @@ enum Capability { CapabilityBFloat16ConversionINTEL = 6115, CapabilitySplitBarrierINTEL = 6141, CapabilityArithmeticFenceEXT = 6144, + CapabilityFPGAClusterAttributesV2ALTERA = 6150, CapabilityFPGAClusterAttributesV2INTEL = 6150, CapabilityFPGAKernelAttributesv2INTEL = 6161, + CapabilityTaskSequenceALTERA = 6162, CapabilityTaskSequenceINTEL = 6162, CapabilityFPMaxErrorINTEL = 6169, + CapabilityFPGALatencyControlALTERA = 6171, CapabilityFPGALatencyControlINTEL = 6171, + CapabilityFPGAArgumentInterfacesALTERA = 6174, CapabilityFPGAArgumentInterfacesINTEL = 6174, CapabilityGlobalVariableHostAccessINTEL = 6187, + CapabilityGlobalVariableFPGADecorationsALTERA = 6189, CapabilityGlobalVariableFPGADecorationsINTEL = 6189, CapabilitySubgroupBufferPrefetchINTEL = 6220, CapabilitySubgroup2DBlockIOINTEL = 6228, @@ -1290,6 +1398,10 @@ enum Capability { CapabilityCacheControlsINTEL = 6441, CapabilityRegisterLimitsINTEL = 6460, CapabilityBindlessImagesINTEL = 6528, + CapabilityDotProductFloat16AccFloat32VALVE = 6912, + CapabilityDotProductFloat16AccFloat16VALVE = 6913, + CapabilityDotProductBFloat16AccVALVE = 6914, + CapabilityDotProductFloat8AccFloat32VALVE = 6915, CapabilityMax = 0x7fffffff, }; @@ -1485,7 +1597,9 @@ enum TensorOperandsMask { }; enum InitializationModeQualifier { + InitializationModeQualifierInitOnDeviceReprogramALTERA = 0, InitializationModeQualifierInitOnDeviceReprogramINTEL = 0, + InitializationModeQualifierInitOnDeviceResetALTERA = 1, InitializationModeQualifierInitOnDeviceResetINTEL = 1, InitializationModeQualifierMax = 0x7fffffff, }; @@ -2045,6 +2159,16 @@ enum Op { OpSpecConstantStringAMDX = 5104, OpGroupNonUniformQuadAllKHR = 5110, OpGroupNonUniformQuadAnyKHR = 5111, + OpTypeBufferEXT = 5115, + OpBufferPointerEXT = 5119, + OpAbortKHR = 5121, + OpUntypedImageTexelPointerEXT = 5126, + OpMemberDecorateIdEXT = 5127, + OpConstantSizeOfEXT = 5129, + OpConstantDataKHR = 5147, + OpSpecConstantDataKHR = 5148, + OpPoisonKHR = 5158, + OpFreezeKHR = 5159, OpHitObjectRecordHitMotionNV = 5249, OpHitObjectRecordHitWithIndexMotionNV = 5250, OpHitObjectRecordMissMotionNV = 5251, @@ -2080,6 +2204,7 @@ enum Op { OpTypeHitObjectNV = 5281, OpImageSampleFootprintNV = 5283, OpTypeCooperativeVectorNV = 5288, + OpTypeVectorIdEXT = 5288, OpCooperativeVectorMatrixMulNV = 5289, OpCooperativeVectorOuterProductAccumulateNV = 5290, OpCooperativeVectorReduceSumAccumulateNV = 5291, @@ -2087,12 +2212,43 @@ enum Op { OpCooperativeMatrixConvertNV = 5293, OpEmitMeshTasksEXT = 5294, OpSetMeshOutputsEXT = 5295, + OpGroupNonUniformPartitionEXT = 5296, OpGroupNonUniformPartitionNV = 5296, OpWritePackedPrimitiveIndices4x8NV = 5299, OpFetchMicroTriangleVertexPositionNV = 5300, OpFetchMicroTriangleVertexBarycentricNV = 5301, OpCooperativeVectorLoadNV = 5302, OpCooperativeVectorStoreNV = 5303, + OpHitObjectRecordFromQueryEXT = 5304, + OpHitObjectRecordMissEXT = 5305, + OpHitObjectRecordMissMotionEXT = 5306, + OpHitObjectGetIntersectionTriangleVertexPositionsEXT = 5307, + OpHitObjectGetRayFlagsEXT = 5308, + OpHitObjectSetShaderBindingTableRecordIndexEXT = 5309, + OpHitObjectReorderExecuteShaderEXT = 5310, + OpHitObjectTraceReorderExecuteEXT = 5311, + OpHitObjectTraceMotionReorderExecuteEXT = 5312, + OpTypeHitObjectEXT = 5313, + OpReorderThreadWithHintEXT = 5314, + OpReorderThreadWithHitObjectEXT = 5315, + OpHitObjectTraceRayEXT = 5316, + OpHitObjectTraceRayMotionEXT = 5317, + OpHitObjectRecordEmptyEXT = 5318, + OpHitObjectExecuteShaderEXT = 5319, + OpHitObjectGetCurrentTimeEXT = 5320, + OpHitObjectGetAttributesEXT = 5321, + OpHitObjectGetHitKindEXT = 5322, + OpHitObjectGetPrimitiveIndexEXT = 5323, + OpHitObjectGetGeometryIndexEXT = 5324, + OpHitObjectGetInstanceIdEXT = 5325, + OpHitObjectGetInstanceCustomIndexEXT = 5326, + OpHitObjectGetObjectRayOriginEXT = 5327, + OpHitObjectGetObjectRayDirectionEXT = 5328, + OpHitObjectGetWorldRayDirectionEXT = 5329, + OpHitObjectGetWorldRayOriginEXT = 5330, + OpHitObjectGetObjectToWorldEXT = 5331, + OpHitObjectGetWorldToObjectEXT = 5332, + OpHitObjectGetRayTMaxEXT = 5333, OpReportIntersectionKHR = 5334, OpReportIntersectionNV = 5334, OpIgnoreIntersectionNV = 5335, @@ -2107,6 +2263,12 @@ enum Op { OpRayQueryGetClusterIdNV = 5345, OpRayQueryGetIntersectionClusterIdNV = 5345, OpHitObjectGetClusterIdNV = 5346, + OpHitObjectGetRayTMinEXT = 5347, + OpHitObjectGetShaderBindingTableRecordIndexEXT = 5348, + OpHitObjectGetShaderRecordBufferHandleEXT = 5349, + OpHitObjectIsEmptyEXT = 5350, + OpHitObjectIsHitEXT = 5351, + OpHitObjectIsMissEXT = 5352, OpTypeCooperativeMatrixNV = 5358, OpCooperativeMatrixLoadNV = 5359, OpCooperativeMatrixStoreNV = 5360, @@ -2313,23 +2475,41 @@ enum Op { OpVariableLengthArrayINTEL = 5818, OpSaveMemoryINTEL = 5819, OpRestoreMemoryINTEL = 5820, + OpArbitraryFloatSinCosPiALTERA = 5840, OpArbitraryFloatSinCosPiINTEL = 5840, + OpArbitraryFloatCastALTERA = 5841, OpArbitraryFloatCastINTEL = 5841, + OpArbitraryFloatCastFromIntALTERA = 5842, OpArbitraryFloatCastFromIntINTEL = 5842, + OpArbitraryFloatCastToIntALTERA = 5843, OpArbitraryFloatCastToIntINTEL = 5843, + OpArbitraryFloatAddALTERA = 5846, OpArbitraryFloatAddINTEL = 5846, + OpArbitraryFloatSubALTERA = 5847, OpArbitraryFloatSubINTEL = 5847, + OpArbitraryFloatMulALTERA = 5848, OpArbitraryFloatMulINTEL = 5848, + OpArbitraryFloatDivALTERA = 5849, OpArbitraryFloatDivINTEL = 5849, + OpArbitraryFloatGTALTERA = 5850, OpArbitraryFloatGTINTEL = 5850, + OpArbitraryFloatGEALTERA = 5851, OpArbitraryFloatGEINTEL = 5851, + OpArbitraryFloatLTALTERA = 5852, OpArbitraryFloatLTINTEL = 5852, + OpArbitraryFloatLEALTERA = 5853, OpArbitraryFloatLEINTEL = 5853, + OpArbitraryFloatEQALTERA = 5854, OpArbitraryFloatEQINTEL = 5854, + OpArbitraryFloatRecipALTERA = 5855, OpArbitraryFloatRecipINTEL = 5855, + OpArbitraryFloatRSqrtALTERA = 5856, OpArbitraryFloatRSqrtINTEL = 5856, + OpArbitraryFloatCbrtALTERA = 5857, OpArbitraryFloatCbrtINTEL = 5857, + OpArbitraryFloatHypotALTERA = 5858, OpArbitraryFloatHypotINTEL = 5858, + OpArbitraryFloatSqrtALTERA = 5859, OpArbitraryFloatSqrtINTEL = 5859, OpArbitraryFloatLogINTEL = 5860, OpArbitraryFloatLog2INTEL = 5861, @@ -2358,21 +2538,37 @@ enum Op { OpAliasDomainDeclINTEL = 5911, OpAliasScopeDeclINTEL = 5912, OpAliasScopeListDeclINTEL = 5913, + OpFixedSqrtALTERA = 5923, OpFixedSqrtINTEL = 5923, + OpFixedRecipALTERA = 5924, OpFixedRecipINTEL = 5924, + OpFixedRsqrtALTERA = 5925, OpFixedRsqrtINTEL = 5925, + OpFixedSinALTERA = 5926, OpFixedSinINTEL = 5926, + OpFixedCosALTERA = 5927, OpFixedCosINTEL = 5927, + OpFixedSinCosALTERA = 5928, OpFixedSinCosINTEL = 5928, + OpFixedSinPiALTERA = 5929, OpFixedSinPiINTEL = 5929, + OpFixedCosPiALTERA = 5930, OpFixedCosPiINTEL = 5930, + OpFixedSinCosPiALTERA = 5931, OpFixedSinCosPiINTEL = 5931, + OpFixedLogALTERA = 5932, OpFixedLogINTEL = 5932, + OpFixedExpALTERA = 5933, OpFixedExpINTEL = 5933, + OpPtrCastToCrossWorkgroupALTERA = 5934, OpPtrCastToCrossWorkgroupINTEL = 5934, + OpCrossWorkgroupCastToPtrALTERA = 5938, OpCrossWorkgroupCastToPtrINTEL = 5938, + OpReadPipeBlockingALTERA = 5946, OpReadPipeBlockingINTEL = 5946, + OpWritePipeBlockingALTERA = 5947, OpWritePipeBlockingINTEL = 5947, + OpFPGARegALTERA = 5949, OpFPGARegINTEL = 5949, OpRayQueryGetRayTMinKHR = 6016, OpRayQueryGetRayFlagsKHR = 6017, @@ -2402,10 +2598,15 @@ enum Op { OpControlBarrierArriveINTEL = 6142, OpControlBarrierWaitINTEL = 6143, OpArithmeticFenceEXT = 6145, + OpTaskSequenceCreateALTERA = 6163, OpTaskSequenceCreateINTEL = 6163, + OpTaskSequenceAsyncALTERA = 6164, OpTaskSequenceAsyncINTEL = 6164, + OpTaskSequenceGetALTERA = 6165, OpTaskSequenceGetINTEL = 6165, + OpTaskSequenceReleaseALTERA = 6166, OpTaskSequenceReleaseINTEL = 6166, + OpTypeTaskSequenceALTERA = 6199, OpTypeTaskSequenceINTEL = 6199, OpSubgroupBlockPrefetchINTEL = 6221, OpSubgroup2DBlockLoadINTEL = 6231, @@ -2437,6 +2638,9 @@ enum Op { OpConvertHandleToImageINTEL = 6529, OpConvertHandleToSamplerINTEL = 6530, OpConvertHandleToSampledImageINTEL = 6531, + OpFDot2MixAcc32VALVE = 6916, + OpFDot2MixAcc16VALVE = 6917, + OpFDot4MixAcc32VALVE = 6918, OpMax = 0x7fffffff, }; @@ -2884,6 +3088,16 @@ inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) { case OpSpecConstantStringAMDX: *hasResult = true; *hasResultType = false; break; case OpGroupNonUniformQuadAllKHR: *hasResult = true; *hasResultType = true; break; case OpGroupNonUniformQuadAnyKHR: *hasResult = true; *hasResultType = true; break; + case OpTypeBufferEXT: *hasResult = true; *hasResultType = false; break; + case OpBufferPointerEXT: *hasResult = true; *hasResultType = true; break; + case OpAbortKHR: *hasResult = false; *hasResultType = false; break; + case OpUntypedImageTexelPointerEXT: *hasResult = true; *hasResultType = true; break; + case OpMemberDecorateIdEXT: *hasResult = false; *hasResultType = false; break; + case OpConstantSizeOfEXT: *hasResult = true; *hasResultType = true; break; + case OpConstantDataKHR: *hasResult = true; *hasResultType = true; break; + case OpSpecConstantDataKHR: *hasResult = true; *hasResultType = true; break; + case OpPoisonKHR: *hasResult = true; *hasResultType = true; break; + case OpFreezeKHR: *hasResult = true; *hasResultType = true; break; case OpHitObjectRecordHitMotionNV: *hasResult = false; *hasResultType = false; break; case OpHitObjectRecordHitWithIndexMotionNV: *hasResult = false; *hasResultType = false; break; case OpHitObjectRecordMissMotionNV: *hasResult = false; *hasResultType = false; break; @@ -2918,7 +3132,7 @@ inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) { case OpReorderThreadWithHintNV: *hasResult = false; *hasResultType = false; break; case OpTypeHitObjectNV: *hasResult = true; *hasResultType = false; break; case OpImageSampleFootprintNV: *hasResult = true; *hasResultType = true; break; - case OpTypeCooperativeVectorNV: *hasResult = true; *hasResultType = false; break; + case OpTypeVectorIdEXT: *hasResult = true; *hasResultType = false; break; case OpCooperativeVectorMatrixMulNV: *hasResult = true; *hasResultType = true; break; case OpCooperativeVectorOuterProductAccumulateNV: *hasResult = false; *hasResultType = false; break; case OpCooperativeVectorReduceSumAccumulateNV: *hasResult = false; *hasResultType = false; break; @@ -2926,12 +3140,42 @@ inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) { case OpCooperativeMatrixConvertNV: *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 OpGroupNonUniformPartitionEXT: *hasResult = true; *hasResultType = true; break; case OpWritePackedPrimitiveIndices4x8NV: *hasResult = false; *hasResultType = false; break; case OpFetchMicroTriangleVertexPositionNV: *hasResult = true; *hasResultType = true; break; case OpFetchMicroTriangleVertexBarycentricNV: *hasResult = true; *hasResultType = true; break; case OpCooperativeVectorLoadNV: *hasResult = true; *hasResultType = true; break; case OpCooperativeVectorStoreNV: *hasResult = false; *hasResultType = false; break; + case OpHitObjectRecordFromQueryEXT: *hasResult = false; *hasResultType = false; break; + case OpHitObjectRecordMissEXT: *hasResult = false; *hasResultType = false; break; + case OpHitObjectRecordMissMotionEXT: *hasResult = false; *hasResultType = false; break; + case OpHitObjectGetIntersectionTriangleVertexPositionsEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectGetRayFlagsEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectSetShaderBindingTableRecordIndexEXT: *hasResult = false; *hasResultType = false; break; + case OpHitObjectReorderExecuteShaderEXT: *hasResult = false; *hasResultType = false; break; + case OpHitObjectTraceReorderExecuteEXT: *hasResult = false; *hasResultType = false; break; + case OpHitObjectTraceMotionReorderExecuteEXT: *hasResult = false; *hasResultType = false; break; + case OpTypeHitObjectEXT: *hasResult = true; *hasResultType = false; break; + case OpReorderThreadWithHintEXT: *hasResult = false; *hasResultType = false; break; + case OpReorderThreadWithHitObjectEXT: *hasResult = false; *hasResultType = false; break; + case OpHitObjectTraceRayEXT: *hasResult = false; *hasResultType = false; break; + case OpHitObjectTraceRayMotionEXT: *hasResult = false; *hasResultType = false; break; + case OpHitObjectRecordEmptyEXT: *hasResult = false; *hasResultType = false; break; + case OpHitObjectExecuteShaderEXT: *hasResult = false; *hasResultType = false; break; + case OpHitObjectGetCurrentTimeEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectGetAttributesEXT: *hasResult = false; *hasResultType = false; break; + case OpHitObjectGetHitKindEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectGetPrimitiveIndexEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectGetGeometryIndexEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectGetInstanceIdEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectGetInstanceCustomIndexEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectGetObjectRayOriginEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectGetObjectRayDirectionEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectGetWorldRayDirectionEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectGetWorldRayOriginEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectGetObjectToWorldEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectGetWorldToObjectEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectGetRayTMaxEXT: *hasResult = true; *hasResultType = true; break; case OpReportIntersectionKHR: *hasResult = true; *hasResultType = true; break; case OpIgnoreIntersectionNV: *hasResult = false; *hasResultType = false; break; case OpTerminateRayNV: *hasResult = false; *hasResultType = false; break; @@ -2943,6 +3187,12 @@ inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) { case OpExecuteCallableNV: *hasResult = false; *hasResultType = false; break; case OpRayQueryGetIntersectionClusterIdNV: *hasResult = true; *hasResultType = true; break; case OpHitObjectGetClusterIdNV: *hasResult = true; *hasResultType = true; break; + case OpHitObjectGetRayTMinEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectGetShaderBindingTableRecordIndexEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectGetShaderRecordBufferHandleEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectIsEmptyEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectIsHitEXT: *hasResult = true; *hasResultType = true; break; + case OpHitObjectIsMissEXT: *hasResult = true; *hasResultType = true; break; case OpTypeCooperativeMatrixNV: *hasResult = true; *hasResultType = false; break; case OpCooperativeMatrixLoadNV: *hasResult = true; *hasResultType = true; break; case OpCooperativeMatrixStoreNV: *hasResult = false; *hasResultType = false; break; @@ -3146,24 +3396,24 @@ inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) { 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 OpArbitraryFloatSinCosPiALTERA: *hasResult = true; *hasResultType = true; break; + case OpArbitraryFloatCastALTERA: *hasResult = true; *hasResultType = true; break; + case OpArbitraryFloatCastFromIntALTERA: *hasResult = true; *hasResultType = true; break; + case OpArbitraryFloatCastToIntALTERA: *hasResult = true; *hasResultType = true; break; + case OpArbitraryFloatAddALTERA: *hasResult = true; *hasResultType = true; break; + case OpArbitraryFloatSubALTERA: *hasResult = true; *hasResultType = true; break; + case OpArbitraryFloatMulALTERA: *hasResult = true; *hasResultType = true; break; + case OpArbitraryFloatDivALTERA: *hasResult = true; *hasResultType = true; break; + case OpArbitraryFloatGTALTERA: *hasResult = true; *hasResultType = true; break; + case OpArbitraryFloatGEALTERA: *hasResult = true; *hasResultType = true; break; + case OpArbitraryFloatLTALTERA: *hasResult = true; *hasResultType = true; break; + case OpArbitraryFloatLEALTERA: *hasResult = true; *hasResultType = true; break; + case OpArbitraryFloatEQALTERA: *hasResult = true; *hasResultType = true; break; + case OpArbitraryFloatRecipALTERA: *hasResult = true; *hasResultType = true; break; + case OpArbitraryFloatRSqrtALTERA: *hasResult = true; *hasResultType = true; break; + case OpArbitraryFloatCbrtALTERA: *hasResult = true; *hasResultType = true; break; + case OpArbitraryFloatHypotALTERA: *hasResult = true; *hasResultType = true; break; + case OpArbitraryFloatSqrtALTERA: *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; @@ -3191,22 +3441,22 @@ inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) { 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; + case OpFixedSqrtALTERA: *hasResult = true; *hasResultType = true; break; + case OpFixedRecipALTERA: *hasResult = true; *hasResultType = true; break; + case OpFixedRsqrtALTERA: *hasResult = true; *hasResultType = true; break; + case OpFixedSinALTERA: *hasResult = true; *hasResultType = true; break; + case OpFixedCosALTERA: *hasResult = true; *hasResultType = true; break; + case OpFixedSinCosALTERA: *hasResult = true; *hasResultType = true; break; + case OpFixedSinPiALTERA: *hasResult = true; *hasResultType = true; break; + case OpFixedCosPiALTERA: *hasResult = true; *hasResultType = true; break; + case OpFixedSinCosPiALTERA: *hasResult = true; *hasResultType = true; break; + case OpFixedLogALTERA: *hasResult = true; *hasResultType = true; break; + case OpFixedExpALTERA: *hasResult = true; *hasResultType = true; break; + case OpPtrCastToCrossWorkgroupALTERA: *hasResult = true; *hasResultType = true; break; + case OpCrossWorkgroupCastToPtrALTERA: *hasResult = true; *hasResultType = true; break; + case OpReadPipeBlockingALTERA: *hasResult = true; *hasResultType = true; break; + case OpWritePipeBlockingALTERA: *hasResult = true; *hasResultType = true; break; + case OpFPGARegALTERA: *hasResult = true; *hasResultType = true; break; case OpRayQueryGetRayTMinKHR: *hasResult = true; *hasResultType = true; break; case OpRayQueryGetRayFlagsKHR: *hasResult = true; *hasResultType = true; break; case OpRayQueryGetIntersectionTKHR: *hasResult = true; *hasResultType = true; break; @@ -3235,11 +3485,11 @@ inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) { case OpControlBarrierArriveINTEL: *hasResult = false; *hasResultType = false; break; case OpControlBarrierWaitINTEL: *hasResult = false; *hasResultType = false; break; case OpArithmeticFenceEXT: *hasResult = true; *hasResultType = true; break; - case OpTaskSequenceCreateINTEL: *hasResult = true; *hasResultType = true; break; - case OpTaskSequenceAsyncINTEL: *hasResult = false; *hasResultType = false; break; - case OpTaskSequenceGetINTEL: *hasResult = true; *hasResultType = true; break; - case OpTaskSequenceReleaseINTEL: *hasResult = false; *hasResultType = false; break; - case OpTypeTaskSequenceINTEL: *hasResult = true; *hasResultType = false; break; + case OpTaskSequenceCreateALTERA: *hasResult = true; *hasResultType = true; break; + case OpTaskSequenceAsyncALTERA: *hasResult = false; *hasResultType = false; break; + case OpTaskSequenceGetALTERA: *hasResult = true; *hasResultType = true; break; + case OpTaskSequenceReleaseALTERA: *hasResult = false; *hasResultType = false; break; + case OpTypeTaskSequenceALTERA: *hasResult = true; *hasResultType = false; break; case OpSubgroupBlockPrefetchINTEL: *hasResult = false; *hasResultType = false; break; case OpSubgroup2DBlockLoadINTEL: *hasResult = false; *hasResultType = false; break; case OpSubgroup2DBlockLoadTransformINTEL: *hasResult = false; *hasResultType = false; break; @@ -3270,6 +3520,9 @@ inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) { case OpConvertHandleToImageINTEL: *hasResult = true; *hasResultType = true; break; case OpConvertHandleToSamplerINTEL: *hasResult = true; *hasResultType = true; break; case OpConvertHandleToSampledImageINTEL: *hasResult = true; *hasResultType = true; break; + case OpFDot2MixAcc32VALVE: *hasResult = true; *hasResultType = true; break; + case OpFDot2MixAcc16VALVE: *hasResult = true; *hasResultType = true; break; + case OpFDot4MixAcc32VALVE: *hasResult = true; *hasResultType = true; break; } } inline const char* SourceLanguageToString(SourceLanguage value) { @@ -3404,6 +3657,7 @@ inline const char* ExecutionModeToString(ExecutionMode value) { case ExecutionModeQuadDerivativesKHR: return "QuadDerivativesKHR"; case ExecutionModeRequireFullQuadsKHR: return "RequireFullQuadsKHR"; case ExecutionModeSharesInputWithAMDX: return "SharesInputWithAMDX"; + case ExecutionModeArithmeticPoisonKHR: return "ArithmeticPoisonKHR"; case ExecutionModeOutputLinesEXT: return "OutputLinesEXT"; case ExecutionModeOutputPrimitivesEXT: return "OutputPrimitivesEXT"; case ExecutionModeDerivativeGroupQuadsKHR: return "DerivativeGroupQuadsKHR"; @@ -3415,6 +3669,7 @@ inline const char* ExecutionModeToString(ExecutionMode value) { case ExecutionModeSampleInterlockUnorderedEXT: return "SampleInterlockUnorderedEXT"; case ExecutionModeShadingRateInterlockOrderedEXT: return "ShadingRateInterlockOrderedEXT"; case ExecutionModeShadingRateInterlockUnorderedEXT: return "ShadingRateInterlockUnorderedEXT"; + case ExecutionModeShader64BitIndexingEXT: return "Shader64BitIndexingEXT"; case ExecutionModeSharedLocalMemorySizeINTEL: return "SharedLocalMemorySizeINTEL"; case ExecutionModeRoundingModeRTPINTEL: return "RoundingModeRTPINTEL"; case ExecutionModeRoundingModeRTNINTEL: return "RoundingModeRTNINTEL"; @@ -3464,9 +3719,10 @@ inline const char* StorageClassToString(StorageClass value) { case StorageClassPhysicalStorageBuffer: return "PhysicalStorageBuffer"; case StorageClassHitObjectAttributeNV: return "HitObjectAttributeNV"; case StorageClassTaskPayloadWorkgroupEXT: return "TaskPayloadWorkgroupEXT"; + case StorageClassHitObjectAttributeEXT: return "HitObjectAttributeEXT"; case StorageClassCodeSectionINTEL: return "CodeSectionINTEL"; - case StorageClassDeviceOnlyINTEL: return "DeviceOnlyINTEL"; - case StorageClassHostOnlyINTEL: return "HostOnlyINTEL"; + case StorageClassDeviceOnlyALTERA: return "DeviceOnlyALTERA"; + case StorageClassHostOnlyALTERA: return "HostOnlyALTERA"; default: return "Unknown"; } } @@ -3648,7 +3904,7 @@ inline const char* FunctionParameterAttributeToString(FunctionParameterAttribute case FunctionParameterAttributeNoCapture: return "NoCapture"; case FunctionParameterAttributeNoWrite: return "NoWrite"; case FunctionParameterAttributeNoReadWrite: return "NoReadWrite"; - case FunctionParameterAttributeRuntimeAlignedINTEL: return "RuntimeAlignedINTEL"; + case FunctionParameterAttributeRuntimeAlignedALTERA: return "RuntimeAlignedALTERA"; default: return "Unknown"; } } @@ -3717,6 +3973,9 @@ inline const char* DecorationToString(Decoration value) { case DecorationPayloadNodeSparseArrayAMDX: return "PayloadNodeSparseArrayAMDX"; case DecorationPayloadNodeArraySizeAMDX: return "PayloadNodeArraySizeAMDX"; case DecorationPayloadDispatchIndirectAMDX: return "PayloadDispatchIndirectAMDX"; + case DecorationArrayStrideIdEXT: return "ArrayStrideIdEXT"; + case DecorationOffsetIdEXT: return "OffsetIdEXT"; + case DecorationUTFEncodedKHR: return "UTFEncodedKHR"; case DecorationOverrideCoverageNV: return "OverrideCoverageNV"; case DecorationPassthroughNV: return "PassthroughNV"; case DecorationViewportRelativeNV: return "ViewportRelativeNV"; @@ -3728,7 +3987,10 @@ inline const char* DecorationToString(Decoration value) { case DecorationNonUniform: return "NonUniform"; case DecorationRestrictPointer: return "RestrictPointer"; case DecorationAliasedPointer: return "AliasedPointer"; + case DecorationMemberOffsetNV: return "MemberOffsetNV"; case DecorationHitObjectShaderRecordBufferNV: return "HitObjectShaderRecordBufferNV"; + case DecorationHitObjectShaderRecordBufferEXT: return "HitObjectShaderRecordBufferEXT"; + case DecorationBankNV: return "BankNV"; case DecorationBindlessSamplerNV: return "BindlessSamplerNV"; case DecorationBindlessImageNV: return "BindlessImageNV"; case DecorationBoundSamplerNV: return "BoundSamplerNV"; @@ -3747,55 +4009,55 @@ inline const char* DecorationToString(Decoration value) { case DecorationUserTypeGOOGLE: return "UserTypeGOOGLE"; case DecorationFunctionRoundingModeINTEL: return "FunctionRoundingModeINTEL"; case DecorationFunctionDenormModeINTEL: return "FunctionDenormModeINTEL"; - case DecorationRegisterINTEL: return "RegisterINTEL"; - case DecorationMemoryINTEL: return "MemoryINTEL"; - case DecorationNumbanksINTEL: return "NumbanksINTEL"; - case DecorationBankwidthINTEL: return "BankwidthINTEL"; - case DecorationMaxPrivateCopiesINTEL: return "MaxPrivateCopiesINTEL"; - case DecorationSinglepumpINTEL: return "SinglepumpINTEL"; - case DecorationDoublepumpINTEL: return "DoublepumpINTEL"; - case DecorationMaxReplicatesINTEL: return "MaxReplicatesINTEL"; - case DecorationSimpleDualPortINTEL: return "SimpleDualPortINTEL"; - case DecorationMergeINTEL: return "MergeINTEL"; - case DecorationBankBitsINTEL: return "BankBitsINTEL"; - case DecorationForcePow2DepthINTEL: return "ForcePow2DepthINTEL"; - case DecorationStridesizeINTEL: return "StridesizeINTEL"; - case DecorationWordsizeINTEL: return "WordsizeINTEL"; - case DecorationTrueDualPortINTEL: return "TrueDualPortINTEL"; - case DecorationBurstCoalesceINTEL: return "BurstCoalesceINTEL"; - case DecorationCacheSizeINTEL: return "CacheSizeINTEL"; - case DecorationDontStaticallyCoalesceINTEL: return "DontStaticallyCoalesceINTEL"; - case DecorationPrefetchINTEL: return "PrefetchINTEL"; - case DecorationStallEnableINTEL: return "StallEnableINTEL"; - case DecorationFuseLoopsInFunctionINTEL: return "FuseLoopsInFunctionINTEL"; - case DecorationMathOpDSPModeINTEL: return "MathOpDSPModeINTEL"; + case DecorationRegisterALTERA: return "RegisterALTERA"; + case DecorationMemoryALTERA: return "MemoryALTERA"; + case DecorationNumbanksALTERA: return "NumbanksALTERA"; + case DecorationBankwidthALTERA: return "BankwidthALTERA"; + case DecorationMaxPrivateCopiesALTERA: return "MaxPrivateCopiesALTERA"; + case DecorationSinglepumpALTERA: return "SinglepumpALTERA"; + case DecorationDoublepumpALTERA: return "DoublepumpALTERA"; + case DecorationMaxReplicatesALTERA: return "MaxReplicatesALTERA"; + case DecorationSimpleDualPortALTERA: return "SimpleDualPortALTERA"; + case DecorationMergeALTERA: return "MergeALTERA"; + case DecorationBankBitsALTERA: return "BankBitsALTERA"; + case DecorationForcePow2DepthALTERA: return "ForcePow2DepthALTERA"; + case DecorationStridesizeALTERA: return "StridesizeALTERA"; + case DecorationWordsizeALTERA: return "WordsizeALTERA"; + case DecorationTrueDualPortALTERA: return "TrueDualPortALTERA"; + case DecorationBurstCoalesceALTERA: return "BurstCoalesceALTERA"; + case DecorationCacheSizeALTERA: return "CacheSizeALTERA"; + case DecorationDontStaticallyCoalesceALTERA: return "DontStaticallyCoalesceALTERA"; + case DecorationPrefetchALTERA: return "PrefetchALTERA"; + case DecorationStallEnableALTERA: return "StallEnableALTERA"; + case DecorationFuseLoopsInFunctionALTERA: return "FuseLoopsInFunctionALTERA"; + case DecorationMathOpDSPModeALTERA: return "MathOpDSPModeALTERA"; case DecorationAliasScopeINTEL: return "AliasScopeINTEL"; case DecorationNoAliasINTEL: return "NoAliasINTEL"; - case DecorationInitiationIntervalINTEL: return "InitiationIntervalINTEL"; - case DecorationMaxConcurrencyINTEL: return "MaxConcurrencyINTEL"; - case DecorationPipelineEnableINTEL: return "PipelineEnableINTEL"; - case DecorationBufferLocationINTEL: return "BufferLocationINTEL"; - case DecorationIOPipeStorageINTEL: return "IOPipeStorageINTEL"; + case DecorationInitiationIntervalALTERA: return "InitiationIntervalALTERA"; + case DecorationMaxConcurrencyALTERA: return "MaxConcurrencyALTERA"; + case DecorationPipelineEnableALTERA: return "PipelineEnableALTERA"; + case DecorationBufferLocationALTERA: return "BufferLocationALTERA"; + case DecorationIOPipeStorageALTERA: return "IOPipeStorageALTERA"; case DecorationFunctionFloatingPointModeINTEL: return "FunctionFloatingPointModeINTEL"; case DecorationSingleElementVectorINTEL: return "SingleElementVectorINTEL"; case DecorationVectorComputeCallableFunctionINTEL: return "VectorComputeCallableFunctionINTEL"; case DecorationMediaBlockIOINTEL: return "MediaBlockIOINTEL"; - case DecorationStallFreeINTEL: return "StallFreeINTEL"; + case DecorationStallFreeALTERA: return "StallFreeALTERA"; case DecorationFPMaxErrorDecorationINTEL: return "FPMaxErrorDecorationINTEL"; - case DecorationLatencyControlLabelINTEL: return "LatencyControlLabelINTEL"; - case DecorationLatencyControlConstraintINTEL: return "LatencyControlConstraintINTEL"; - case DecorationConduitKernelArgumentINTEL: return "ConduitKernelArgumentINTEL"; - case DecorationRegisterMapKernelArgumentINTEL: return "RegisterMapKernelArgumentINTEL"; - case DecorationMMHostInterfaceAddressWidthINTEL: return "MMHostInterfaceAddressWidthINTEL"; - case DecorationMMHostInterfaceDataWidthINTEL: return "MMHostInterfaceDataWidthINTEL"; - case DecorationMMHostInterfaceLatencyINTEL: return "MMHostInterfaceLatencyINTEL"; - case DecorationMMHostInterfaceReadWriteModeINTEL: return "MMHostInterfaceReadWriteModeINTEL"; - case DecorationMMHostInterfaceMaxBurstINTEL: return "MMHostInterfaceMaxBurstINTEL"; - case DecorationMMHostInterfaceWaitRequestINTEL: return "MMHostInterfaceWaitRequestINTEL"; - case DecorationStableKernelArgumentINTEL: return "StableKernelArgumentINTEL"; + case DecorationLatencyControlLabelALTERA: return "LatencyControlLabelALTERA"; + case DecorationLatencyControlConstraintALTERA: return "LatencyControlConstraintALTERA"; + case DecorationConduitKernelArgumentALTERA: return "ConduitKernelArgumentALTERA"; + case DecorationRegisterMapKernelArgumentALTERA: return "RegisterMapKernelArgumentALTERA"; + case DecorationMMHostInterfaceAddressWidthALTERA: return "MMHostInterfaceAddressWidthALTERA"; + case DecorationMMHostInterfaceDataWidthALTERA: return "MMHostInterfaceDataWidthALTERA"; + case DecorationMMHostInterfaceLatencyALTERA: return "MMHostInterfaceLatencyALTERA"; + case DecorationMMHostInterfaceReadWriteModeALTERA: return "MMHostInterfaceReadWriteModeALTERA"; + case DecorationMMHostInterfaceMaxBurstALTERA: return "MMHostInterfaceMaxBurstALTERA"; + case DecorationMMHostInterfaceWaitRequestALTERA: return "MMHostInterfaceWaitRequestALTERA"; + case DecorationStableKernelArgumentALTERA: return "StableKernelArgumentALTERA"; case DecorationHostAccessINTEL: return "HostAccessINTEL"; - case DecorationInitModeINTEL: return "InitModeINTEL"; - case DecorationImplementInRegisterMapINTEL: return "ImplementInRegisterMapINTEL"; + case DecorationInitModeALTERA: return "InitModeALTERA"; + case DecorationImplementInRegisterMapALTERA: return "ImplementInRegisterMapALTERA"; case DecorationConditionalINTEL: return "ConditionalINTEL"; case DecorationCacheControlLoadINTEL: return "CacheControlLoadINTEL"; case DecorationCacheControlStoreINTEL: return "CacheControlStoreINTEL"; @@ -3876,6 +4138,8 @@ inline const char* BuiltInToString(BuiltIn value) { case BuiltInFragStencilRefEXT: return "FragStencilRefEXT"; case BuiltInRemainingRecursionLevelsAMDX: return "RemainingRecursionLevelsAMDX"; case BuiltInShaderIndexAMDX: return "ShaderIndexAMDX"; + case BuiltInSamplerHeapEXT: return "SamplerHeapEXT"; + case BuiltInResourceHeapEXT: return "ResourceHeapEXT"; case BuiltInViewportMaskNV: return "ViewportMaskNV"; case BuiltInSecondaryPositionNV: return "SecondaryPositionNV"; case BuiltInSecondaryViewportMaskNV: return "SecondaryViewportMaskNV"; @@ -3954,9 +4218,9 @@ inline const char* GroupOperationToString(GroupOperation value) { case GroupOperationInclusiveScan: return "InclusiveScan"; case GroupOperationExclusiveScan: return "ExclusiveScan"; case GroupOperationClusteredReduce: return "ClusteredReduce"; - case GroupOperationPartitionedReduceNV: return "PartitionedReduceNV"; - case GroupOperationPartitionedInclusiveScanNV: return "PartitionedInclusiveScanNV"; - case GroupOperationPartitionedExclusiveScanNV: return "PartitionedExclusiveScanNV"; + case GroupOperationPartitionedReduceEXT: return "PartitionedReduceEXT"; + case GroupOperationPartitionedInclusiveScanEXT: return "PartitionedInclusiveScanEXT"; + case GroupOperationPartitionedExclusiveScanEXT: return "PartitionedExclusiveScanEXT"; default: return "Unknown"; } } @@ -4103,6 +4367,10 @@ inline const char* CapabilityToString(Capability value) { case CapabilityBFloat16TypeKHR: return "BFloat16TypeKHR"; case CapabilityBFloat16DotProductKHR: return "BFloat16DotProductKHR"; case CapabilityBFloat16CooperativeMatrixKHR: return "BFloat16CooperativeMatrixKHR"; + case CapabilityAbortKHR: return "AbortKHR"; + case CapabilityDescriptorHeapEXT: return "DescriptorHeapEXT"; + case CapabilityConstantDataKHR: return "ConstantDataKHR"; + case CapabilityPoisonFreezeKHR: return "PoisonFreezeKHR"; case CapabilitySampleMaskOverrideCoverageNV: return "SampleMaskOverrideCoverageNV"; case CapabilityGeometryShaderPassthroughNV: return "GeometryShaderPassthroughNV"; case CapabilityShaderViewportIndexLayerEXT: return "ShaderViewportIndexLayerEXT"; @@ -4116,7 +4384,7 @@ inline const char* CapabilityToString(Capability value) { case CapabilityFragmentBarycentricKHR: return "FragmentBarycentricKHR"; case CapabilityComputeDerivativeGroupQuadsKHR: return "ComputeDerivativeGroupQuadsKHR"; case CapabilityFragmentDensityEXT: return "FragmentDensityEXT"; - case CapabilityGroupNonUniformPartitionedNV: return "GroupNonUniformPartitionedNV"; + case CapabilityGroupNonUniformPartitionedEXT: return "GroupNonUniformPartitionedEXT"; case CapabilityShaderNonUniform: return "ShaderNonUniform"; case CapabilityRuntimeDescriptorArray: return "RuntimeDescriptorArray"; case CapabilityInputAttachmentArrayDynamicIndexing: return "InputAttachmentArrayDynamicIndexing"; @@ -4146,6 +4414,7 @@ inline const char* CapabilityToString(Capability value) { case CapabilityDisplacementMicromapNV: return "DisplacementMicromapNV"; case CapabilityRayTracingOpacityMicromapEXT: return "RayTracingOpacityMicromapEXT"; case CapabilityShaderInvocationReorderNV: return "ShaderInvocationReorderNV"; + case CapabilityShaderInvocationReorderEXT: return "ShaderInvocationReorderEXT"; case CapabilityBindlessTextureNV: return "BindlessTextureNV"; case CapabilityRayQueryPositionFetchKHR: return "RayQueryPositionFetchKHR"; case CapabilityCooperativeVectorNV: return "CooperativeVectorNV"; @@ -4154,6 +4423,9 @@ inline const char* CapabilityToString(Capability value) { case CapabilityRawAccessChainsNV: return "RawAccessChainsNV"; case CapabilityRayTracingSpheresGeometryNV: return "RayTracingSpheresGeometryNV"; case CapabilityRayTracingLinearSweptSpheresGeometryNV: return "RayTracingLinearSweptSpheresGeometryNV"; + case CapabilityPushConstantBanksNV: return "PushConstantBanksNV"; + case CapabilityLongVectorEXT: return "LongVectorEXT"; + case CapabilityShader64BitIndexingEXT: return "Shader64BitIndexingEXT"; case CapabilityCooperativeMatrixReductionsNV: return "CooperativeMatrixReductionsNV"; case CapabilityCooperativeMatrixConversionsNV: return "CooperativeMatrixConversionsNV"; case CapabilityCooperativeMatrixPerElementOperationsNV: return "CooperativeMatrixPerElementOperationsNV"; @@ -4183,27 +4455,27 @@ inline const char* CapabilityToString(Capability value) { case CapabilitySubgroupAvcMotionEstimationChromaINTEL: return "SubgroupAvcMotionEstimationChromaINTEL"; case CapabilityVariableLengthArrayINTEL: return "VariableLengthArrayINTEL"; case CapabilityFunctionFloatControlINTEL: return "FunctionFloatControlINTEL"; - case CapabilityFPGAMemoryAttributesINTEL: return "FPGAMemoryAttributesINTEL"; + case CapabilityFPGAMemoryAttributesALTERA: return "FPGAMemoryAttributesALTERA"; case CapabilityFPFastMathModeINTEL: return "FPFastMathModeINTEL"; - case CapabilityArbitraryPrecisionIntegersINTEL: return "ArbitraryPrecisionIntegersINTEL"; - case CapabilityArbitraryPrecisionFloatingPointINTEL: return "ArbitraryPrecisionFloatingPointINTEL"; + case CapabilityArbitraryPrecisionIntegersALTERA: return "ArbitraryPrecisionIntegersALTERA"; + case CapabilityArbitraryPrecisionFloatingPointALTERA: return "ArbitraryPrecisionFloatingPointALTERA"; case CapabilityUnstructuredLoopControlsINTEL: return "UnstructuredLoopControlsINTEL"; - case CapabilityFPGALoopControlsINTEL: return "FPGALoopControlsINTEL"; + case CapabilityFPGALoopControlsALTERA: return "FPGALoopControlsALTERA"; case CapabilityKernelAttributesINTEL: return "KernelAttributesINTEL"; case CapabilityFPGAKernelAttributesINTEL: return "FPGAKernelAttributesINTEL"; - case CapabilityFPGAMemoryAccessesINTEL: return "FPGAMemoryAccessesINTEL"; - case CapabilityFPGAClusterAttributesINTEL: return "FPGAClusterAttributesINTEL"; - case CapabilityLoopFuseINTEL: return "LoopFuseINTEL"; - case CapabilityFPGADSPControlINTEL: return "FPGADSPControlINTEL"; + case CapabilityFPGAMemoryAccessesALTERA: return "FPGAMemoryAccessesALTERA"; + case CapabilityFPGAClusterAttributesALTERA: return "FPGAClusterAttributesALTERA"; + case CapabilityLoopFuseALTERA: return "LoopFuseALTERA"; + case CapabilityFPGADSPControlALTERA: return "FPGADSPControlALTERA"; case CapabilityMemoryAccessAliasingINTEL: return "MemoryAccessAliasingINTEL"; - case CapabilityFPGAInvocationPipeliningAttributesINTEL: return "FPGAInvocationPipeliningAttributesINTEL"; - case CapabilityFPGABufferLocationINTEL: return "FPGABufferLocationINTEL"; - case CapabilityArbitraryPrecisionFixedPointINTEL: return "ArbitraryPrecisionFixedPointINTEL"; - case CapabilityUSMStorageClassesINTEL: return "USMStorageClassesINTEL"; - case CapabilityRuntimeAlignedAttributeINTEL: return "RuntimeAlignedAttributeINTEL"; - case CapabilityIOPipesINTEL: return "IOPipesINTEL"; - case CapabilityBlockingPipesINTEL: return "BlockingPipesINTEL"; - case CapabilityFPGARegINTEL: return "FPGARegINTEL"; + case CapabilityFPGAInvocationPipeliningAttributesALTERA: return "FPGAInvocationPipeliningAttributesALTERA"; + case CapabilityFPGABufferLocationALTERA: return "FPGABufferLocationALTERA"; + case CapabilityArbitraryPrecisionFixedPointALTERA: return "ArbitraryPrecisionFixedPointALTERA"; + case CapabilityUSMStorageClassesALTERA: return "USMStorageClassesALTERA"; + case CapabilityRuntimeAlignedAttributeALTERA: return "RuntimeAlignedAttributeALTERA"; + case CapabilityIOPipesALTERA: return "IOPipesALTERA"; + case CapabilityBlockingPipesALTERA: return "BlockingPipesALTERA"; + case CapabilityFPGARegALTERA: return "FPGARegALTERA"; case CapabilityDotProductInputAll: return "DotProductInputAll"; case CapabilityDotProductInput4x8Bit: return "DotProductInput4x8Bit"; case CapabilityDotProductInput4x8BitPacked: return "DotProductInput4x8BitPacked"; @@ -4224,14 +4496,14 @@ inline const char* CapabilityToString(Capability value) { case CapabilityBFloat16ConversionINTEL: return "BFloat16ConversionINTEL"; case CapabilitySplitBarrierINTEL: return "SplitBarrierINTEL"; case CapabilityArithmeticFenceEXT: return "ArithmeticFenceEXT"; - case CapabilityFPGAClusterAttributesV2INTEL: return "FPGAClusterAttributesV2INTEL"; + case CapabilityFPGAClusterAttributesV2ALTERA: return "FPGAClusterAttributesV2ALTERA"; case CapabilityFPGAKernelAttributesv2INTEL: return "FPGAKernelAttributesv2INTEL"; - case CapabilityTaskSequenceINTEL: return "TaskSequenceINTEL"; + case CapabilityTaskSequenceALTERA: return "TaskSequenceALTERA"; case CapabilityFPMaxErrorINTEL: return "FPMaxErrorINTEL"; - case CapabilityFPGALatencyControlINTEL: return "FPGALatencyControlINTEL"; - case CapabilityFPGAArgumentInterfacesINTEL: return "FPGAArgumentInterfacesINTEL"; + case CapabilityFPGALatencyControlALTERA: return "FPGALatencyControlALTERA"; + case CapabilityFPGAArgumentInterfacesALTERA: return "FPGAArgumentInterfacesALTERA"; case CapabilityGlobalVariableHostAccessINTEL: return "GlobalVariableHostAccessINTEL"; - case CapabilityGlobalVariableFPGADecorationsINTEL: return "GlobalVariableFPGADecorationsINTEL"; + case CapabilityGlobalVariableFPGADecorationsALTERA: return "GlobalVariableFPGADecorationsALTERA"; case CapabilitySubgroupBufferPrefetchINTEL: return "SubgroupBufferPrefetchINTEL"; case CapabilitySubgroup2DBlockIOINTEL: return "Subgroup2DBlockIOINTEL"; case CapabilitySubgroup2DBlockTransformINTEL: return "Subgroup2DBlockTransformINTEL"; @@ -4247,6 +4519,10 @@ inline const char* CapabilityToString(Capability value) { case CapabilityCacheControlsINTEL: return "CacheControlsINTEL"; case CapabilityRegisterLimitsINTEL: return "RegisterLimitsINTEL"; case CapabilityBindlessImagesINTEL: return "BindlessImagesINTEL"; + case CapabilityDotProductFloat16AccFloat32VALVE: return "DotProductFloat16AccFloat32VALVE"; + case CapabilityDotProductFloat16AccFloat16VALVE: return "DotProductFloat16AccFloat16VALVE"; + case CapabilityDotProductBFloat16AccVALVE: return "DotProductBFloat16AccVALVE"; + case CapabilityDotProductFloat8AccFloat32VALVE: return "DotProductFloat8AccFloat32VALVE"; default: return "Unknown"; } } @@ -4355,8 +4631,8 @@ inline const char* TensorClampModeToString(TensorClampMode value) { inline const char* InitializationModeQualifierToString(InitializationModeQualifier value) { switch (value) { - case InitializationModeQualifierInitOnDeviceReprogramINTEL: return "InitOnDeviceReprogramINTEL"; - case InitializationModeQualifierInitOnDeviceResetINTEL: return "InitOnDeviceResetINTEL"; + case InitializationModeQualifierInitOnDeviceReprogramALTERA: return "InitOnDeviceReprogramALTERA"; + case InitializationModeQualifierInitOnDeviceResetALTERA: return "InitOnDeviceResetALTERA"; default: return "Unknown"; } } @@ -4877,6 +5153,16 @@ inline const char* OpToString(Op value) { case OpSpecConstantStringAMDX: return "OpSpecConstantStringAMDX"; case OpGroupNonUniformQuadAllKHR: return "OpGroupNonUniformQuadAllKHR"; case OpGroupNonUniformQuadAnyKHR: return "OpGroupNonUniformQuadAnyKHR"; + case OpTypeBufferEXT: return "OpTypeBufferEXT"; + case OpBufferPointerEXT: return "OpBufferPointerEXT"; + case OpAbortKHR: return "OpAbortKHR"; + case OpUntypedImageTexelPointerEXT: return "OpUntypedImageTexelPointerEXT"; + case OpMemberDecorateIdEXT: return "OpMemberDecorateIdEXT"; + case OpConstantSizeOfEXT: return "OpConstantSizeOfEXT"; + case OpConstantDataKHR: return "OpConstantDataKHR"; + case OpSpecConstantDataKHR: return "OpSpecConstantDataKHR"; + case OpPoisonKHR: return "OpPoisonKHR"; + case OpFreezeKHR: return "OpFreezeKHR"; case OpHitObjectRecordHitMotionNV: return "OpHitObjectRecordHitMotionNV"; case OpHitObjectRecordHitWithIndexMotionNV: return "OpHitObjectRecordHitWithIndexMotionNV"; case OpHitObjectRecordMissMotionNV: return "OpHitObjectRecordMissMotionNV"; @@ -4919,12 +5205,42 @@ inline const char* OpToString(Op value) { case OpCooperativeMatrixConvertNV: return "OpCooperativeMatrixConvertNV"; case OpEmitMeshTasksEXT: return "OpEmitMeshTasksEXT"; case OpSetMeshOutputsEXT: return "OpSetMeshOutputsEXT"; - case OpGroupNonUniformPartitionNV: return "OpGroupNonUniformPartitionNV"; + case OpGroupNonUniformPartitionEXT: return "OpGroupNonUniformPartitionEXT"; case OpWritePackedPrimitiveIndices4x8NV: return "OpWritePackedPrimitiveIndices4x8NV"; case OpFetchMicroTriangleVertexPositionNV: return "OpFetchMicroTriangleVertexPositionNV"; case OpFetchMicroTriangleVertexBarycentricNV: return "OpFetchMicroTriangleVertexBarycentricNV"; case OpCooperativeVectorLoadNV: return "OpCooperativeVectorLoadNV"; case OpCooperativeVectorStoreNV: return "OpCooperativeVectorStoreNV"; + case OpHitObjectRecordFromQueryEXT: return "OpHitObjectRecordFromQueryEXT"; + case OpHitObjectRecordMissEXT: return "OpHitObjectRecordMissEXT"; + case OpHitObjectRecordMissMotionEXT: return "OpHitObjectRecordMissMotionEXT"; + case OpHitObjectGetIntersectionTriangleVertexPositionsEXT: return "OpHitObjectGetIntersectionTriangleVertexPositionsEXT"; + case OpHitObjectGetRayFlagsEXT: return "OpHitObjectGetRayFlagsEXT"; + case OpHitObjectSetShaderBindingTableRecordIndexEXT: return "OpHitObjectSetShaderBindingTableRecordIndexEXT"; + case OpHitObjectReorderExecuteShaderEXT: return "OpHitObjectReorderExecuteShaderEXT"; + case OpHitObjectTraceReorderExecuteEXT: return "OpHitObjectTraceReorderExecuteEXT"; + case OpHitObjectTraceMotionReorderExecuteEXT: return "OpHitObjectTraceMotionReorderExecuteEXT"; + case OpTypeHitObjectEXT: return "OpTypeHitObjectEXT"; + case OpReorderThreadWithHintEXT: return "OpReorderThreadWithHintEXT"; + case OpReorderThreadWithHitObjectEXT: return "OpReorderThreadWithHitObjectEXT"; + case OpHitObjectTraceRayEXT: return "OpHitObjectTraceRayEXT"; + case OpHitObjectTraceRayMotionEXT: return "OpHitObjectTraceRayMotionEXT"; + case OpHitObjectRecordEmptyEXT: return "OpHitObjectRecordEmptyEXT"; + case OpHitObjectExecuteShaderEXT: return "OpHitObjectExecuteShaderEXT"; + case OpHitObjectGetCurrentTimeEXT: return "OpHitObjectGetCurrentTimeEXT"; + case OpHitObjectGetAttributesEXT: return "OpHitObjectGetAttributesEXT"; + case OpHitObjectGetHitKindEXT: return "OpHitObjectGetHitKindEXT"; + case OpHitObjectGetPrimitiveIndexEXT: return "OpHitObjectGetPrimitiveIndexEXT"; + case OpHitObjectGetGeometryIndexEXT: return "OpHitObjectGetGeometryIndexEXT"; + case OpHitObjectGetInstanceIdEXT: return "OpHitObjectGetInstanceIdEXT"; + case OpHitObjectGetInstanceCustomIndexEXT: return "OpHitObjectGetInstanceCustomIndexEXT"; + case OpHitObjectGetObjectRayOriginEXT: return "OpHitObjectGetObjectRayOriginEXT"; + case OpHitObjectGetObjectRayDirectionEXT: return "OpHitObjectGetObjectRayDirectionEXT"; + case OpHitObjectGetWorldRayDirectionEXT: return "OpHitObjectGetWorldRayDirectionEXT"; + case OpHitObjectGetWorldRayOriginEXT: return "OpHitObjectGetWorldRayOriginEXT"; + case OpHitObjectGetObjectToWorldEXT: return "OpHitObjectGetObjectToWorldEXT"; + case OpHitObjectGetWorldToObjectEXT: return "OpHitObjectGetWorldToObjectEXT"; + case OpHitObjectGetRayTMaxEXT: return "OpHitObjectGetRayTMaxEXT"; case OpReportIntersectionKHR: return "OpReportIntersectionKHR"; case OpIgnoreIntersectionNV: return "OpIgnoreIntersectionNV"; case OpTerminateRayNV: return "OpTerminateRayNV"; @@ -4936,6 +5252,12 @@ inline const char* OpToString(Op value) { case OpExecuteCallableNV: return "OpExecuteCallableNV"; case OpRayQueryGetClusterIdNV: return "OpRayQueryGetClusterIdNV"; case OpHitObjectGetClusterIdNV: return "OpHitObjectGetClusterIdNV"; + case OpHitObjectGetRayTMinEXT: return "OpHitObjectGetRayTMinEXT"; + case OpHitObjectGetShaderBindingTableRecordIndexEXT: return "OpHitObjectGetShaderBindingTableRecordIndexEXT"; + case OpHitObjectGetShaderRecordBufferHandleEXT: return "OpHitObjectGetShaderRecordBufferHandleEXT"; + case OpHitObjectIsEmptyEXT: return "OpHitObjectIsEmptyEXT"; + case OpHitObjectIsHitEXT: return "OpHitObjectIsHitEXT"; + case OpHitObjectIsMissEXT: return "OpHitObjectIsMissEXT"; case OpTypeCooperativeMatrixNV: return "OpTypeCooperativeMatrixNV"; case OpCooperativeMatrixLoadNV: return "OpCooperativeMatrixLoadNV"; case OpCooperativeMatrixStoreNV: return "OpCooperativeMatrixStoreNV"; @@ -5139,24 +5461,24 @@ inline const char* OpToString(Op value) { case OpVariableLengthArrayINTEL: return "OpVariableLengthArrayINTEL"; case OpSaveMemoryINTEL: return "OpSaveMemoryINTEL"; case OpRestoreMemoryINTEL: return "OpRestoreMemoryINTEL"; - case OpArbitraryFloatSinCosPiINTEL: return "OpArbitraryFloatSinCosPiINTEL"; - case OpArbitraryFloatCastINTEL: return "OpArbitraryFloatCastINTEL"; - case OpArbitraryFloatCastFromIntINTEL: return "OpArbitraryFloatCastFromIntINTEL"; - case OpArbitraryFloatCastToIntINTEL: return "OpArbitraryFloatCastToIntINTEL"; - case OpArbitraryFloatAddINTEL: return "OpArbitraryFloatAddINTEL"; - case OpArbitraryFloatSubINTEL: return "OpArbitraryFloatSubINTEL"; - case OpArbitraryFloatMulINTEL: return "OpArbitraryFloatMulINTEL"; - case OpArbitraryFloatDivINTEL: return "OpArbitraryFloatDivINTEL"; - case OpArbitraryFloatGTINTEL: return "OpArbitraryFloatGTINTEL"; - case OpArbitraryFloatGEINTEL: return "OpArbitraryFloatGEINTEL"; - case OpArbitraryFloatLTINTEL: return "OpArbitraryFloatLTINTEL"; - case OpArbitraryFloatLEINTEL: return "OpArbitraryFloatLEINTEL"; - case OpArbitraryFloatEQINTEL: return "OpArbitraryFloatEQINTEL"; - case OpArbitraryFloatRecipINTEL: return "OpArbitraryFloatRecipINTEL"; - case OpArbitraryFloatRSqrtINTEL: return "OpArbitraryFloatRSqrtINTEL"; - case OpArbitraryFloatCbrtINTEL: return "OpArbitraryFloatCbrtINTEL"; - case OpArbitraryFloatHypotINTEL: return "OpArbitraryFloatHypotINTEL"; - case OpArbitraryFloatSqrtINTEL: return "OpArbitraryFloatSqrtINTEL"; + case OpArbitraryFloatSinCosPiALTERA: return "OpArbitraryFloatSinCosPiALTERA"; + case OpArbitraryFloatCastALTERA: return "OpArbitraryFloatCastALTERA"; + case OpArbitraryFloatCastFromIntALTERA: return "OpArbitraryFloatCastFromIntALTERA"; + case OpArbitraryFloatCastToIntALTERA: return "OpArbitraryFloatCastToIntALTERA"; + case OpArbitraryFloatAddALTERA: return "OpArbitraryFloatAddALTERA"; + case OpArbitraryFloatSubALTERA: return "OpArbitraryFloatSubALTERA"; + case OpArbitraryFloatMulALTERA: return "OpArbitraryFloatMulALTERA"; + case OpArbitraryFloatDivALTERA: return "OpArbitraryFloatDivALTERA"; + case OpArbitraryFloatGTALTERA: return "OpArbitraryFloatGTALTERA"; + case OpArbitraryFloatGEALTERA: return "OpArbitraryFloatGEALTERA"; + case OpArbitraryFloatLTALTERA: return "OpArbitraryFloatLTALTERA"; + case OpArbitraryFloatLEALTERA: return "OpArbitraryFloatLEALTERA"; + case OpArbitraryFloatEQALTERA: return "OpArbitraryFloatEQALTERA"; + case OpArbitraryFloatRecipALTERA: return "OpArbitraryFloatRecipALTERA"; + case OpArbitraryFloatRSqrtALTERA: return "OpArbitraryFloatRSqrtALTERA"; + case OpArbitraryFloatCbrtALTERA: return "OpArbitraryFloatCbrtALTERA"; + case OpArbitraryFloatHypotALTERA: return "OpArbitraryFloatHypotALTERA"; + case OpArbitraryFloatSqrtALTERA: return "OpArbitraryFloatSqrtALTERA"; case OpArbitraryFloatLogINTEL: return "OpArbitraryFloatLogINTEL"; case OpArbitraryFloatLog2INTEL: return "OpArbitraryFloatLog2INTEL"; case OpArbitraryFloatLog10INTEL: return "OpArbitraryFloatLog10INTEL"; @@ -5184,22 +5506,22 @@ inline const char* OpToString(Op value) { case OpAliasDomainDeclINTEL: return "OpAliasDomainDeclINTEL"; case OpAliasScopeDeclINTEL: return "OpAliasScopeDeclINTEL"; case OpAliasScopeListDeclINTEL: return "OpAliasScopeListDeclINTEL"; - case OpFixedSqrtINTEL: return "OpFixedSqrtINTEL"; - case OpFixedRecipINTEL: return "OpFixedRecipINTEL"; - case OpFixedRsqrtINTEL: return "OpFixedRsqrtINTEL"; - case OpFixedSinINTEL: return "OpFixedSinINTEL"; - case OpFixedCosINTEL: return "OpFixedCosINTEL"; - case OpFixedSinCosINTEL: return "OpFixedSinCosINTEL"; - case OpFixedSinPiINTEL: return "OpFixedSinPiINTEL"; - case OpFixedCosPiINTEL: return "OpFixedCosPiINTEL"; - case OpFixedSinCosPiINTEL: return "OpFixedSinCosPiINTEL"; - case OpFixedLogINTEL: return "OpFixedLogINTEL"; - case OpFixedExpINTEL: return "OpFixedExpINTEL"; - case OpPtrCastToCrossWorkgroupINTEL: return "OpPtrCastToCrossWorkgroupINTEL"; - case OpCrossWorkgroupCastToPtrINTEL: return "OpCrossWorkgroupCastToPtrINTEL"; - case OpReadPipeBlockingINTEL: return "OpReadPipeBlockingINTEL"; - case OpWritePipeBlockingINTEL: return "OpWritePipeBlockingINTEL"; - case OpFPGARegINTEL: return "OpFPGARegINTEL"; + case OpFixedSqrtALTERA: return "OpFixedSqrtALTERA"; + case OpFixedRecipALTERA: return "OpFixedRecipALTERA"; + case OpFixedRsqrtALTERA: return "OpFixedRsqrtALTERA"; + case OpFixedSinALTERA: return "OpFixedSinALTERA"; + case OpFixedCosALTERA: return "OpFixedCosALTERA"; + case OpFixedSinCosALTERA: return "OpFixedSinCosALTERA"; + case OpFixedSinPiALTERA: return "OpFixedSinPiALTERA"; + case OpFixedCosPiALTERA: return "OpFixedCosPiALTERA"; + case OpFixedSinCosPiALTERA: return "OpFixedSinCosPiALTERA"; + case OpFixedLogALTERA: return "OpFixedLogALTERA"; + case OpFixedExpALTERA: return "OpFixedExpALTERA"; + case OpPtrCastToCrossWorkgroupALTERA: return "OpPtrCastToCrossWorkgroupALTERA"; + case OpCrossWorkgroupCastToPtrALTERA: return "OpCrossWorkgroupCastToPtrALTERA"; + case OpReadPipeBlockingALTERA: return "OpReadPipeBlockingALTERA"; + case OpWritePipeBlockingALTERA: return "OpWritePipeBlockingALTERA"; + case OpFPGARegALTERA: return "OpFPGARegALTERA"; case OpRayQueryGetRayTMinKHR: return "OpRayQueryGetRayTMinKHR"; case OpRayQueryGetRayFlagsKHR: return "OpRayQueryGetRayFlagsKHR"; case OpRayQueryGetIntersectionTKHR: return "OpRayQueryGetIntersectionTKHR"; @@ -5228,11 +5550,11 @@ inline const char* OpToString(Op value) { case OpControlBarrierArriveINTEL: return "OpControlBarrierArriveINTEL"; case OpControlBarrierWaitINTEL: return "OpControlBarrierWaitINTEL"; case OpArithmeticFenceEXT: return "OpArithmeticFenceEXT"; - case OpTaskSequenceCreateINTEL: return "OpTaskSequenceCreateINTEL"; - case OpTaskSequenceAsyncINTEL: return "OpTaskSequenceAsyncINTEL"; - case OpTaskSequenceGetINTEL: return "OpTaskSequenceGetINTEL"; - case OpTaskSequenceReleaseINTEL: return "OpTaskSequenceReleaseINTEL"; - case OpTypeTaskSequenceINTEL: return "OpTypeTaskSequenceINTEL"; + case OpTaskSequenceCreateALTERA: return "OpTaskSequenceCreateALTERA"; + case OpTaskSequenceAsyncALTERA: return "OpTaskSequenceAsyncALTERA"; + case OpTaskSequenceGetALTERA: return "OpTaskSequenceGetALTERA"; + case OpTaskSequenceReleaseALTERA: return "OpTaskSequenceReleaseALTERA"; + case OpTypeTaskSequenceALTERA: return "OpTypeTaskSequenceALTERA"; case OpSubgroupBlockPrefetchINTEL: return "OpSubgroupBlockPrefetchINTEL"; case OpSubgroup2DBlockLoadINTEL: return "OpSubgroup2DBlockLoadINTEL"; case OpSubgroup2DBlockLoadTransformINTEL: return "OpSubgroup2DBlockLoadTransformINTEL"; @@ -5263,6 +5585,9 @@ inline const char* OpToString(Op value) { case OpConvertHandleToImageINTEL: return "OpConvertHandleToImageINTEL"; case OpConvertHandleToSamplerINTEL: return "OpConvertHandleToSamplerINTEL"; case OpConvertHandleToSampledImageINTEL: return "OpConvertHandleToSampledImageINTEL"; + case OpFDot2MixAcc32VALVE: return "OpFDot2MixAcc32VALVE"; + case OpFDot2MixAcc16VALVE: return "OpFDot2MixAcc16VALVE"; + case OpFDot4MixAcc32VALVE: return "OpFDot4MixAcc32VALVE"; default: return "Unknown"; } } diff --git a/3rdparty/spirv-cross/spirv_common.hpp b/3rdparty/spirv-cross/spirv_common.hpp index cd06f754c..b003d4db1 100644 --- a/3rdparty/spirv-cross/spirv_common.hpp +++ b/3rdparty/spirv-cross/spirv_common.hpp @@ -608,7 +608,8 @@ struct SPIRType : IVariant FloatE4M3, FloatE5M2, - Tensor + Tensor, + DescriptorHeapBuffer }; // Scalar/vector/matrix support. @@ -655,6 +656,11 @@ struct SPIRType : IVariant uint32_t rank; uint32_t shape; } tensor; + + struct + { + spv::StorageClass storage; + } descriptor_heap_buffer; } ext; spv::StorageClass storage = spv::StorageClassGeneric; @@ -808,6 +814,13 @@ struct SPIRExpression : IVariant // Whether or not gl_MeshVerticesEXT[].gl_Position (as a whole or .y) is referenced bool access_meshlet_position_y = false; + // If this expression represents a OpBufferPointerEXT cast. + bool buffer_pointer = false; + + // Temporaries which can remain forwarded as long as this variable is not modified. + // Only used for buffer pointers. + SmallVector buffer_pointer_dependees; + // A list of expressions which this expression depends on. SmallVector expression_dependencies; @@ -1200,6 +1213,12 @@ struct SPIRVariable : IVariant // Used to find global LUTs bool is_written_to = false; + // Untyped pointer. The pointer of the variable is effectively void. + // The underlying payload for allocation is in alloca_type, but may be 0 too. + // This is mostly here to support descriptor heap proxy. + bool untyped = false; + ID untyped_alloca_type = 0; + SPIRFunction::Parameter *parameter = nullptr; SPIRV_CROSS_DECLARE_CLONE(SPIRVariable) @@ -1535,6 +1554,9 @@ struct SPIRConstant : IVariant // preprocessor directives before compiling the shader. std::string specialization_constant_macro_name; + // ConstantSizeOfEXT. + ID size_of_type = 0; + SPIRV_CROSS_DECLARE_CLONE(SPIRConstant) }; @@ -1815,10 +1837,12 @@ struct Meta uint32_t set = 0; uint32_t binding = 0; uint32_t offset = 0; + uint32_t offset_id = 0; uint32_t xfb_buffer = 0; uint32_t xfb_stride = 0; uint32_t stream = 0; uint32_t array_stride = 0; + uint32_t array_stride_id = 0; uint32_t matrix_stride = 0; uint32_t input_attachment = 0; uint32_t spec_id = 0; diff --git a/3rdparty/spirv-cross/spirv_cross.cpp b/3rdparty/spirv-cross/spirv_cross.cpp index ee4c2afc1..810084e96 100644 --- a/3rdparty/spirv-cross/spirv_cross.cpp +++ b/3rdparty/spirv-cross/spirv_cross.cpp @@ -78,6 +78,11 @@ string Compiler::compile() bool Compiler::variable_storage_is_aliased(const SPIRVariable &v) { auto &type = get(v.basetype); + + // Untyped pointer, assume full aliasing. + if (type.basetype == SPIRType::Void) + return true; + bool ssbo = v.storage == StorageClassStorageBuffer || ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); bool image = type.basetype == SPIRType::Image; @@ -429,10 +434,19 @@ SPIRVariable *Compiler::maybe_get_backing_variable(uint32_t chain) return var; } +SPIRExpression *Compiler::maybe_get_backing_buffer_pointer(uint32_t chain) +{ + auto *expr = maybe_get(chain); + while (expr && !expr->buffer_pointer && expr->loaded_from) + expr = maybe_get(expr->loaded_from); + return expr && expr->buffer_pointer ? expr : nullptr; +} + void Compiler::register_read(uint32_t expr, uint32_t chain, bool forwarded) { auto &e = get(expr); auto *var = maybe_get_backing_variable(chain); + auto *buffer_pointer = maybe_get_backing_buffer_pointer(chain); if (var) { @@ -447,6 +461,13 @@ void Compiler::register_read(uint32_t expr, uint32_t chain, bool forwarded) if (var && var->parameter) var->parameter->read_count++; } + else if (buffer_pointer) + { + e.loaded_from = buffer_pointer->self; + // If the backing variable is immutable, we do not need to depend on the variable. + if (forwarded && !is_immutable(buffer_pointer->self)) + buffer_pointer->buffer_pointer_dependees.push_back(e.self); + } } void Compiler::register_write(uint32_t chain) @@ -464,6 +485,8 @@ void Compiler::register_write(uint32_t chain) var = maybe_get(access_chain->loaded_from); } + auto *buffer_pointer = maybe_get_backing_buffer_pointer(chain); + auto &chain_type = expression_type(chain); if (var) @@ -508,6 +531,10 @@ void Compiler::register_write(uint32_t chain) force_recompile(); } } + else if (buffer_pointer) + { + flush_dependees(*buffer_pointer); + } else if (chain_type.pointer) { // If we stored through a variable pointer, then we don't know which @@ -529,6 +556,16 @@ void Compiler::flush_dependees(SPIRVariable &var) var.dependees.clear(); } +void Compiler::flush_dependees(SPIRExpression &expr) +{ + // A little ugly to split things up like this since BufferPointerEXT is a weird case + // where it's both an expression (chain into global heap) and a memory declaration at the same time ... + assert(expr.buffer_pointer); + for (auto dep : expr.buffer_pointer_dependees) + invalid_expressions.insert(dep); + expr.buffer_pointer_dependees.clear(); +} + void Compiler::flush_all_aliased_variables() { for (auto aliased : aliased_variables) @@ -539,6 +576,8 @@ void Compiler::flush_all_atomic_capable_variables() { for (auto global : global_variables) flush_dependees(get(global)); + for (auto global : buffer_pointer_variables) + flush_dependees(get(global)); flush_all_aliased_variables(); } @@ -560,6 +599,8 @@ void Compiler::flush_all_active_variables() flush_dependees(get(arg.id)); for (auto global : global_variables) flush_dependees(get(global)); + for (auto global : buffer_pointer_variables) + flush_dependees(get(global)); flush_all_aliased_variables(); } @@ -668,8 +709,9 @@ bool Compiler::is_hidden_variable(const SPIRVariable &var, bool include_builtins } // In SPIR-V 1.4 and up we must also use the active variable interface to disable global variables - // which are not part of the entry point. - if (ir.get_spirv_version() >= 0x10400 && var.storage != StorageClassGeneric && + // which are not part of the entry point. Library modules have no real entry point so the filter + // would hide every global so skip it in that case. + if (ir.get_spirv_version() >= 0x10400 && !ir.is_library_module && var.storage != StorageClassGeneric && var.storage != StorageClassFunction && !interface_variable_exists_in_entry_point(var.self)) { return true; @@ -742,17 +784,17 @@ bool Compiler::is_array(const SPIRType &type) const bool Compiler::is_pointer(const SPIRType &type) const { - return type.op == OpTypePointer && type.basetype != SPIRType::Unknown; // Ignore function pointers. + return (type.op == OpTypePointer || type.op == OpTypeUntypedPointerKHR) && type.basetype != SPIRType::Unknown; // Ignore function pointers. } bool Compiler::is_physical_pointer(const SPIRType &type) const { - return type.op == OpTypePointer && type.storage == StorageClassPhysicalStorageBuffer; + return (type.op == OpTypePointer || type.op == OpTypeUntypedPointerKHR) && type.storage == StorageClassPhysicalStorageBuffer; } bool Compiler::is_physical_or_buffer_pointer(const SPIRType &type) const { - return type.op == OpTypePointer && + return (type.op == OpTypePointer || type.op == OpTypeUntypedPointerKHR) && (type.storage == StorageClassPhysicalStorageBuffer || type.storage == StorageClassUniform || type.storage == StorageClassStorageBuffer || type.storage == StorageClassWorkgroup || type.storage == StorageClassPushConstant); @@ -2068,7 +2110,7 @@ size_t Compiler::get_declared_struct_size_runtime_array(const SPIRType &type, si size_t size = get_declared_struct_size(type); auto &last_type = get(type.member_types.back()); - if (!last_type.array.empty() && last_type.array_size_literal[0] && last_type.array[0] == 0) // Runtime array + if (!last_type.array.empty() && last_type.array_size_literal.back() && last_type.array.back() == 0) // Runtime array size += array_size * type_struct_member_array_stride(type, uint32_t(type.member_types.size() - 1)); return size; @@ -4799,6 +4841,18 @@ void Compiler::build_function_control_flow_graphs_and_analyze() CFGBuilder handler(*this); handler.function_cfgs[ir.default_entry_point].reset(new CFG(*this, get(ir.default_entry_point))); traverse_all_reachable_opcodes(get(ir.default_entry_point), handler); + if (ir.is_library_module) + { + // In library mode, default_entry_point is just the first exported + // function. Build a CFG for every other exported function (and its + // callees) so per-function analyses below cover all of them. + for (auto export_id : ir.library_exported_functions) + { + auto &func = get(export_id); + if (handler.follow_function_call(func)) + traverse_all_reachable_opcodes(func, handler); + } + } function_cfgs = std::move(handler.function_cfgs); bool single_function = function_cfgs.size() <= 1; @@ -5429,6 +5483,212 @@ void Compiler::analyze_non_block_pointer_types() physical_storage_type_to_alignment = std::move(handler.physical_block_type_meta); } +void Compiler::analyze_descriptor_heap_types() +{ + struct HeapHandler : OpcodeHandler + { + bool handle(Op opcode, const uint32_t *args, uint32_t) override + { + switch (opcode) + { + case OpBufferPointerEXT: + { + auto &ptr_type = compiler.get(args[0]); + // BufferPointerEXT can return untyped or typed pointers. + // If it's typed, we resolve it here. + if (ptr_type.basetype == SPIRType::Struct) + { + DescriptorHeapMeta meta = {}; + meta.type = ptr_type.self; + meta.hlsl_style_stride = hlsl_style_stride_access_chains.count(args[2]); + meta.buffer_pointer_id = args[1]; + meta.storage = ptr_type.storage; + meta.nonreadable = compiler.has_decoration(args[1], DecorationNonReadable); + meta.nonwritable = compiler.has_decoration(args[1], DecorationNonWritable); + meta.coherent = compiler.has_decoration(args[1], DecorationCoherent); + meta.is_restrict = compiler.has_decoration(args[1], DecorationRestrict); + meta.is_volatile = compiler.has_decoration(args[1], DecorationVolatile); + add_unique_type(meta); + } + buffer_pointers[args[1]] = { args[0], hlsl_style_stride_access_chains.count(args[2]) != 0 }; + break; + } + + case OpUntypedAccessChainKHR: + case OpUntypedInBoundsAccessChainKHR: + case OpUntypedArrayLengthKHR: + { + auto &data_type = compiler.get(args[2]); + + if (compiler.is_pointer(data_type)) + SPIRV_CROSS_THROW("pointer type not allowed."); + + bool hlsl_style_stride = false; + + // Need to validate the array stride and types. HLLs are not flexible enough to support the full flexibility of SPIR-V. + if (BuiltIn(compiler.get_decoration(args[3], DecorationBuiltIn)) == BuiltInResourceHeapEXT) + { + if (!compiler.is_runtime_size_array(data_type)) + SPIRV_CROSS_THROW("Descriptor heap must be accessed as a runtime array."); + + // The only meaningful use of this is ArrayStride equal to sizeof(type) right now. + uint32_t array_stride_id = compiler.get_decoration(args[2], DecorationArrayStrideIdEXT); + if (!array_stride_id) + SPIRV_CROSS_THROW("Expected ArrayStrideIdEXT to be set for resource heap."); + + auto *spec_c = compiler.maybe_get(array_stride_id); + auto *c = compiler.maybe_get(array_stride_id); + + if (!spec_c && !c) + SPIRV_CROSS_THROW("Array stride must be some constant expression."); + + if (spec_c) + { + // This gets potentially infinitely weird, but if we get HLSL-style shaders + // we expect the array stride to be max(buffer, image) since all descriptors have equal size in D3D12. + // We just have to be a bit loose here since it's impossible to anticipate every theoretical formulation. + // Anything non-conforming to strict GLSL is flagged in the codegen output. + if (spec_c->opcode == OpSelect) + { + auto *true_value = compiler.maybe_get(spec_c->arguments[1]); + auto *false_value = compiler.maybe_get(spec_c->arguments[2]); + hlsl_style_stride = true_value && true_value->size_of_type && + false_value && false_value->size_of_type; + } + + if (!hlsl_style_stride) + SPIRV_CROSS_THROW("Unusual pattern of descriptor stride detected. This probably cannot be expressed in current GLSL."); + } + + if (c && !c->size_of_type) + SPIRV_CROSS_THROW("Resource heap array stride must be ConstantSizeOfEXT for high level languages."); + + auto &element_type = compiler.get(data_type.parent_type); + + if (element_type.basetype == SPIRType::DescriptorHeapBuffer) + { + if (c && compiler.get(c->size_of_type).basetype != SPIRType::DescriptorHeapBuffer) + SPIRV_CROSS_THROW("Buffer descriptors in heap must be ConstantSizeOfEXT(OpTypeBufferEXT) for GLSL."); + } + else if (data_type.basetype == SPIRType::Image) + { + if (c && compiler.get(c->size_of_type).basetype != SPIRType::Image) + SPIRV_CROSS_THROW("Image descriptors in heap must be ConstantSizeOfEXT(OpTypeImage) for GLSL."); + } + else if (data_type.basetype == SPIRType::AccelerationStructure) + { + if (c && compiler.get(c->size_of_type).basetype != SPIRType::AccelerationStructure) + SPIRV_CROSS_THROW("RTAS descriptors in heap must be ConstantSizeOfEXT(OpTypeAccelerationStructure) for GLSL."); + } + } + else if (BuiltIn(compiler.get_decoration(args[3], DecorationBuiltIn)) == BuiltInSamplerHeapEXT) + { + if (!compiler.is_runtime_size_array(data_type)) + SPIRV_CROSS_THROW("Descriptor heap must be accessed as a runtime array."); + + // The only meaningful use of this is ArrayStride equal to sizeof(sampler) right now. + uint32_t array_stride_id = compiler.get_decoration(args[2], DecorationArrayStrideIdEXT); + if (!array_stride_id) + SPIRV_CROSS_THROW("Expected ArrayStrideIdEXT to be set for sampler heap."); + + auto *c = compiler.maybe_get(array_stride_id); + if (!c || !c->size_of_type || compiler.get(c->size_of_type).basetype != SPIRType::Sampler) + SPIRV_CROSS_THROW("Sampler heap array stride must be ConstantSizeOfEXT(OpTypeSampler) for high level languages."); + } + + // Remember this for OpBufferPointerEXT. + if (hlsl_style_stride) + hlsl_style_stride_access_chains.insert(args[1]); + + if (data_type.basetype == SPIRType::SampledImage) + { + SPIRV_CROSS_THROW("Attempting to access heap as combined sampler image. This does not make sense."); + } + else if (data_type.basetype == SPIRType::Image || + data_type.basetype == SPIRType::AccelerationStructure || + data_type.basetype == SPIRType::Sampler) + { + DescriptorHeapMeta meta = {}; + meta.type = data_type.self; + meta.hlsl_style_stride = hlsl_style_stride; + add_unique_type(meta); + } + else if (buffer_pointers.count(args[3]) != 0) + { + if (!compiler.has_decoration(data_type.self, DecorationBlock) && + !compiler.has_decoration(data_type.self, DecorationBufferBlock)) + { + SPIRV_CROSS_THROW("BufferPointerEXT must reference a block type."); + } + + auto &pointer_meta = buffer_pointers[args[3]]; + auto &buffer_type = compiler.get(pointer_meta.type); + if (buffer_type.basetype == SPIRType::Void) + { + // This is where the pointer becomes typed, so register it here. + DescriptorHeapMeta meta = {}; + meta.type = data_type.self; + meta.hlsl_style_stride = pointer_meta.hlsl_style_stride; + meta.buffer_pointer_id = args[3]; + meta.storage = buffer_type.storage; + meta.nonreadable = compiler.has_decoration(args[3], DecorationNonReadable); + meta.nonwritable = compiler.has_decoration(args[3], DecorationNonWritable); + meta.coherent = compiler.has_decoration(args[3], DecorationCoherent); + meta.is_volatile = compiler.has_decoration(args[3], DecorationVolatile); + meta.is_restrict = compiler.has_decoration(args[3], DecorationRestrict); + add_unique_type(meta); + } + } + break; + } + + default: + break; + } + + return true; + } + + explicit HeapHandler(Compiler &compiler_) : OpcodeHandler(compiler_) {} + + std::vector heap_types; + + struct BufferPointerMeta + { + TypeID type; + bool hlsl_style_stride; + }; + std::unordered_map buffer_pointers; + std::unordered_set hlsl_style_stride_access_chains; + + void add_unique_type(const DescriptorHeapMeta &meta) + { + assert(meta.type != 0); + + for (auto &type : heap_types) + { + if (type.type == meta.type && type.storage == meta.storage && + type.buffer_pointer_id == meta.buffer_pointer_id && + type.nonreadable == meta.nonreadable && + type.nonwritable == meta.nonwritable && + type.coherent == meta.coherent && + type.is_restrict == meta.is_restrict && + type.hlsl_style_stride == meta.hlsl_style_stride && + type.is_volatile == meta.is_volatile) + { + return; + } + } + + heap_types.push_back(meta); + } + }; + + HeapHandler handler(*this); + traverse_all_reachable_opcodes(get(ir.default_entry_point), handler); + descriptor_heap_types = std::move(handler.heap_types); +} + bool Compiler::InterlockedResourceAccessPrepassHandler::handle(Op op, const uint32_t *, uint32_t) { if (op == OpBeginInvocationInterlockEXT || op == OpEndInvocationInterlockEXT) diff --git a/3rdparty/spirv-cross/spirv_cross.hpp b/3rdparty/spirv-cross/spirv_cross.hpp index f72d79979..d8b05169f 100644 --- a/3rdparty/spirv-cross/spirv_cross.hpp +++ b/3rdparty/spirv-cross/spirv_cross.hpp @@ -91,6 +91,28 @@ struct BuiltInResource Resource resource; }; +// Needs to stay in sync 1:1 with C API. +enum ResourceType +{ + ResourceTypeUnknown = 0, + ResourceTypeUniformBuffer = 1, + ResourceTypeStorageBuffer = 2, + ResourceTypeStageInput = 3, + ResourceTypeStageOutput = 4, + ResourceTypeSubpassInput = 5, + ResourceTypeStorageImage = 6, + ResourceTypeSampledImage = 7, + ResourceTypeAtomicCounter = 8, + ResourceTypePushConstant = 9, + ResourceTypeSeparateImage = 10, + ResourceTypeSeparateSamplers = 11, + ResourceTypeAccelerationStructure = 12, + ResourceTypeRayQuery = 13, + ResourceTypeShaderRecordBuffer = 14, + ResourceTypeGLPlainUniform = 15, + ResourceTypeTensor = 16 +}; + struct ShaderResources { SmallVector uniform_buffers; @@ -585,6 +607,7 @@ protected: // (SSBO, image load store, etc) SmallVector global_variables; SmallVector aliased_variables; + SmallVector buffer_pointer_variables; SPIRFunction *current_function = nullptr; SPIRBlock *current_block = nullptr; @@ -705,6 +728,7 @@ protected: bool expression_is_lvalue(uint32_t id) const; bool variable_storage_is_aliased(const SPIRVariable &var); SPIRVariable *maybe_get_backing_variable(uint32_t chain); + SPIRExpression *maybe_get_backing_buffer_pointer(uint32_t chain); void register_read(uint32_t expr, uint32_t chain, bool forwarded); void register_write(uint32_t chain); @@ -739,6 +763,7 @@ protected: // Dependency tracking for temporaries read from variables. void flush_dependees(SPIRVariable &var); + void flush_dependees(SPIRExpression &expr); void flush_all_active_variables(); void flush_control_dependent_expressions(uint32_t block); void flush_all_atomic_capable_variables(); @@ -1090,6 +1115,23 @@ protected: SmallVector physical_storage_non_block_pointer_types; std::unordered_map physical_storage_type_to_alignment; + struct DescriptorHeapMeta + { + TypeID type; + bool hlsl_style_stride; + + // For buffers + ID buffer_pointer_id; + StorageClass storage; + bool nonwritable; + bool nonreadable; + bool coherent; + bool is_volatile; + bool is_restrict; + }; + std::vector descriptor_heap_types; + void analyze_descriptor_heap_types(); + void analyze_variable_scope(SPIRFunction &function, AnalyzeVariableScopeAccessHandler &handler); void find_function_local_luts(SPIRFunction &function, const AnalyzeVariableScopeAccessHandler &handler, bool single_function); diff --git a/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp b/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp index bb9a5f58e..47cd1dbe1 100644 --- a/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp +++ b/3rdparty/spirv-cross/spirv_cross_parsed_ir.cpp @@ -79,6 +79,9 @@ ParsedIR &ParsedIR::operator=(ParsedIR &&other) SPIRV_CROSS_NOEXCEPT memory_model = other.memory_model; default_entry_point = other.default_entry_point; + is_library_module = other.is_library_module; + library_exports = std::move(other.library_exports); + library_exported_functions = std::move(other.library_exported_functions); sources = std::move(other.sources); loop_iteration_depth_hard = other.loop_iteration_depth_hard; loop_iteration_depth_soft = other.loop_iteration_depth_soft; @@ -111,6 +114,9 @@ ParsedIR &ParsedIR::operator=(const ParsedIR &other) continue_block_to_loop_header = other.continue_block_to_loop_header; entry_points = other.entry_points; default_entry_point = other.default_entry_point; + is_library_module = other.is_library_module; + library_exports = other.library_exports; + library_exported_functions = other.library_exported_functions; sources = other.sources; loop_iteration_depth_hard = other.loop_iteration_depth_hard; loop_iteration_depth_soft = other.loop_iteration_depth_soft; @@ -404,6 +410,10 @@ void ParsedIR::set_decoration(ID id, Decoration decoration, uint32_t argument) dec.offset = argument; break; + case DecorationOffsetIdEXT: + dec.offset_id = argument; + break; + case DecorationXfbBuffer: dec.xfb_buffer = argument; break; @@ -420,6 +430,10 @@ void ParsedIR::set_decoration(ID id, Decoration decoration, uint32_t argument) dec.array_stride = argument; break; + case DecorationArrayStrideIdEXT: + dec.array_stride_id = argument; + break; + case DecorationMatrixStride: dec.matrix_stride = argument; break; @@ -492,6 +506,10 @@ void ParsedIR::set_member_decoration(TypeID id, uint32_t index, Decoration decor dec.offset = argument; break; + case DecorationOffsetIdEXT: + dec.offset_id = argument; + break; + case DecorationXfbBuffer: dec.xfb_buffer = argument; break; @@ -645,6 +663,8 @@ uint32_t ParsedIR::get_decoration(ID id, Decoration decoration) const return dec.component; case DecorationOffset: return dec.offset; + case DecorationOffsetIdEXT: + return dec.offset_id; case DecorationXfbBuffer: return dec.xfb_buffer; case DecorationXfbStride: @@ -661,6 +681,8 @@ uint32_t ParsedIR::get_decoration(ID id, Decoration decoration) const return dec.spec_id; case DecorationArrayStride: return dec.array_stride; + case DecorationArrayStrideIdEXT: + return dec.array_stride_id; case DecorationMatrixStride: return dec.matrix_stride; case DecorationIndex: @@ -720,6 +742,10 @@ void ParsedIR::unset_decoration(ID id, Decoration decoration) dec.offset = 0; break; + case DecorationOffsetIdEXT: + dec.offset_id = 0; + break; + case DecorationXfbBuffer: dec.xfb_buffer = 0; break; @@ -806,6 +832,8 @@ uint32_t ParsedIR::get_member_decoration(TypeID id, uint32_t index, Decoration d return dec.binding; case DecorationOffset: return dec.offset; + case DecorationOffsetIdEXT: + return dec.offset_id; case DecorationXfbBuffer: return dec.xfb_buffer; case DecorationXfbStride: @@ -903,6 +931,10 @@ void ParsedIR::unset_member_decoration(TypeID id, uint32_t index, Decoration dec dec.offset = 0; break; + case DecorationOffsetIdEXT: + dec.offset_id = 0; + break; + case DecorationXfbBuffer: dec.xfb_buffer = 0; break; diff --git a/3rdparty/spirv-cross/spirv_cross_parsed_ir.hpp b/3rdparty/spirv-cross/spirv_cross_parsed_ir.hpp index b1e76f205..efe9d4a4b 100644 --- a/3rdparty/spirv-cross/spirv_cross_parsed_ir.hpp +++ b/3rdparty/spirv-cross/spirv_cross_parsed_ir.hpp @@ -110,6 +110,13 @@ public: std::unordered_map entry_points; FunctionID default_entry_point = 0; + // A "library" module has no OpEntryPoint and instead exports symbols via + // OpDecorate ... LinkageAttributes ... Export. These vectors keep track + // of all these exports and specifically the function exports. + bool is_library_module = false; + SmallVector library_exports; + SmallVector library_exported_functions; + struct Source { SourceLanguage lang = SourceLanguageUnknown; diff --git a/3rdparty/spirv-cross/spirv_glsl.cpp b/3rdparty/spirv-cross/spirv_glsl.cpp index a20f4498d..fd7180cde 100644 --- a/3rdparty/spirv-cross/spirv_glsl.cpp +++ b/3rdparty/spirv-cross/spirv_glsl.cpp @@ -368,6 +368,7 @@ void CompilerGLSL::reset(uint32_t iteration_count) expression_usage_counts.clear(); forwarded_temporaries.clear(); suppressed_usage_tracking.clear(); + buffer_pointer_variables.clear(); // Ensure that we declare phi-variable copies even if the original declaration isn't deferred flushed_phi_variables.clear(); @@ -611,7 +612,7 @@ void CompilerGLSL::find_static_extensions() { switch (cap) { - case CapabilityShaderNonUniformEXT: + case CapabilityShaderNonUniform: if (!options.vulkan_semantics) require_extension_internal("GL_NV_gpu_shader5"); else @@ -689,6 +690,15 @@ void CompilerGLSL::find_static_extensions() require_extension_internal("GL_ARM_tensors"); break; + case CapabilityDescriptorHeapEXT: + if (!options.vulkan_semantics) + SPIRV_CROSS_THROW("DescriptorHeapEXT requires Vulkan semantics."); + require_extension_internal("GL_EXT_descriptor_heap"); + require_extension_internal("GL_EXT_nonuniform_qualifier"); + // We lose information about writeonly/readonly in SPIR-V. Just pre-empt this to avoid complicating code later. + require_extension_internal("GL_EXT_shader_image_load_formatted"); + break; + default: break; } @@ -783,6 +793,17 @@ string CompilerGLSL::compile() if (ir.addressing_model == AddressingModelPhysicalStorageBuffer64) analyze_non_block_pointer_types(); + if (std::find(ir.declared_capabilities.begin(), ir.declared_capabilities.end(), + CapabilityDescriptorHeapEXT) != ir.declared_capabilities.end()) + { + // Need to figure out all the aliased types that view the heap. + // In GLSL, each unique type must be declared with layout(descriptor_heap) type-decl spvSomeIdentResourceHeap[]; + // During untyped access chain traversal, we prefix the name to match the aliases. + // HLSL has more direct native support and will not need these, but we still need to call this function + // to verify that descriptor strides make sense. + analyze_descriptor_heap_types(); + } + uint32_t pass_count = 0; do { @@ -799,14 +820,23 @@ string CompilerGLSL::compile() if ((options.es || options.vulkan_semantics) && required_polyfills_relaxed != 0) emit_polyfills(required_polyfills_relaxed, true); - emit_function(get(ir.default_entry_point), Bitset()); + if (ir.is_library_module) + { + // Emit each exported function as a normal free function. + // emit_function recursively emits callees, so internal helpers + // are picked up too. + for (auto export_id : ir.library_exported_functions) + emit_function(get(export_id), Bitset()); + } + else + emit_function(get(ir.default_entry_point), Bitset()); pass_count++; } while (is_forcing_recompilation()); // Implement the interlocked wrapper function at the end. // The body was implemented in lieu of main(). - if (interlocked_is_complex) + if (interlocked_is_complex && !ir.is_library_module) { if (options.use_entry_point_name) statement("void ", get_entry_point().name, "()"); @@ -820,8 +850,9 @@ string CompilerGLSL::compile() end_scope(); } - // Entry point in GLSL is always main(). - if (!options.use_entry_point_name) + // Entry point in GLSL is always main(). Skip the rename for library + // modules; their exports keep their declared names. + if (!options.use_entry_point_name && !ir.is_library_module) get_entry_point().name = "main"; return buffer.str(); @@ -894,6 +925,16 @@ void CompilerGLSL::request_subgroup_feature(ShaderSubgroupSupportHelper::Feature void CompilerGLSL::emit_header() { auto &execution = get_entry_point(); + + // Library modules have no entry point. The emitted GLSL is meant to be #include'd or appended + // rather than compiled standalone, so the version and extension directives that follow are + // wrapped in `#ifdef SPIRV_CROSS_LIBRARY_HEADER ... #endif`. By default they are skipped (the + // consuming translation unit provides its own preamble); a caller that wants to compile the + // library standalone defines SPIRV_CROSS_LIBRARY_HEADER to opt in. The stage-specific layout + // block at the end of this function is skipped entirely in library mode. + if (ir.is_library_module) + statement("#ifdef SPIRV_CROSS_LIBRARY_HEADER"); + statement("#version ", options.version, options.es && options.version > 100 ? " es" : ""); if (!options.es && options.version < 420) @@ -1103,6 +1144,13 @@ void CompilerGLSL::emit_header() for (auto &header : header_lines) statement(header); + if (ir.is_library_module) + { + statement("#endif"); + statement(""); + return; + } + SmallVector inputs; SmallVector outputs; @@ -2380,7 +2428,7 @@ void CompilerGLSL::emit_push_constant_block(const SPIRVariable &var) else if (options.vulkan_semantics) emit_push_constant_block_vulkan(var); else if (options.emit_push_constant_as_uniform_buffer) - emit_buffer_block_native(var); + emit_buffer_block_native(&var, nullptr); else emit_push_constant_block_glsl(var); } @@ -2429,7 +2477,7 @@ void CompilerGLSL::emit_buffer_block(const SPIRVariable &var) (ubo_block && options.emit_uniform_buffer_as_plain_uniforms)) emit_buffer_block_legacy(var); else - emit_buffer_block_native(var); + emit_buffer_block_native(&var, nullptr); } void CompilerGLSL::emit_buffer_block_legacy(const SPIRVariable &var) @@ -2575,30 +2623,98 @@ void CompilerGLSL::emit_buffer_reference_block(uint32_t type_id, bool forward_de } } -void CompilerGLSL::emit_buffer_block_native(const SPIRVariable &var) +std::string CompilerGLSL::heap_meta_to_prefix(const DescriptorHeapMeta &meta) { - auto &type = get(var.basetype); + std::string prefix; + + if (meta.nonreadable) + prefix += "NoRead"; + if (meta.nonwritable) + prefix += "NoWrite"; + if (meta.coherent) + prefix += "Coherent"; + if (meta.is_volatile) + prefix += "Volatile"; + if (meta.is_restrict) + prefix += "Restrict"; + + return prefix; +} + +std::string CompilerGLSL::to_buffer_pointer_name_prefix(uint32_t ptr_id) const +{ + auto itr = std::find_if(descriptor_heap_types.begin(), descriptor_heap_types.end(), + [&](const DescriptorHeapMeta &meta) { return meta.buffer_pointer_id == ptr_id; }); + + assert(itr != descriptor_heap_types.end()); + + auto name = to_name(itr->type); + + // The same block type can be instantiated with different read-write decorations. + name += heap_meta_to_prefix(*itr); + + // Disambiguate since we can create multiple buffer pointers with same types. + name += to_name(itr->buffer_pointer_id); + + return join("spv", name); +} + +void CompilerGLSL::emit_buffer_block_native(const SPIRVariable *var, const DescriptorHeapMeta *heap_meta) +{ + assert(var || heap_meta); + + SPIRType *type; + if (var) + type = &get(var->basetype); + else + type = &get(heap_meta->type); + + Bitset flags = var ? ir.get_buffer_block_flags(*var) : ir.get_buffer_block_type_flags(*type); + auto storage = var ? var->storage : heap_meta->storage; + + if (heap_meta) + { + if (heap_meta->nonreadable) + flags.set(DecorationNonReadable); + if (heap_meta->nonwritable) + flags.set(DecorationNonWritable); + if (heap_meta->coherent) + flags.set(DecorationCoherent); + if (heap_meta->is_volatile) + flags.set(DecorationVolatile); + if (heap_meta->is_restrict) + flags.set(DecorationRestrict); + } + + bool ssbo = storage == StorageClassStorageBuffer || storage == StorageClassShaderRecordBufferKHR || + has_decoration(type->self, DecorationBufferBlock); - Bitset flags = ir.get_buffer_block_flags(var); - bool ssbo = var.storage == StorageClassStorageBuffer || var.storage == StorageClassShaderRecordBufferKHR || - ir.meta[type.self].decoration.decoration_flags.get(DecorationBufferBlock); bool is_restrict = ssbo && flags.get(DecorationRestrict); bool is_writeonly = ssbo && flags.get(DecorationNonReadable); bool is_readonly = ssbo && flags.get(DecorationNonWritable); bool is_coherent = ssbo && flags.get(DecorationCoherent); // Block names should never alias, but from HLSL input they kind of can because block types are reused for UAVs ... - auto buffer_name = to_name(type.self, false); + auto buffer_name = to_name(type->self, false); + + if (heap_meta) + { + // The same block type can be instantiated with different read-write decorations. + buffer_name += heap_meta_to_prefix(*heap_meta); + } auto &block_namespace = ssbo ? block_ssbo_names : block_ubo_names; // Shaders never use the block by interface name, so we don't // have to track this other than updating name caches. // If we have a collision for any reason, just fallback immediately. - if (ir.meta[type.self].decoration.alias.empty() || block_namespace.find(buffer_name) != end(block_namespace) || - resource_names.find(buffer_name) != end(resource_names)) + if (var) { - buffer_name = get_block_fallback_name(var.self); + if (ir.meta[type->self].decoration.alias.empty() || block_namespace.find(buffer_name) != end(block_namespace) || + resource_names.find(buffer_name) != end(resource_names)) + { + buffer_name = get_block_fallback_name(var->self); + } } // Make sure we get something unique for both global name scope and block name scope. @@ -2609,40 +2725,68 @@ void CompilerGLSL::emit_buffer_block_native(const SPIRVariable &var) // This cannot conflict with anything else, so we're safe now. // We cannot reuse this fallback name in neither global scope (blocked by block_names) nor block name scope. if (buffer_name.empty()) - buffer_name = join("_", get(var.basetype).self, "_", var.self); + { + if (var) + buffer_name = join("_", get(var->basetype).self, "_", var->self); + else + buffer_name = join("_", type->self); + } block_names.insert(buffer_name); block_namespace.insert(buffer_name); // Save for post-reflection later. - declared_block_names[var.self] = buffer_name; + if (var) + declared_block_names[var->self] = buffer_name; - statement(layout_for_variable(var), is_coherent ? "coherent " : "", is_restrict ? "restrict " : "", + string layout; + + if (var) + { + layout = layout_for_variable(*var); + } + else + { + auto packing_standard = buffer_to_packing_standard(*type, ssbo, true); + layout = join("layout(", + to_descriptor_heap_layout(*type, ssbo ? StorageClassStorageBuffer : StorageClassUniform), + ", ", packing_standard, ") "); + } + + statement(layout, is_coherent ? "coherent " : "", is_restrict ? "restrict " : "", is_writeonly ? "writeonly " : "", is_readonly ? "readonly " : "", ssbo ? "buffer " : "uniform ", buffer_name); begin_scope(); - type.member_name_cache.clear(); + type->member_name_cache.clear(); uint32_t i = 0; - for (auto &member : type.member_types) + for (auto &member : type->member_types) { - add_member_name(type, i); - emit_struct_member(type, member, i); + add_member_name(*type, i); + emit_struct_member(*type, member, i); i++; } // Don't declare empty blocks in GLSL, this is not allowed. - if (type_is_empty(type) && !backend.supports_empty_struct) + if (type_is_empty(*type) && !backend.supports_empty_struct) statement("int empty_struct_member;"); // var.self can be used as a backup name for the block name, // so we need to make sure we don't disturb the name here on a recompile. // It will need to be reset if we have to recompile. - preserve_alias_on_reset(var.self); - add_resource_name(var.self); - end_scope_decl(to_name(var.self) + type_to_array_glsl(type, var.self)); + if (var) + { + preserve_alias_on_reset(var->self); + add_resource_name(var->self); + end_scope_decl(to_name(var->self) + type_to_array_glsl(*type, var->self)); + } + else + { + end_scope_decl(join(to_buffer_pointer_name_prefix(heap_meta->buffer_pointer_id), "ResourceHeap[]")); + } + statement(""); } @@ -3782,6 +3926,49 @@ void CompilerGLSL::emit_resources() statement(""); emitted = false; + SmallVector spec_const_dependencies; + bool legacy_spec_constant_workgroup = execution.model == ExecutionModelGLCompute && !options.vulkan_semantics && + (execution.workgroup_size.constant != 0 || execution.flags.get( + ExecutionModeLocalSizeId)); + if (legacy_spec_constant_workgroup) + { + SpecializationConstant wg_x, wg_y, wg_z; + get_work_group_size_specialization_constants(wg_x, wg_y, wg_z); + + if (wg_x.id != ConstantID(0)) + spec_const_dependencies.push_back(wg_x.id); + if (wg_y.id != ConstantID(0)) + spec_const_dependencies.push_back(wg_y.id); + if (wg_z.id != ConstantID(0)) + spec_const_dependencies.push_back(wg_z.id); + } + + const auto notify_spec_constant = [&](ConstantID id) + { + if (legacy_spec_constant_workgroup) + { + auto itr = std::find(spec_const_dependencies.begin(), spec_const_dependencies.end(), id); + + if (itr == spec_const_dependencies.end()) + return; + + spec_const_dependencies.erase(itr); + if (spec_const_dependencies.empty()) + { + SpecializationConstant wg_x, wg_y, wg_z; + // We have declared all dependencies. We must delcare the workgroup size immediately + // as subsequent spec constant ops may depend on the declaration. + // Newer glslang does not allow gl_WorkGroupSize to be accessed before layout(local_size) in; + get_work_group_size_specialization_constants(wg_x, wg_y, wg_z); + SmallVector inputs; + build_workgroup_size(inputs, wg_x, wg_y, wg_z); + statement("layout(", merge(inputs), ") in;"); + statement(""); + legacy_spec_constant_workgroup = false; + } + } + }; + // If emitted Vulkan GLSL, // emit specialization constants as actual floats, // spec op expressions will redirect to the constant name. @@ -3812,11 +3999,15 @@ void CompilerGLSL::emit_resources() emit_constant(c); emitted = true; } + + if (c.specialization) + notify_spec_constant(ConstantID(c.self)); } else if (id.get_type() == TypeConstantOp) { emit_specialization_constant_op(id.get()); emitted = true; + notify_spec_constant(ConstantID(id.get_id())); } else if (id.get_type() == TypeType) { @@ -3871,24 +4062,6 @@ void CompilerGLSL::emit_resources() if (emitted) statement(""); - // If we needed to declare work group size late, check here. - // If the work group size depends on a specialization constant, we need to declare the layout() block - // after constants (and their macros) have been declared. - if (execution.model == ExecutionModelGLCompute && !options.vulkan_semantics && - (execution.workgroup_size.constant != 0 || execution.flags.get(ExecutionModeLocalSizeId))) - { - SpecializationConstant wg_x, wg_y, wg_z; - get_work_group_size_specialization_constants(wg_x, wg_y, wg_z); - - if ((wg_x.id != ConstantID(0)) || (wg_y.id != ConstantID(0)) || (wg_z.id != ConstantID(0))) - { - SmallVector inputs; - build_workgroup_size(inputs, wg_x, wg_y, wg_z); - statement("layout(", merge(inputs), ") in;"); - statement(""); - } - } - emitted = false; if (ir.addressing_model == AddressingModelPhysicalStorageBuffer64) @@ -4062,6 +4235,45 @@ void CompilerGLSL::emit_resources() } } + bool needs_hlsl_warning = false; + + for (const auto &heap_type : descriptor_heap_types) + { + auto &type = get(heap_type.type); + + if (heap_type.hlsl_style_stride) + needs_hlsl_warning = true; + + if (type.basetype == SPIRType::Image || type.basetype == SPIRType::AccelerationStructure) + { + string type_layout; + + // We lose NonWritable / NonReadable information, glslang issue, no good way to plumb it through either ... + if (type.basetype == SPIRType::Image && type.image.sampled == 2 && type.image.format != ImageFormatUnknown) + type_layout = join("layout(", to_descriptor_heap_layout(type), ", ", format_to_glsl(type.image.format), ") uniform "); + else + type_layout = join("layout(", to_descriptor_heap_layout(type), ") uniform "); + + statement(type_layout, variable_decl(type, join("spv", to_name(type.self), "ResourceHeap")), "[];"); + } + else if (type.basetype == SPIRType::Sampler) + { + statement("layout(", to_descriptor_heap_layout(type), ") uniform ", + variable_decl(type, join("spv", to_name(type.self), "SamplerHeap")), "[];"); + } + else + { + emit_buffer_block_native(nullptr, &heap_type); + } + } + + if (needs_hlsl_warning) + { + statement("// WARNING: HLSL style descriptor heap stride is assumed for one or more descriptors. Allowing for compatibility with HLSL shaders."); + statement("// This may be not strictly be compatible with GLSL if sizeof(buffer) != sizeof(image)."); + statement("// Application side can convert bindless indices accordingly to compensate or use explicit mapping API to configure strides outside SPIRV-Cross."); + } + if (emitted) statement(""); } @@ -4151,19 +4363,19 @@ void CompilerGLSL::emit_output_variable_initializer(const SPIRVariable &var) if (type_is_array && !is_control_point) { uint32_t indices[2] = { j, i }; - auto chain = access_chain_internal(var.self, indices, 2, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &meta); + auto chain = access_chain_internal(var.self, indices, 2, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &meta, nullptr); statement(chain, " = ", lut_name, "[", j, "];"); } else if (is_control_point) { uint32_t indices[2] = { invocation_id, member_index_id }; - auto chain = access_chain_internal(var.self, indices, 2, 0, &meta); + auto chain = access_chain_internal(var.self, indices, 2, 0, &meta, nullptr); statement(chain, " = ", lut_name, "[", builtin_to_glsl(BuiltInInvocationId, StorageClassInput), "];"); } else { auto chain = - access_chain_internal(var.self, &i, 1, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &meta); + access_chain_internal(var.self, &i, 1, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &meta, nullptr); statement(chain, " = ", to_expression(c.subconstants[i]), ";"); } @@ -5527,7 +5739,7 @@ string CompilerGLSL::to_non_uniform_aware_expression(uint32_t id) { string expr = to_expression(id); - if (has_decoration(id, DecorationNonUniform)) + if (is_descriptor_non_uniform(id)) convert_non_uniform_expression(expr, id); return expr; @@ -5919,7 +6131,7 @@ string CompilerGLSL::constant_op_expression(const SPIRConstantOp &cop) else { expr = access_chain_internal(cop.arguments[0], &cop.arguments[1], uint32_t(cop.arguments.size() - 1), - ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, nullptr); + ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, nullptr, nullptr); } return expr; } @@ -8082,7 +8294,7 @@ std::string CompilerGLSL::to_texture_op(const Instruction &i, bool sparse, bool auto &result_type = get(result_type_id); inherited_expressions.push_back(coord); - if (has_decoration(img, DecorationNonUniform) && !maybe_get_backing_variable(img)) + if (is_descriptor_non_uniform(img) && !maybe_get_backing_variable(img)) nonuniform_expression = true; switch (op) @@ -8136,6 +8348,9 @@ std::string CompilerGLSL::to_texture_op(const Instruction &i, bool sparse, bool case OpImageFetch: case OpImageSparseFetch: + if (options.vulkan_semantics && !dummy_sampler_id && (op == OpImageFetch || op == OpImageSparseFetch)) + require_extension_internal("GL_EXT_samplerless_texture_functions"); + // fallthrough case OpImageRead: // Reads == fetches in Metal (other langs will not get here) opt = &ops[4]; length -= 4; @@ -8236,6 +8451,17 @@ std::string CompilerGLSL::to_texture_op(const Instruction &i, bool sparse, bool base_args.is_proj = proj != 0; string expr; + + // texture() with bias on sampler2DArrayShadow or samplerCubeArrayShadow requires GL_EXT_texture_shadow_lod. + // textureOffset() with bias on sampler2DArrayShadow also requires it. + if (bias != 0 && dref != 0 && !fetch && !gather && + ((imgtype.image.arrayed && imgtype.image.dim == Dim2D) || + (imgtype.image.arrayed && imgtype.image.dim == DimCube)) && + is_depth_image(imgtype, img)) + { + require_extension_internal("GL_EXT_texture_shadow_lod"); + } + TextureFunctionNameArguments name_args = {}; name_args.base = base_args; @@ -8329,20 +8555,25 @@ bool CompilerGLSL::expression_is_constant_null(uint32_t id) const return c->constant_is_null(); } -bool CompilerGLSL::expression_is_non_value_type_array(uint32_t ptr) +bool CompilerGLSL::expression_is_non_value_type_array(uint32_t value_type_id, uint32_t ptr) { - auto &type = expression_type(ptr); - if (!is_array(get_pointee_type(type))) + auto &type = get(value_type_id); + if (!is_array(type)) return false; if (!backend.array_is_value_type) return true; + if (!backend.array_is_value_type_in_buffer_blocks && maybe_get_backing_buffer_pointer(ptr)) + return true; + auto *var = maybe_get_backing_variable(ptr); if (!var) return false; auto &backed_type = get(var->basetype); + + // Only consider explicitly laid out types here, not IO blocks. return !backend.array_is_value_type_in_buffer_blocks && backed_type.basetype == SPIRType::Struct && has_member_decoration(backed_type.self, 0, DecorationOffset); } @@ -8372,10 +8603,12 @@ string CompilerGLSL::to_function_name(const TextureFunctionNameArguments &args) { if (!expression_is_constant_null(args.lod)) { - SPIRV_CROSS_THROW("textureLod on sampler2DArrayShadow is not constant 0.0. This cannot be " - "expressed in GLSL."); + require_extension_internal("GL_EXT_texture_shadow_lod"); + } + else + { + workaround_lod_array_shadow_as_grad = true; } - workaround_lod_array_shadow_as_grad = true; } if (args.is_sparse_feedback) @@ -8510,9 +8743,11 @@ string CompilerGLSL::to_function_args(const TextureFunctionArguments &args, bool // To emulate this, we will have to use textureGrad with a constant gradient of 0. // The workaround will assert that the LOD is in fact constant 0, or we cannot emit correct code. // This happens for HLSL SampleCmpLevelZero on Texture2DArray and TextureCube. + // If GL_EXT_texture_shadow_lod is in use, textureLod is available directly with arbitrary LOD. bool workaround_lod_array_shadow_as_grad = ((imgtype.image.arrayed && imgtype.image.dim == Dim2D) || imgtype.image.dim == DimCube) && - is_depth_image(imgtype, img) && args.lod != 0 && !args.base.is_fetch; + is_depth_image(imgtype, img) && args.lod != 0 && !args.base.is_fetch && + !has_extension("GL_EXT_texture_shadow_lod"); if (args.dref) { @@ -10515,6 +10750,15 @@ string CompilerGLSL::builtin_to_glsl(BuiltIn builtin, StorageClass storage) return "gl_ClusterIDNV"; } + case BuiltInResourceHeapEXT: + // This builtin name is a placeholder. + // We will override this name later with prefix per actual type. + // However, this allows untyped access chain to index into the heap directly. + return "ResourceHeap"; + + case BuiltInSamplerHeapEXT: + return "SamplerHeap"; + default: return join("gl_BuiltIn_", convert_to_string(builtin)); } @@ -10576,7 +10820,8 @@ bool CompilerGLSL::access_chain_needs_stage_io_builtin_translation(uint32_t) } string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indices, uint32_t count, - AccessChainFlags flags, AccessChainMeta *meta) + AccessChainFlags flags, AccessChainMeta *meta, + const SPIRType *untyped_data_type) { string expr; @@ -10602,7 +10847,12 @@ string CompilerGLSL::access_chain_internal(uint32_t base, const uint32_t *indice // Start traversing type hierarchy at the proper non-pointer types, // but keep type_id referencing the original pointer for use below. uint32_t type_id = expression_type_id(base); - const auto *type = &get_pointee_type(type_id); + + // If nullptr we're doing untyped pointers. + // For now we don't really care about types since we're just doing a single index into the heap. + // If we intend to support complete untyped pointers usage later, we need to pass down the base type + // and override chain type based on that. + const auto *type = untyped_data_type ? untyped_data_type : &get_pointee_type(type_id); if (!backend.native_pointers) { @@ -11186,10 +11436,13 @@ uint32_t CompilerGLSL::get_physical_type_id_stride(TypeID) const } string CompilerGLSL::access_chain(uint32_t base, const uint32_t *indices, uint32_t count, const SPIRType &target_type, - AccessChainMeta *meta, bool ptr_chain) + AccessChainMeta *meta, bool ptr_chain, const SPIRType *untyped_data_type) { if (flattened_buffer_blocks.count(base)) { + if (untyped_data_type) + SPIRV_CROSS_THROW("Flattening not compatible with untyped pointers."); + uint32_t matrix_stride = 0; uint32_t array_stride = 0; bool need_transpose = false; @@ -11207,6 +11460,9 @@ string CompilerGLSL::access_chain(uint32_t base, const uint32_t *indices, uint32 } else if (flattened_structs.count(base) && count > 0) { + if (untyped_data_type) + SPIRV_CROSS_THROW("Flattening not compatible with untyped pointers."); + AccessChainFlags flags = ACCESS_CHAIN_CHAIN_ONLY_BIT | ACCESS_CHAIN_SKIP_REGISTER_EXPRESSION_READ_BIT; if (ptr_chain) flags |= ACCESS_CHAIN_PTR_CHAIN_BIT; @@ -11218,7 +11474,7 @@ string CompilerGLSL::access_chain(uint32_t base, const uint32_t *indices, uint32 meta->flattened_struct = target_type.basetype == SPIRType::Struct; } - auto chain = access_chain_internal(base, indices, count, flags, nullptr).substr(1); + auto chain = access_chain_internal(base, indices, count, flags, nullptr, nullptr).substr(1); if (meta) { meta->need_transpose = false; @@ -11255,7 +11511,7 @@ string CompilerGLSL::access_chain(uint32_t base, const uint32_t *indices, uint32 } } - return access_chain_internal(base, indices, count, flags, meta); + return access_chain_internal(base, indices, count, flags, meta, untyped_data_type); } } @@ -11755,6 +12011,9 @@ bool CompilerGLSL::should_forward(uint32_t id) const if (is_immutable(id)) return true; + if (expr && expr->buffer_pointer) + return true; + return false; } @@ -11843,6 +12102,8 @@ void CompilerGLSL::register_impure_function_call() flush_dependees(get(global)); for (auto aliased : aliased_variables) flush_dependees(get(aliased)); + for (auto ptr : buffer_pointer_variables) + flush_dependees(get(ptr)); } void CompilerGLSL::register_call_out_argument(uint32_t id) @@ -12258,7 +12519,7 @@ void CompilerGLSL::emit_store_statement(uint32_t lhs_expression, uint32_t rhs_ex if (!unroll_array_to_complex_store(lhs_expression, rhs_expression)) { auto lhs = to_dereferenced_expression(lhs_expression); - if (has_decoration(lhs_expression, DecorationNonUniform)) + if (is_descriptor_non_uniform(lhs_expression)) convert_non_uniform_expression(lhs, lhs_expression); // We might need to cast in order to store to a builtin. @@ -12476,6 +12737,18 @@ static bool opcode_is_precision_sensitive_operation(Op op) case OpConvertUToF: case OpConvertFToU: case OpConvertFToS: + case OpShiftLeftLogical: + case OpShiftRightLogical: + case OpShiftRightArithmetic: + case OpBitwiseOr: + case OpBitwiseXor: + case OpBitwiseAnd: + case OpNot: + case OpBitFieldInsert: + case OpBitFieldSExtract: + case OpBitFieldUExtract: + case OpBitReverse: + case OpBitCount: return true; default: @@ -12649,6 +12922,14 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) // If an expression is mutable and forwardable, we speculate that it is immutable. bool forward = should_forward(ptr) && forced_temporaries.find(id) == end(forced_temporaries); + // Volatile memory access requires the value be read exactly once from + // memory. Do not forward the expression so that re-evaluation at each + // use site cannot re-read potentially modified memory. + // FIXME: To force implementations to actually respect the volatile nature of the load, + // the block itself must be marked volatile, or VulkanMM is used to do an explicit volatile load. + if (forward && length >= 4 && (ops[3] & MemoryAccessVolatileMask) != 0) + forward = false; + // If loading a non-native row-major matrix, mark the expression as need_transpose. bool need_transpose = false; bool old_need_transpose = false; @@ -12712,7 +12993,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) // Also, loading from gl_SampleMask array needs special unroll. unroll_array_from_complex_load(id, ptr, expr); - if (!type_is_opaque_value(type) && has_decoration(ptr, DecorationNonUniform)) + if (!type_is_opaque_value(type) && is_descriptor_non_uniform(ptr)) { // If we're loading something non-opaque, we need to handle non-uniform descriptor access. convert_non_uniform_expression(expr, ptr); @@ -12732,7 +13013,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) bool usage_tracking = flattened && (type.basetype == SPIRType::Struct || (type.columns > 1)); SPIRExpression *e = nullptr; - if (!forward && expression_is_non_value_type_array(ptr)) + if (!forward && expression_is_non_value_type_array(result_type, ptr)) { // Complicated load case where we need to make a copy of ptr, but we cannot, because // it is an array, and our backend does not support arrays as value types. @@ -12770,11 +13051,41 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) break; } + case OpUntypedPtrAccessChainKHR: + SPIRV_CROSS_THROW("OpUntypedPtrAccessChainKHR is not supported."); + break; + + case OpUntypedAccessChainKHR: + case OpUntypedInBoundsAccessChainKHR: case OpInBoundsAccessChain: case OpAccessChain: case OpPtrAccessChain: { - auto *var = maybe_get(ops[2]); + bool untyped = opcode == OpUntypedAccessChainKHR || opcode == OpUntypedInBoundsAccessChainKHR; + + uint32_t type_id = ops[0]; + uint32_t result_id = ops[1]; + uint32_t ptr_id = ops[untyped ? 3 : 2]; + uint32_t indices_start = untyped ? 4 : 3; + + if (untyped) + { + auto *var = maybe_get_backing_variable(ptr_id); + // Buffer pointers stop the loaded from chain to deal with aliasing better, so carve that out specifically. + auto *expr = maybe_get_backing_buffer_pointer(ptr_id); + + if (!expr) + { + if (!var || !has_decoration(var->self, DecorationBuiltIn) || + (BuiltIn(get_decoration(var->self, DecorationBuiltIn)) != BuiltInResourceHeapEXT && + BuiltIn(get_decoration(var->self, DecorationBuiltIn)) != BuiltInSamplerHeapEXT)) + { + SPIRV_CROSS_THROW("Untyped pointer access chains are currently only supported for descriptor heap access."); + } + } + } + + auto *var = maybe_get(ptr_id); if (var) flush_variable_declaration(var->self); @@ -12782,57 +13093,78 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) // If an expression is mutable and forwardable, we speculate that it is immutable. AccessChainMeta meta; bool ptr_chain = opcode == OpPtrAccessChain; - auto &target_type = get(ops[0]); - auto e = access_chain(ops[2], &ops[3], length - 3, target_type, &meta, ptr_chain); + auto &target_type = get(type_id); + auto e = access_chain(ptr_id, &ops[indices_start], length - indices_start, target_type, &meta, ptr_chain, + untyped ? &get(ops[2]) : nullptr); + + if (untyped) + { + auto &data_type = get(ops[2]); + auto *ptr_expr = maybe_get(ptr_id); + if (data_type.basetype == SPIRType::Image || data_type.basetype == SPIRType::Sampler || + data_type.basetype == SPIRType::AccelerationStructure || + (ptr_expr && ptr_expr->buffer_pointer)) + { + // We can resolve this type now. + // For further buffer access chains, we don't do any fixups since we have resolved to proper types. + // For buffer types we only prepend when the access chain starts from a BufferPointerEXT base. + // Multi-stage access chains are not possible for image types. + if (ptr_expr && ptr_expr->buffer_pointer) + e = join(to_buffer_pointer_name_prefix(ptr_expr->self), e); + else + e = join("spv", to_name(data_type.self), e); + } + } // If the base is flattened UBO of struct type, the expression has to be a composite. // In that case, backends which do not support inline syntax need it to be bound to a temporary. // Otherwise, invalid expressions like ({UBO[0].xyz, UBO[0].w, UBO[1]}).member are emitted. bool requires_temporary = false; - if (flattened_buffer_blocks.count(ops[2]) && target_type.basetype == SPIRType::Struct) + if (flattened_buffer_blocks.count(ptr_id) && target_type.basetype == SPIRType::Struct) requires_temporary = !backend.can_declare_struct_inline; auto &expr = requires_temporary ? - emit_op(ops[0], ops[1], std::move(e), false) : - set(ops[1], std::move(e), ops[0], should_forward(ops[2])); + emit_op(type_id, result_id, std::move(e), false) : + set(result_id, std::move(e), type_id, should_forward(ptr_id)); - auto *backing_variable = maybe_get_backing_variable(ops[2]); - expr.loaded_from = backing_variable ? backing_variable->self : ID(ops[2]); + auto *backing_variable = maybe_get_backing_variable(ptr_id); + expr.loaded_from = backing_variable ? backing_variable->self : ID(ptr_id); expr.need_transpose = meta.need_transpose; expr.access_chain = true; expr.access_meshlet_position_y = meta.access_meshlet_position_y; // Mark the result as being packed. Some platforms handled packed vectors differently than non-packed. if (meta.storage_is_packed) - set_extended_decoration(ops[1], SPIRVCrossDecorationPhysicalTypePacked); + set_extended_decoration(result_id, SPIRVCrossDecorationPhysicalTypePacked); if (meta.storage_physical_type != 0) - set_extended_decoration(ops[1], SPIRVCrossDecorationPhysicalTypeID, meta.storage_physical_type); + set_extended_decoration(result_id, SPIRVCrossDecorationPhysicalTypeID, meta.storage_physical_type); if (meta.storage_is_invariant) - set_decoration(ops[1], DecorationInvariant); + set_decoration(result_id, DecorationInvariant); if (meta.flattened_struct) - flattened_structs[ops[1]] = true; + flattened_structs[result_id] = true; if (meta.relaxed_precision && backend.requires_relaxed_precision_analysis) - set_decoration(ops[1], DecorationRelaxedPrecision); + set_decoration(result_id, DecorationRelaxedPrecision); if (meta.chain_is_builtin) - set_decoration(ops[1], DecorationBuiltIn, meta.builtin); + set_decoration(result_id, DecorationBuiltIn, meta.builtin); // If we have some expression dependencies in our access chain, this access chain is technically a forwarded // temporary which could be subject to invalidation. // Need to assume we're forwarded while calling inherit_expression_depdendencies. - forwarded_temporaries.insert(ops[1]); + forwarded_temporaries.insert(result_id); // The access chain itself is never forced to a temporary, but its dependencies might. - suppressed_usage_tracking.insert(ops[1]); + suppressed_usage_tracking.insert(result_id); - for (uint32_t i = 2; i < length; i++) + // Include the base pointer. + for (uint32_t i = indices_start - 1; i < length; i++) { - inherit_expression_dependencies(ops[1], ops[i]); + inherit_expression_dependencies(result_id, ops[i]); add_implied_read_expression(expr, ops[i]); } // If we have no dependencies after all, i.e., all indices in the access chain are immutable temporaries, // we're not forwarded after all. if (expr.expression_dependencies.empty()) - forwarded_temporaries.erase(ops[1]); + forwarded_temporaries.erase(result_id); break; } @@ -12866,15 +13198,80 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) break; } + case OpUntypedArrayLengthKHR: case OpArrayLength: { + bool untyped = opcode == OpUntypedArrayLengthKHR; uint32_t result_type = ops[0]; uint32_t id = ops[1]; - auto e = access_chain_internal(ops[2], &ops[3], length - 3, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, nullptr); - if (has_decoration(ops[2], DecorationNonUniform)) - convert_non_uniform_expression(e, ops[2]); - set(id, join(type_to_glsl(get(result_type)), "(", e, ".length())"), result_type, - true); + + const SPIRType *untyped_data_type = untyped ? &get(ops[2]) : nullptr; + uint32_t ptr_id = ops[untyped ? 3 : 2]; + uint32_t index_offset = untyped ? 4 : 3; + + auto e = access_chain_internal(ptr_id, &ops[index_offset], length - index_offset, + ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, + nullptr, untyped_data_type); + + if (untyped) + { + auto *ptr_expr = maybe_get(ptr_id); + if (ptr_expr && ptr_expr->buffer_pointer) + e = join(to_buffer_pointer_name_prefix(ptr_expr->self), e); + } + + if (is_descriptor_non_uniform(ptr_id)) + convert_non_uniform_expression(e, ptr_id); + set(id, join(type_to_glsl(get(result_type)), "(", e, ".length())"), result_type, true); + break; + } + + case OpBufferPointerEXT: + { + uint32_t type_id = ops[0]; + uint32_t result_id = ops[1]; + uint32_t ptr_id = ops[2]; + + auto *backing_variable = maybe_get_backing_variable(ptr_id); + if (!backing_variable) + SPIRV_CROSS_THROW("There is no backing variable for BufferPointerEXT."); + + auto *chain_expr = maybe_get(ptr_id); + if (!chain_expr || !chain_expr->access_chain) + SPIRV_CROSS_THROW("Expected to see access chain for BufferPointerEXT."); + + auto e = to_expression(ptr_id); + + // BufferPointerEXT can return a typed pointer, in which case we need to resolve the heap alias now. + auto &type = get(type_id); + if (type.basetype == SPIRType::Struct) + e = join(to_buffer_pointer_name_prefix(result_id), e); + + auto &expr = set(result_id, std::move(e), type_id, true); + // There isn't any backing variable here. OpBufferPointerEXT is meant to be a memory declaration instruction. + expr.loaded_from = 0; + expr.access_chain = true; + expr.buffer_pointer = true; + expr.implied_read_expressions = chain_expr->implied_read_expressions; + expr.expression_dependencies = chain_expr->expression_dependencies; + expr.immutable = false; + + // If the buffer pointer is marked non-writable, ignore alias tracking by flagging the expression as immutable. + for (auto &heap : descriptor_heap_types) + { + if (heap.buffer_pointer_id == result_id) + { + if (heap.nonwritable) + expr.immutable = true; + break; + } + } + + if (!expr.immutable && ir.get_buffer_block_type_flags(get(type_id)).get(DecorationNonWritable)) + expr.immutable = true; + + // Used for load-store tracking. + buffer_pointer_variables.push_back(result_id); break; } @@ -13108,7 +13505,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) // Make a copy, then use access chain to store the variable. statement(declare_temporary(result_type, id), to_expression(vec), ";"); set(id, to_name(id), result_type, true); - auto chain = access_chain_internal(id, &index, 1, 0, nullptr); + auto chain = access_chain_internal(id, &index, 1, 0, nullptr, nullptr); statement(chain, " = ", to_unpacked_expression(comp), ";"); break; } @@ -13118,7 +13515,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) uint32_t result_type = ops[0]; uint32_t id = ops[1]; - auto expr = access_chain_internal(ops[2], &ops[3], 1, 0, nullptr); + auto expr = access_chain_internal(ops[2], &ops[3], 1, 0, nullptr, nullptr); emit_op(result_type, id, expr, should_forward(ops[2])); inherit_expression_dependencies(id, ops[2]); inherit_expression_dependencies(id, ops[3]); @@ -13182,7 +13579,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) // from expression causing it to be forced to an actual temporary in GLSL. auto expr = access_chain_internal(ops[2], &ops[3], length, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_CHAIN_ONLY_BIT | - ACCESS_CHAIN_FORCE_COMPOSITE_BIT, &meta); + ACCESS_CHAIN_FORCE_COMPOSITE_BIT, &meta, nullptr); e = &emit_op(result_type, id, expr, true, should_suppress_usage_tracking(ops[2])); inherit_expression_dependencies(id, ops[2]); e->base_expression = ops[2]; @@ -13193,7 +13590,8 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) else { auto expr = access_chain_internal(ops[2], &ops[3], length, - ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_FORCE_COMPOSITE_BIT, &meta); + ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_FORCE_COMPOSITE_BIT, + &meta, nullptr); e = &emit_op(result_type, id, expr, should_forward(ops[2]), should_suppress_usage_tracking(ops[2])); inherit_expression_dependencies(id, ops[2]); } @@ -13261,7 +13659,8 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) if (!forced_temporaries.count(composite)) force_temporary_and_recompile(composite); - auto chain = access_chain_internal(composite, elems, length, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, nullptr); + auto chain = access_chain_internal(composite, elems, length, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, + nullptr, nullptr); statement(chain, " = ", to_unpacked_expression(obj), ";"); set(id, to_expression(composite), result_type, true); invalid_expressions.insert(composite); @@ -13280,7 +13679,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) set(id, to_name(id), result_type, true); } - auto chain = access_chain_internal(id, elems, length, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, nullptr); + auto chain = access_chain_internal(id, elems, length, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, nullptr, nullptr); statement(chain, " = ", to_unpacked_expression(obj), ";"); } @@ -14574,7 +14973,7 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) op = "textureQueryLod"; auto sampler_expr = to_expression(ops[2]); - if (has_decoration(ops[2], DecorationNonUniform)) + if (is_descriptor_non_uniform(ops[2])) { if (maybe_get_backing_variable(ops[2])) convert_non_uniform_expression(sampler_expr, ops[2]); @@ -14849,23 +15248,28 @@ void CompilerGLSL::emit_instruction(const Instruction &instruction) break; } + case OpUntypedImageTexelPointerEXT: case OpImageTexelPointer: { + bool untyped = opcode == OpUntypedImageTexelPointerEXT; uint32_t result_type = ops[0]; uint32_t id = ops[1]; - auto coord_expr = to_expression(ops[3]); - auto target_coord_type = expression_type(ops[3]); - target_coord_type.basetype = SPIRType::Int; - coord_expr = bitcast_expression(target_coord_type, expression_type(ops[3]).basetype, coord_expr); + uint32_t image_id = ops[untyped ? 3 : 2]; + uint32_t coord_id = ops[untyped ? 4 : 3]; - auto expr = join(to_expression(ops[2]), ", ", coord_expr); + auto coord_expr = to_expression(coord_id); + auto target_coord_type = expression_type(coord_id); + target_coord_type.basetype = SPIRType::Int; + coord_expr = bitcast_expression(target_coord_type, expression_type(coord_id).basetype, coord_expr); + + auto expr = join(to_expression(image_id), ", ", coord_expr); auto &e = set(id, expr, result_type, true); // When using the pointer, we need to know which variable it is actually loaded from. - auto *var = maybe_get_backing_variable(ops[2]); + auto *var = maybe_get_backing_variable(image_id); e.loaded_from = var ? var->self : ID(0); - inherit_expression_dependencies(id, ops[3]); + inherit_expression_dependencies(id, coord_id); break; } @@ -17368,7 +17772,12 @@ void CompilerGLSL::add_function_overload(const SPIRFunction &func) void CompilerGLSL::emit_function_prototype(SPIRFunction &func, const Bitset &return_flags) { - if (func.self != ir.default_entry_point) + // In library mode default_entry_point points at the first exported + // function; treat every export as a normal function rather than as the + // shader's entry point. + const bool is_entry_point = !ir.is_library_module && func.self == ir.default_entry_point; + + if (!is_entry_point) add_function_overload(func); // Avoid shadow declarations. @@ -17382,7 +17791,7 @@ void CompilerGLSL::emit_function_prototype(SPIRFunction &func, const Bitset &ret decl += type_to_array_glsl(type, 0); decl += " "; - if (func.self == ir.default_entry_point) + if (is_entry_point) { // If we need complex fallback in GLSL, we just wrap main() in a function // and interlock the entire shader ... @@ -17920,13 +18329,30 @@ string CompilerGLSL::emit_continue_block(uint32_t continue_block, bool follow_tr return merge(statements); } +// Loop variable with OpUndef init: zero-init instead of leaving uninitialized (FXC X4555/X4000). +std::string CompilerGLSL::undef_loop_variable_initializer_suffix(const SPIRVariable &var) +{ + if (!backend.requires_phi_undef_zero_init) + return ""; + + uint32_t expr = var.static_expression; + if (expr == 0 || ir.ids[expr].get_type() != TypeUndef) + return ""; + + auto &type = get(var.basetype); + if (!type_can_zero_initialize(type)) + return ""; + + return join(" = ", to_zero_initialized_expression(var.basetype)); +} + void CompilerGLSL::emit_while_loop_initializers(const SPIRBlock &block) { // While loops do not take initializers, so declare all of them outside. for (auto &loop_var : block.loop_variables) { auto &var = get(loop_var); - statement(variable_decl(var), ";"); + statement(variable_decl(var), undef_loop_variable_initializer_suffix(var), ";"); } } @@ -17958,7 +18384,10 @@ string CompilerGLSL::emit_for_loop_initializers(const SPIRBlock &block) else if (!same_types || missing_initializers == uint32_t(block.loop_variables.size())) { for (auto &loop_var : block.loop_variables) - statement(variable_decl(get(loop_var)), ";"); + { + auto &var = get(loop_var); + statement(variable_decl(var), undef_loop_variable_initializer_suffix(var), ";"); + } return ""; } else @@ -17969,10 +18398,11 @@ string CompilerGLSL::emit_for_loop_initializers(const SPIRBlock &block) for (auto &loop_var : block.loop_variables) { - uint32_t static_expr = get(loop_var).static_expression; + auto &var_for_undef = get(loop_var); + uint32_t static_expr = var_for_undef.static_expression; if (static_expr == 0 || ir.ids[static_expr].get_type() == TypeUndef) { - statement(variable_decl(get(loop_var)), ";"); + statement(variable_decl(var_for_undef), undef_loop_variable_initializer_suffix(var_for_undef), ";"); } else { @@ -18226,7 +18656,11 @@ void CompilerGLSL::emit_hoisted_temporaries(SmallVector> &tempo // There are some rare scenarios where we are asked to declare pointer types as hoisted temporaries. // This should be ignored unless we're doing actual variable pointers and backend supports it. // Access chains cannot normally be lowered to temporaries in GLSL and HLSL. - if (type.pointer && !backend.native_pointers) + if (type.pointer && (!backend.native_pointers || type_is_opaque_value(get_pointee_type(type)))) + continue; + + // Anything involving opaque objects cannot be lowered to temporaries ever. + if (type_is_opaque_value(type)) continue; add_local_variable_name(tmp.second); @@ -19348,16 +19782,28 @@ void CompilerGLSL::convert_non_uniform_expression(string &expr, uint32_t ptr_id) return; auto *var = maybe_get_backing_variable(ptr_id); - if (!var) + auto *buffer_pointer = maybe_get_backing_buffer_pointer(ptr_id); + if (!var && !buffer_pointer) return; - if (var->storage != StorageClassUniformConstant && + if (!buffer_pointer && + var->storage != StorageClassUniformConstant && var->storage != StorageClassStorageBuffer && var->storage != StorageClassUniform) return; - auto &backing_type = get(var->basetype); - if (backing_type.array.empty()) + auto &backing_type = get(var ? var->basetype : buffer_pointer->expression_type); + + bool descriptor_heap = false; + if (var) + { + auto builtin = BuiltIn(get_decoration(var->self, DecorationBuiltIn)); + descriptor_heap = builtin == BuiltInResourceHeapEXT || builtin == BuiltInSamplerHeapEXT; + } + else if (buffer_pointer) + descriptor_heap = true; + + if (!descriptor_heap && backing_type.array.empty()) return; // If we get here, we know we're accessing an arrayed resource which @@ -19646,9 +20092,9 @@ void CompilerGLSL::emit_copy_logical_type(uint32_t lhs_id, uint32_t lhs_type_id, AccessChainMeta lhs_meta, rhs_meta; auto lhs = access_chain_internal(lhs_id, chain.data(), uint32_t(chain.size()), - ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &lhs_meta); + ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &lhs_meta, nullptr); auto rhs = access_chain_internal(rhs_id, chain.data(), uint32_t(chain.size()), - ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &rhs_meta); + ACCESS_CHAIN_INDEX_IS_LITERAL_BIT, &rhs_meta, nullptr); uint32_t id = ir.increase_bound_by(2); lhs_id = id; @@ -20338,3 +20784,130 @@ bool CompilerGLSL::has_legacy_nocontract(uint32_t result_type, uint32_t id) cons FPFastMathModeAllowReassocMask; return (get_fp_fast_math_flags_for_op(result_type, id) & fp_flags) != fp_flags; } + +void CompilerGLSL::remap_descriptor_heap(ResourceType type, uint32_t desc_set, uint32_t binding, Dim dim) +{ + for (auto &mapping : descriptor_heap_mappings) + { + if (mapping.type == type) + { + mapping.desc_set = desc_set; + mapping.binding = binding; + mapping.dim = dim; + return; + } + } + + descriptor_heap_mappings.push_back({ type, desc_set, binding, dim }); +} + +bool CompilerGLSL::is_descriptor_non_uniform(uint32_t id) const +{ + if (has_decoration(id, DecorationNonUniform)) + return true; + + // Only infer nonuniform for descriptors. + auto &type = expression_type(id); + + if (is_pointer(type)) + { + if (type.storage != StorageClassUniform && type.storage != StorageClassStorageBuffer && + type.storage != StorageClassUniformConstant && type.storage != StorageClassImage) + return false; + } + else if (!type_is_opaque_value(type)) + { + return false; + } + + if (descriptor_heap_mappings.empty()) + return false; + + if (has_decoration(id, DecorationUniform)) + return false; + + if (std::find(ir.declared_capabilities.begin(), ir.declared_capabilities.end(), + CapabilityDescriptorHeapEXT) == ir.declared_capabilities.end()) + return false; + + // Definitely not. + if (maybe_get(id) || maybe_get(id)) + return false; + + // DescriptorHeapEXT requires that nonuniformEXT is implied, + // but if we're remapping to legacy set/binding model, glslang will not emit the cap in cross compiled source, + // so we have to enforce it. We don't have compiler-infra to deduce subgroup uniformity statically, + // so just slap it on everything. Compilers generally figure this stuff out. + return true; +} + +std::string CompilerGLSL::to_descriptor_heap_layout(const SPIRType &type, StorageClass storage) const +{ + auto resource = ResourceTypeUnknown; + Dim dim = DimMax; + + switch (type.basetype) + { + case SPIRType::Sampler: + resource = ResourceTypeSeparateSamplers; + break; + + case SPIRType::Image: + dim = type.image.dim == DimBuffer ? DimBuffer : Dim2D; + resource = type.image.sampled == 2 ? ResourceTypeStorageImage : ResourceTypeSeparateImage; + break; + + case SPIRType::SampledImage: + resource = ResourceTypeSampledImage; + break; + + case SPIRType::AccelerationStructure: + resource = ResourceTypeAccelerationStructure; + break; + + case SPIRType::AtomicCounter: + resource = ResourceTypeAtomicCounter; + break; + + case SPIRType::Struct: + { + bool ssbo = storage == StorageClassStorageBuffer || has_decoration(type.self, DecorationBufferBlock); + resource = ssbo ? ResourceTypeStorageBuffer : ResourceTypeUniformBuffer; + break; + } + + default: + break; + } + + for (auto &mapping : descriptor_heap_mappings) + { + if (mapping.type == resource) + { + bool has_match = false; + + if (type.basetype == SPIRType::Image) + { + if (dim == DimBuffer && (mapping.dim == DimMax || mapping.dim == DimBuffer)) + has_match = true; + if (dim != DimBuffer && mapping.dim != DimBuffer) + has_match = true; + } + else + { + has_match = true; + } + + if (has_match) + return join("set = ", mapping.desc_set, ", binding = ", mapping.binding); + } + } + + // Fallback to unknown mapping. + for (auto &mapping : descriptor_heap_mappings) + if (mapping.type == ResourceTypeUnknown) + return join("set = ", mapping.desc_set, ", binding = ", mapping.binding); + + return "descriptor_heap"; +} + diff --git a/3rdparty/spirv-cross/spirv_glsl.hpp b/3rdparty/spirv-cross/spirv_glsl.hpp index d6c20247a..0b7590577 100644 --- a/3rdparty/spirv-cross/spirv_glsl.hpp +++ b/3rdparty/spirv-cross/spirv_glsl.hpp @@ -306,6 +306,19 @@ public: // Returns the macro name corresponding to constant id std::string constant_value_macro_name(uint32_t id) const; + // Rather than using layout(descriptor_heap), emit layout(set, binding). + // This intended to be compatible with descriptor buffers, legacy descriptor indexing, + // or when the heap descriptors require unusual kinds of mapping in the Vulkan API + // which is not expressible by GLSL directly. + // + // ResourceTypeUnknown can be used as a default catch-all mapping. + // dim can be used to disambiguate between texel buffers and images since they are both image types, + // but use different descriptor types in the Vulkan API. + // No distinction is made between 1D/2D/3D/Cube textures. + // The default argument of DimMax maps to both texel buffers and images. + // dim is ignored for ResourceTypeUnknown. + void remap_descriptor_heap(ResourceType type, uint32_t desc_set, uint32_t binding, Dim dim = DimMax); + protected: struct ShaderSubgroupSupportHelper { @@ -670,6 +683,7 @@ protected: bool requires_relaxed_precision_analysis = false; bool implicit_c_integer_promotion_rules = false; bool supports_spec_constant_array_size = true; + bool requires_phi_undef_zero_init = false; } backend; void emit_struct(SPIRType &type); @@ -677,7 +691,9 @@ protected: void emit_extension_workarounds(ExecutionModel model); void emit_subgroup_arithmetic_workaround(const std::string &func, Op op, GroupOperation group_op); void emit_polyfills(uint32_t polyfills, bool relaxed); - void emit_buffer_block_native(const SPIRVariable &var); + void emit_buffer_block_native(const SPIRVariable *var, const DescriptorHeapMeta *heap_meta = nullptr); + std::string to_buffer_pointer_name_prefix(uint32_t ptr_id) const; + static std::string heap_meta_to_prefix(const DescriptorHeapMeta &meta); 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); @@ -771,7 +787,7 @@ protected: AccessChainFlags flags, bool &access_chain_is_arrayed, uint32_t index); std::string access_chain_internal(uint32_t base, const uint32_t *indices, uint32_t count, AccessChainFlags flags, - AccessChainMeta *meta); + AccessChainMeta *meta, const SPIRType *untyped_data_type); // Only meaningful on backends with physical pointer support ala MSL. // Relevant for PtrAccessChain / BDA. @@ -785,7 +801,8 @@ protected: StorageClass storage, bool &is_packed); std::string access_chain(uint32_t base, const uint32_t *indices, uint32_t count, const SPIRType &target_type, - AccessChainMeta *meta = nullptr, bool ptr_chain = false); + AccessChainMeta *meta = nullptr, bool ptr_chain = false, + const SPIRType *untyped_data_type = nullptr); std::string flattened_access_chain(uint32_t base, const uint32_t *indices, uint32_t count, const SPIRType &target_type, uint32_t offset, uint32_t matrix_stride, @@ -1023,6 +1040,7 @@ protected: std::string emit_for_loop_initializers(const SPIRBlock &block); void emit_while_loop_initializers(const SPIRBlock &block); + std::string undef_loop_variable_initializer_suffix(const SPIRVariable &var); bool for_loop_initializers_are_same_type(const SPIRBlock &block); bool optimize_read_modify_write(const SPIRType &type, const std::string &lhs, const std::string &rhs); void fixup_image_load_store_access(); @@ -1055,7 +1073,7 @@ protected: void disallow_forwarding_in_expression_chain(const SPIRExpression &expr); bool expression_is_constant_null(uint32_t id) const; - bool expression_is_non_value_type_array(uint32_t ptr); + bool expression_is_non_value_type_array(uint32_t value_type_id, uint32_t ptr); virtual void emit_store_statement(uint32_t lhs_expression, uint32_t rhs_expression); uint32_t get_integer_width_for_instruction(const Instruction &instr) const; @@ -1089,6 +1107,17 @@ protected: uint32_t get_fp_fast_math_flags_for_op(uint32_t result_type, uint32_t id) const; bool has_legacy_nocontract(uint32_t result_type, uint32_t id) const; + struct DescriptorHeapMapping + { + ResourceType type; + uint32_t desc_set; + uint32_t binding; + Dim dim; + }; + SmallVector descriptor_heap_mappings; + bool is_descriptor_non_uniform(uint32_t id) const; + std::string to_descriptor_heap_layout(const SPIRType &type, StorageClass storage = StorageClassUniformConstant) const; + private: void init(); diff --git a/3rdparty/spirv-cross/spirv_hlsl.cpp b/3rdparty/spirv-cross/spirv_hlsl.cpp index f4e18b41c..5785c2394 100644 --- a/3rdparty/spirv-cross/spirv_hlsl.cpp +++ b/3rdparty/spirv-cross/spirv_hlsl.cpp @@ -3117,7 +3117,12 @@ uint32_t CompilerHLSL::input_vertices_from_execution_mode(SPIREntryPoint &execut void CompilerHLSL::emit_function_prototype(SPIRFunction &func, const Bitset &return_flags) { - if (func.self != ir.default_entry_point) + // In library mode default_entry_point points at the first exported + // function; treat every export as a normal function rather than as the + // shader's entry point. + const bool is_entry_point = !ir.is_library_module && func.self == ir.default_entry_point; + + if (!is_entry_point) add_function_overload(func); // Avoid shadow declarations. @@ -3138,7 +3143,7 @@ void CompilerHLSL::emit_function_prototype(SPIRFunction &func, const Bitset &ret decl = "void "; } - if (func.self == ir.default_entry_point) + if (is_entry_point) { decl += get_inner_entry_point_name(); processing_entry_point = true; @@ -5237,7 +5242,8 @@ string CompilerHLSL::write_access_chain_value(uint32_t value, const SmallVector< { AccessChainMeta meta; ret = access_chain_internal(value, composite_chain.data(), uint32_t(composite_chain.size()), - ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_LITERAL_MSB_FORCE_ID, &meta); + ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_LITERAL_MSB_FORCE_ID, &meta, + nullptr); } if (enclose) @@ -7101,6 +7107,7 @@ string CompilerHLSL::compile() backend.can_return_array = false; backend.nonuniform_qualifier = "NonUniformResourceIndex"; backend.support_case_fallthrough = false; + backend.requires_phi_undef_zero_init = true; backend.force_merged_mesh_block = get_execution_model() == ExecutionModelMeshEXT; backend.force_gl_in_out_block = backend.force_merged_mesh_block; backend.supports_empty_struct = hlsl_options.shader_model <= 30; @@ -7146,14 +7153,27 @@ string CompilerHLSL::compile() emit_header(); emit_resources(); - emit_function(get(ir.default_entry_point), Bitset()); - emit_hlsl_entry_point(); + if (ir.is_library_module) + { + // Emit each exported function as a normal free function. + // emit_function recursively emits callees, so internal helpers + // are picked up too. + for (auto export_id : ir.library_exported_functions) + emit_function(get(export_id), Bitset()); + } + else + { + emit_function(get(ir.default_entry_point), Bitset()); + emit_hlsl_entry_point(); + } pass_count++; } while (is_forcing_recompilation()); // Entry point in HLSL is always main() for the time being. - get_entry_point().name = "main"; + // Skip the rename for library modules; their exports keep their declared names. + if (!ir.is_library_module) + get_entry_point().name = "main"; return buffer.str(); } diff --git a/3rdparty/spirv-cross/spirv_msl.cpp b/3rdparty/spirv-cross/spirv_msl.cpp index 6e332cad0..bd32b37bf 100644 --- a/3rdparty/spirv-cross/spirv_msl.cpp +++ b/3rdparty/spirv-cross/spirv_msl.cpp @@ -1087,6 +1087,18 @@ void CompilerMSL::build_implicit_builtins() dynamic_offsets_buffer_id = var_id; } + if (active_input_builtins.get(BuiltInDrawIndex)) + { + // This is always emulated. + uint32_t var_id = build_constant_uint_array_pointer(); + set_name(var_id, "spvDrawIndex"); + // This should never match anything. + set_decoration(var_id, DecorationDescriptorSet, ~(6u)); + set_decoration(var_id, DecorationBinding, msl_options.draw_id_buffer_index); + set_extended_decoration(var_id, SPIRVCrossDecorationResourceIndexPrimary, msl_options.draw_id_buffer_index); + draw_index_buffer_id = var_id; + } + // If we're returning a struct from a vertex-like entry point, we must return a position attribute. bool need_position = (get_execution_model() == ExecutionModelVertex || is_tese_shader()) && !capture_output_to_buffer && !get_is_rasterization_disabled() && @@ -1766,6 +1778,8 @@ string CompilerMSL::compile() add_active_interface_variable(view_mask_buffer_id); if (dynamic_offsets_buffer_id) add_active_interface_variable(dynamic_offsets_buffer_id); + if (draw_index_buffer_id) + add_active_interface_variable(draw_index_buffer_id); if (builtin_layer_id) add_active_interface_variable(builtin_layer_id); if (builtin_dispatch_base_id && !msl_options.supports_msl_version(1, 2)) @@ -8727,7 +8741,7 @@ bool CompilerMSL::emit_tessellation_io_load(uint32_t result_type_id, uint32_t id const uint32_t indices[2] = { i, interface_index }; AccessChainMeta meta; expr += access_chain_internal(stage_in_ptr_var_id, indices, 2, - ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_PTR_CHAIN_BIT, &meta); + ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_PTR_CHAIN_BIT, &meta, nullptr); if (i + 1 < num_control_points) expr += ", "; } @@ -8763,7 +8777,8 @@ bool CompilerMSL::emit_tessellation_io_load(uint32_t result_type_id, uint32_t id AccessChainMeta meta; expr += access_chain_internal(stage_in_ptr_var_id, indices, 2, - ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_PTR_CHAIN_BIT, &meta); + ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_PTR_CHAIN_BIT, + &meta, nullptr); if (!is_matrix(sub_type) && sub_type.basetype != SPIRType::Struct && expr_type.vecsize > sub_type.vecsize) expr += vector_swizzle(sub_type.vecsize, 0); @@ -8821,7 +8836,8 @@ bool CompilerMSL::emit_tessellation_io_load(uint32_t result_type_id, uint32_t id AccessChainMeta meta; expr += access_chain_internal( stage_in_ptr_var_id, indices, 2, - ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_PTR_CHAIN_BIT, &meta); + ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_PTR_CHAIN_BIT, + &meta, nullptr); } else expr += to_expression(ptr) + "." + to_member_name(iface_type, interface_index); @@ -8845,7 +8861,8 @@ bool CompilerMSL::emit_tessellation_io_load(uint32_t result_type_id, uint32_t id AccessChainMeta meta; expr += access_chain_internal( stage_in_ptr_var_id, indices, 2, - ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_PTR_CHAIN_BIT, &meta); + ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_PTR_CHAIN_BIT, + &meta, nullptr); } else expr += to_expression(ptr) + "." + to_member_name(iface_type, interface_index); @@ -8865,7 +8882,7 @@ bool CompilerMSL::emit_tessellation_io_load(uint32_t result_type_id, uint32_t id AccessChainMeta meta; expr += access_chain_internal(stage_in_ptr_var_id, indices, 2, ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_PTR_CHAIN_BIT, - &meta); + &meta, nullptr); } else expr += to_expression(ptr) + "." + to_member_name(iface_type, interface_index); @@ -8910,7 +8927,8 @@ bool CompilerMSL::emit_tessellation_io_load(uint32_t result_type_id, uint32_t id AccessChainMeta meta; expr += access_chain_internal(stage_in_ptr_var_id, indices, 2, - ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_PTR_CHAIN_BIT, &meta); + ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_PTR_CHAIN_BIT, + &meta, nullptr); if (expr_type.vecsize > result_type.vecsize) expr += vector_swizzle(result_type.vecsize, 0); if (j + 1 < result_type.columns) @@ -8955,7 +8973,8 @@ bool CompilerMSL::emit_tessellation_io_load(uint32_t result_type_id, uint32_t id AccessChainMeta meta; expr += access_chain_internal(stage_in_ptr_var_id, indices, 2, - ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_PTR_CHAIN_BIT, &meta); + ACCESS_CHAIN_INDEX_IS_LITERAL_BIT | ACCESS_CHAIN_PTR_CHAIN_BIT, + &meta, nullptr); if (expr_type.vecsize > result_type.vecsize) expr += vector_swizzle(result_type.vecsize, 0); @@ -9000,7 +9019,6 @@ bool CompilerMSL::emit_tessellation_access_chain(const uint32_t *ops, uint32_t l auto *var = maybe_get_backing_variable(ops[2]); bool patch = false; bool flat_data = false; - bool ptr_is_chain = false; bool flatten_composites = false; bool is_block = false; @@ -9021,12 +9039,6 @@ bool CompilerMSL::emit_tessellation_access_chain(const uint32_t *ops, uint32_t l // Patch inputs are treated as normal block IO variables, so they don't deal with this path at all. if (patch && (!is_block || is_arrayed || var->storage == StorageClassInput)) flat_data = false; - - // We might have a chained access chain, where - // we first take the access chain to the control point, and then we chain into a member or something similar. - // In this case, we need to skip gl_in/gl_out remapping. - // Also, skip ptr chain for patches. - ptr_is_chain = var->self != ID(ops[2]); } bool builtin_variable = false; @@ -9047,10 +9059,27 @@ bool CompilerMSL::emit_tessellation_access_chain(const uint32_t *ops, uint32_t l if (variable_is_flat) { + if (auto *ptr_expr = maybe_get(ops[2])) + { + // Too many edge cases in incrementally resolving tessellation access chains. + // Only reasonable option is to completely rematerialize the chain from the start. + SmallVector rematerialize_ops; + rematerialize_ops.push_back(ops[0]); + rematerialize_ops.push_back(ops[1]); + + for (auto expr : ptr_expr->implied_read_expressions) + rematerialize_ops.push_back(expr); + + for (uint32_t i = 3; i < length; i++) + rematerialize_ops.push_back(ops[i]); + + return emit_tessellation_access_chain(rematerialize_ops.data(), uint32_t(rematerialize_ops.size())); + } + // If output is masked, it is emitted as a "normal" variable, just go through normal code paths. // Only check this for the first level of access chain. // Dealing with this for partial access chains should be possible, but awkward. - if (var->storage == StorageClassOutput && !ptr_is_chain) + if (var->storage == StorageClassOutput) { bool masked = false; if (is_block) @@ -9077,7 +9106,7 @@ bool CompilerMSL::emit_tessellation_access_chain(const uint32_t *ops, uint32_t l indices.reserve(length - 3 + 1); - uint32_t first_non_array_index = (ptr_is_chain ? 3 : 4) - (patch ? 1 : 0); + uint32_t first_non_array_index = 4 - (patch ? 1 : 0); VariableID stage_var_id; if (patch) @@ -9085,8 +9114,9 @@ bool CompilerMSL::emit_tessellation_access_chain(const uint32_t *ops, uint32_t l else stage_var_id = var->storage == StorageClassInput ? stage_in_ptr_var_id : stage_out_ptr_var_id; - VariableID ptr = ptr_is_chain ? VariableID(ops[2]) : stage_var_id; - if (!ptr_is_chain && !patch) + VariableID ptr = stage_var_id; + + if (!patch) { // Index into gl_in/gl_out with first array index. indices.push_back(ops[first_non_array_index - 1]); @@ -9097,17 +9127,7 @@ bool CompilerMSL::emit_tessellation_access_chain(const uint32_t *ops, uint32_t l uint32_t const_mbr_id = next_id++; uint32_t index = get_extended_decoration(ops[2], SPIRVCrossDecorationInterfaceMemberIndex); - // If we have a pointer chain expression, and we are no longer pointing to a composite - // object, we are in the clear. There is no longer a need to flatten anything. - bool further_access_chain_is_trivial = false; - if (ptr_is_chain && flatten_composites) - { - auto &ptr_type = expression_type(ptr); - if (!is_array(ptr_type) && !is_matrix(ptr_type) && ptr_type.basetype != SPIRType::Struct) - further_access_chain_is_trivial = true; - } - - if (!further_access_chain_is_trivial && (flatten_composites || is_block)) + if (flatten_composites || is_block) { uint32_t i = first_non_array_index; auto *type = &get_variable_element_type(*var); @@ -9203,42 +9223,8 @@ bool CompilerMSL::emit_tessellation_access_chain(const uint32_t *ops, uint32_t l // We use the pointer to the base of the input/output array here, // so this is always a pointer chain. - string e; - - if (!ptr_is_chain) - { - // This is the start of an access chain, use ptr_chain to index into control point array. - e = access_chain(ptr, indices.data(), uint32_t(indices.size()), result_ptr_type, &meta, !patch); - } - else - { - // If we're accessing a struct, we need to use member indices which are based on the IO block, - // not actual struct type, so we have to use a split access chain here where - // first path resolves the control point index, i.e. gl_in[index], and second half deals with - // looking up flattened member name. - - // However, it is possible that we partially accessed a struct, - // by taking pointer to member inside the control-point array. - // For this case, we fall back to a natural access chain since we have already dealt with remapping struct members. - // One way to check this here is if we have 2 implied read expressions. - // First one is the gl_in/gl_out struct itself, then an index into that array. - // If we have traversed further, we use a normal access chain formulation. - auto *ptr_expr = maybe_get(ptr); - bool split_access_chain_formulation = flatten_composites && ptr_expr && - ptr_expr->implied_read_expressions.size() == 2 && - !further_access_chain_is_trivial; - - if (split_access_chain_formulation) - { - e = join(to_expression(ptr), - access_chain_internal(stage_var_id, indices.data(), uint32_t(indices.size()), - ACCESS_CHAIN_CHAIN_ONLY_BIT, &meta)); - } - else - { - e = access_chain_internal(ptr, indices.data(), uint32_t(indices.size()), 0, &meta); - } - } + // This is the start of an access chain, use ptr_chain to index into control point array. + auto e = access_chain(ptr, indices.data(), uint32_t(indices.size()), result_ptr_type, &meta, !patch); // Get the actual type of the object that was accessed. If it's a vector type and we changed it, // then we'll need to add a swizzle. @@ -9498,10 +9484,31 @@ void CompilerMSL::emit_instruction(const Instruction &instruction) } else { - // Sample mask input for Metal is not an array - if (BuiltIn(get_decoration(ptr, DecorationBuiltIn)) == BuiltInSampleMask) - set_decoration(id, DecorationBuiltIn, BuiltInSampleMask); - CompilerGLSL::emit_instruction(instruction); + auto is_sample_mask = BuiltIn(get_decoration(ptr, DecorationBuiltIn)) == BuiltInSampleMask; + auto ptr_storage = get_expression_effective_storage_class(ptr); + auto *ptr_var = maybe_get_backing_variable(ptr); + + // More edge cases ... Normally composite outputs are lowered at the end, + // but that's not the case for clip-cull arrays. + if (ptr_var && ptr_storage == StorageClassOutput && is_builtin_variable(*ptr_var) && + !is_sample_mask && is_array(get(ops[0]))) + { + emit_uninitialized_temporary_expression(ops[0], id); + auto &type = get(ops[0]); + if (type.array.size() != 1) + SPIRV_CROSS_THROW("Cannot load array of clip-cull distances from array of array."); + if (!type.array_size_literal.front()) + SPIRV_CROSS_THROW("Cannot load array of clip-cull distances from spec constant array size."); + for (uint32_t i = 0; i < type.array[0]; i++) + statement(to_expression(id), "[", i, "] = ", to_expression(ptr), "[", i, "];"); + } + else + { + // Sample mask input for Metal is not an array + if (is_sample_mask) + set_decoration(id, DecorationBuiltIn, BuiltInSampleMask); + CompilerGLSL::emit_instruction(instruction); + } } break; } @@ -11174,6 +11181,12 @@ bool CompilerMSL::emit_array_copy(const char *expr, uint32_t lhs_id, uint32_t rh else if (rhs_var && rhs_storage != StorageClassGeneric && type_is_explicit_layout(get(rhs_var->basetype))) rhs_is_array_template = false; + // Special consideration for clip/culldistance. Normally composites are lowered, but clip/cull is special for reasons ... + if (lhs_var && lhs_storage == StorageClassOutput && is_builtin_variable(*lhs_var)) + lhs_is_array_template = false; + if (rhs_var && rhs_storage == StorageClassOutput && is_builtin_variable(*rhs_var)) + rhs_is_array_template = false; + // If threadgroup storage qualifiers are *not* used: // Avoid spvCopy* wrapper functions; Otherwise, spvUnsafeArray<> template cannot be used with that storage qualifier. if (lhs_is_array_template && rhs_is_array_template && !using_builtin_array()) @@ -13821,9 +13834,6 @@ string CompilerMSL::member_attribute_qualifier(const SPIRType &type, uint32_t in return ""; return string(" [[") + builtin_qualifier(builtin) + "]]"; - case BuiltInDrawIndex: - SPIRV_CROSS_THROW("DrawIndex is not supported in MSL."); - default: return ""; } @@ -14672,6 +14682,9 @@ bool CompilerMSL::is_direct_input_builtin(BuiltIn bi_type) /* fallthrough */ case BuiltInSubgroupLocalInvocationId: return !msl_options.emulate_subgroups; + case BuiltInDrawIndex: + // Emulated + return false; default: return true; } @@ -15970,6 +15983,12 @@ void CompilerMSL::fix_up_shader_inputs_outputs() to_expression(builtin_dispatch_base_id), ".y;"); }); break; + case BuiltInDrawIndex: + entry_func.fixup_hooks_in.push_back([=]() { + statement(builtin_type_decl(bi_type), " ", to_expression(var_id), " = *", + to_expression(draw_index_buffer_id), ";"); + }); + break; default: break; } @@ -18120,8 +18139,9 @@ string CompilerMSL::builtin_to_glsl(BuiltIn builtin, StorageClass storage) { SPIRV_CROSS_THROW("BaseInstance requires Metal 1.1 and Mac or Apple A9+ hardware."); } + case BuiltInDrawIndex: - SPIRV_CROSS_THROW("DrawIndex is not supported in MSL."); + return "gl_DrawID"; // When used in the entry function, output builtins are qualified with output struct name. // Test storage class as NOT Input, as output builtins might be part of generic type. @@ -18232,8 +18252,6 @@ string CompilerMSL::builtin_qualifier(BuiltIn builtin) return "instance_id"; case BuiltInBaseInstance: return "base_instance"; - case BuiltInDrawIndex: - SPIRV_CROSS_THROW("DrawIndex is not supported in MSL."); // Vertex function out case BuiltInClipDistance: @@ -18457,7 +18475,7 @@ string CompilerMSL::builtin_type_decl(BuiltIn builtin, uint32_t id) case BuiltInBaseInstance: return "uint"; case BuiltInDrawIndex: - SPIRV_CROSS_THROW("DrawIndex is not supported in MSL."); + return "uint"; // Vertex function out case BuiltInClipDistance: @@ -19671,6 +19689,7 @@ void CompilerMSL::cast_from_variable_load(uint32_t source_id, std::string &expr, case BuiltInSubgroupSize: case BuiltInSubgroupLocalInvocationId: case BuiltInViewIndex: + case BuiltInDrawIndex: case BuiltInVertexIndex: case BuiltInInstanceIndex: case BuiltInBaseInstance: diff --git a/3rdparty/spirv-cross/spirv_msl.hpp b/3rdparty/spirv-cross/spirv_msl.hpp index 1a28e0585..37eea5ea8 100644 --- a/3rdparty/spirv-cross/spirv_msl.hpp +++ b/3rdparty/spirv-cross/spirv_msl.hpp @@ -317,6 +317,7 @@ public: uint32_t shader_input_buffer_index = 22; uint32_t shader_index_buffer_index = 21; uint32_t shader_patch_input_buffer_index = 20; + uint32_t draw_id_buffer_index = 19; uint32_t shader_input_wg_index = 0; uint32_t device_index = 0; uint32_t enable_frag_output_mask = 0xffffffff; @@ -1180,6 +1181,7 @@ protected: uint32_t swizzle_buffer_id = 0; uint32_t buffer_size_buffer_id = 0; uint32_t view_mask_buffer_id = 0; + uint32_t draw_index_buffer_id = 0; uint32_t dynamic_offsets_buffer_id = 0; uint32_t uint_type_id = 0; uint32_t shared_uint_type_id = 0; diff --git a/3rdparty/spirv-cross/spirv_parser.cpp b/3rdparty/spirv-cross/spirv_parser.cpp index d2e7dfa5f..2cf232a5f 100644 --- a/3rdparty/spirv-cross/spirv_parser.cpp +++ b/3rdparty/spirv-cross/spirv_parser.cpp @@ -151,8 +151,29 @@ void Parser::parse() SPIRV_CROSS_THROW("Function was not terminated."); if (current_block) SPIRV_CROSS_THROW("Block was not terminated."); + + // Now that all definitions are bound to a kind, we can filter the library + // exports and populate the exported functions. + for (uint32_t id : ir.library_exports) + { + if (ir.ids[id].get_type() == TypeFunction) + ir.library_exported_functions.push_back(id); + } + if (ir.default_entry_point == 0) - SPIRV_CROSS_THROW("There is no entry point in the SPIR-V module."); + { + if (ir.library_exported_functions.empty()) + SPIRV_CROSS_THROW("There is no entry point in the SPIR-V module."); + + // No OpEntryPoint, but the module exports functions. Treat as a library + // module: designate the first exported function as the default entry + // point so analyses keyed on default_entry_point can run. + ir.is_library_module = true; + ir.default_entry_point = ir.library_exported_functions.front(); + auto &name = ir.get_name(ir.default_entry_point); + ir.entry_points.insert(std::make_pair(ir.default_entry_point, + SPIREntryPoint(ir.default_entry_point, ExecutionModelGLCompute, name))); + } } const uint32_t *Parser::stream(const Instruction &instr) const @@ -604,6 +625,19 @@ void Parser::parse(const Instruction &instruction) else ir.set_decoration(id, decoration); + // Track exported functions so we can compile library modules that have no OpEntryPoint. + // LinkageAttributes layout: literal-string (variable words) followed by LinkageType. + if (decoration == DecorationLinkageAttributes && length >= 4 && + static_cast(ops[length - 1]) == LinkageTypeExport) + { + ir.library_exports.push_back(id); + + // If OpName was stripped (e.g. by spirv-opt --strip-debug), fall back + // to the linkage name so the emitted function keeps its export name. + if (ir.get_name(id).empty()) + ir.set_name(id, extract_string(ir.spirv, instruction.offset + 2)); + } + break; } @@ -627,6 +661,12 @@ void Parser::parse(const Instruction &instruction) break; } + // MemberDecorateIdEXT only applies to OffsetIdEXT when descriptors are packed in structs. + // This is currently unsupported and will fail in compilation. + // Pass it through in case someone just needs reflection. + case OpMemberDecorateIdEXT: + break; + case OpMemberDecorateStringGOOGLE: { uint32_t id = ops[0]; @@ -859,6 +899,7 @@ void Parser::parse(const Instruction &instruction) break; } + case OpTypeUntypedPointerKHR: case OpTypePointer: { uint32_t id = ops[0]; @@ -866,7 +907,7 @@ void Parser::parse(const Instruction &instruction) // Very rarely, we might receive a FunctionPrototype here. // 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(ops[2]); + auto *base = op == OpTypePointer ? maybe_get(ops[2]) : nullptr; auto &ptrbase = set(id, op); if (base) @@ -885,7 +926,10 @@ void Parser::parse(const Instruction &instruction) if (base && base->forward_pointer) forward_pointer_fixups.push_back({ id, ops[2] }); - ptrbase.parent_type = ops[2]; + if (op == OpTypePointer) + ptrbase.parent_type = ops[2]; + else + ptrbase.basetype = SPIRType::Void; // Do NOT set ptrbase.self! break; @@ -1005,6 +1049,27 @@ void Parser::parse(const Instruction &instruction) break; } + case OpUntypedVariableKHR: + { + uint32_t type = ops[0]; + uint32_t id = ops[1]; + auto storage = static_cast(ops[2]); + uint32_t data_type = length >= 4 ? ops[3] : 0; + uint32_t initializer = length >= 5 ? ops[4] : 0; + + if (storage == StorageClassFunction) + { + if (!current_function) + SPIRV_CROSS_THROW("No function currently in scope"); + current_function->add_local_variable(id); + } + + auto &v = set(id, type, storage, initializer); + v.untyped = true; + v.untyped_alloca_type = data_type; + break; + } + // OpPhi // OpPhi is a fairly magical opcode. // It selects temporary variables based on which parent block we *came from*. @@ -1135,6 +1200,24 @@ void Parser::parse(const Instruction &instruction) break; } + case OpConstantSizeOfEXT: + { + uint32_t id = ops[1]; + uint32_t type = ops[0]; + auto &c = set(id, type); + c.size_of_type = ops[2]; + break; + } + + case OpTypeBufferEXT: + { + uint32_t type = ops[0]; + auto &t = set(type, OpTypeBufferEXT); + t.basetype = SPIRType::DescriptorHeapBuffer; + t.ext.descriptor_heap_buffer.storage = static_cast(ops[1]); + break; + } + // Functions case OpFunction: {