From 113aadddf99377d7205ef840b3292bbf55648d97 Mon Sep 17 00:00:00 2001 From: Lei Zhang Date: Fri, 13 Sep 2019 15:25:56 -0700 Subject: [PATCH] Update SPIR-V symbols and use GLSL450 instead of VulkanKHR SPIR-V recently publishes v1.5, which brings a bunch of symbols into core. So the suffix "KHR"/"EXT"/etc. is removed from the symbols. We use a script to pull information from the spec directly. Also changed conversion and tests to use GLSL450 instead of VulkanKHR memory model. GLSL450 is still the main memory model supported by Vulkan shaders and it does not require extra capability to enable. PiperOrigin-RevId: 268992661 --- mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td | 407 +++++++++++---------- mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp | 2 +- mlir/test/Conversion/GPUToSPIRV/builtins.mlir | 12 +- mlir/test/Conversion/GPUToSPIRV/load_store.mlir | 2 +- mlir/test/Conversion/GPUToSPIRV/simple.mlir | 2 +- .../Dialect/SPIRV/Serialization/access_chain.mlir | 2 +- .../Dialect/SPIRV/Serialization/array_stride.mlir | 2 +- mlir/test/Dialect/SPIRV/Serialization/bin_ops.mlir | 2 +- mlir/test/Dialect/SPIRV/Serialization/entry.mlir | 2 +- .../SPIRV/Serialization/entry_interface.mlir | 2 +- .../SPIRV/Serialization/execution_mode.mlir | 2 +- .../Dialect/SPIRV/Serialization/load_store.mlir | 2 +- .../SPIRV/Serialization/minimal-module.mlir | 4 +- mlir/test/Dialect/SPIRV/Serialization/select.mlir | 2 +- mlir/test/Dialect/SPIRV/Serialization/struct.mlir | 2 +- .../SPIRV/Serialization/variable_reference.mlir | 2 +- .../Dialect/SPIRV/Serialization/variables.mlir | 2 +- .../SPIRV/Serialization/variables_init.mlir | 2 +- mlir/test/Dialect/SPIRV/control-flow-ops.mlir | 2 +- mlir/test/Dialect/SPIRV/ops.mlir | 10 +- mlir/test/Dialect/SPIRV/structure-ops.mlir | 60 +-- mlir/utils/spirv/gen_spirv_dialect.py | 2 +- 22 files changed, 265 insertions(+), 262 deletions(-) diff --git a/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td b/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td index 0accb05..ed9dee6 100644 --- a/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td +++ b/mlir/include/mlir/Dialect/SPIRV/SPIRVBase.td @@ -259,15 +259,15 @@ def SPV_ExtensionAttr : // Begin enum section. Generated from SPIR-V spec; DO NOT MODIFY! -def SPV_AM_Logical : I32EnumAttrCase<"Logical", 0>; -def SPV_AM_Physical32 : I32EnumAttrCase<"Physical32", 1>; -def SPV_AM_Physical64 : I32EnumAttrCase<"Physical64", 2>; -def SPV_AM_PhysicalStorageBuffer64EXT : I32EnumAttrCase<"PhysicalStorageBuffer64EXT", 5348>; +def SPV_AM_Logical : I32EnumAttrCase<"Logical", 0>; +def SPV_AM_Physical32 : I32EnumAttrCase<"Physical32", 1>; +def SPV_AM_Physical64 : I32EnumAttrCase<"Physical64", 2>; +def SPV_AM_PhysicalStorageBuffer64 : I32EnumAttrCase<"PhysicalStorageBuffer64", 5348>; def SPV_AddressingModelAttr : I32EnumAttr<"AddressingModel", "valid SPIR-V AddressingModel", [ SPV_AM_Logical, SPV_AM_Physical32, SPV_AM_Physical64, - SPV_AM_PhysicalStorageBuffer64EXT + SPV_AM_PhysicalStorageBuffer64 ]> { let returnType = "::mlir::spirv::AddressingModel"; let convertFromStorage = "static_cast<::mlir::spirv::AddressingModel>($_self.getInt())"; @@ -410,144 +410,146 @@ def SPV_BuiltInAttr : let cppNamespace = "::mlir::spirv"; } -def SPV_C_Matrix : I32EnumAttrCase<"Matrix", 0>; -def SPV_C_Shader : I32EnumAttrCase<"Shader", 1>; -def SPV_C_Geometry : I32EnumAttrCase<"Geometry", 2>; -def SPV_C_Tessellation : I32EnumAttrCase<"Tessellation", 3>; -def SPV_C_Addresses : I32EnumAttrCase<"Addresses", 4>; -def SPV_C_Linkage : I32EnumAttrCase<"Linkage", 5>; -def SPV_C_Kernel : I32EnumAttrCase<"Kernel", 6>; -def SPV_C_Vector16 : I32EnumAttrCase<"Vector16", 7>; -def SPV_C_Float16Buffer : I32EnumAttrCase<"Float16Buffer", 8>; -def SPV_C_Float16 : I32EnumAttrCase<"Float16", 9>; -def SPV_C_Float64 : I32EnumAttrCase<"Float64", 10>; -def SPV_C_Int64 : I32EnumAttrCase<"Int64", 11>; -def SPV_C_Int64Atomics : I32EnumAttrCase<"Int64Atomics", 12>; -def SPV_C_ImageBasic : I32EnumAttrCase<"ImageBasic", 13>; -def SPV_C_ImageReadWrite : I32EnumAttrCase<"ImageReadWrite", 14>; -def SPV_C_ImageMipmap : I32EnumAttrCase<"ImageMipmap", 15>; -def SPV_C_Pipes : I32EnumAttrCase<"Pipes", 17>; -def SPV_C_Groups : I32EnumAttrCase<"Groups", 18>; -def SPV_C_DeviceEnqueue : I32EnumAttrCase<"DeviceEnqueue", 19>; -def SPV_C_LiteralSampler : I32EnumAttrCase<"LiteralSampler", 20>; -def SPV_C_AtomicStorage : I32EnumAttrCase<"AtomicStorage", 21>; -def SPV_C_Int16 : I32EnumAttrCase<"Int16", 22>; -def SPV_C_TessellationPointSize : I32EnumAttrCase<"TessellationPointSize", 23>; -def SPV_C_GeometryPointSize : I32EnumAttrCase<"GeometryPointSize", 24>; -def SPV_C_ImageGatherExtended : I32EnumAttrCase<"ImageGatherExtended", 25>; -def SPV_C_StorageImageMultisample : I32EnumAttrCase<"StorageImageMultisample", 27>; -def SPV_C_UniformBufferArrayDynamicIndexing : I32EnumAttrCase<"UniformBufferArrayDynamicIndexing", 28>; -def SPV_C_SampledImageArrayDynamicIndexing : I32EnumAttrCase<"SampledImageArrayDynamicIndexing", 29>; -def SPV_C_StorageBufferArrayDynamicIndexing : I32EnumAttrCase<"StorageBufferArrayDynamicIndexing", 30>; -def SPV_C_StorageImageArrayDynamicIndexing : I32EnumAttrCase<"StorageImageArrayDynamicIndexing", 31>; -def SPV_C_ClipDistance : I32EnumAttrCase<"ClipDistance", 32>; -def SPV_C_CullDistance : I32EnumAttrCase<"CullDistance", 33>; -def SPV_C_ImageCubeArray : I32EnumAttrCase<"ImageCubeArray", 34>; -def SPV_C_SampleRateShading : I32EnumAttrCase<"SampleRateShading", 35>; -def SPV_C_ImageRect : I32EnumAttrCase<"ImageRect", 36>; -def SPV_C_SampledRect : I32EnumAttrCase<"SampledRect", 37>; -def SPV_C_GenericPointer : I32EnumAttrCase<"GenericPointer", 38>; -def SPV_C_Int8 : I32EnumAttrCase<"Int8", 39>; -def SPV_C_InputAttachment : I32EnumAttrCase<"InputAttachment", 40>; -def SPV_C_SparseResidency : I32EnumAttrCase<"SparseResidency", 41>; -def SPV_C_MinLod : I32EnumAttrCase<"MinLod", 42>; -def SPV_C_Sampled1D : I32EnumAttrCase<"Sampled1D", 43>; -def SPV_C_Image1D : I32EnumAttrCase<"Image1D", 44>; -def SPV_C_SampledCubeArray : I32EnumAttrCase<"SampledCubeArray", 45>; -def SPV_C_SampledBuffer : I32EnumAttrCase<"SampledBuffer", 46>; -def SPV_C_ImageBuffer : I32EnumAttrCase<"ImageBuffer", 47>; -def SPV_C_ImageMSArray : I32EnumAttrCase<"ImageMSArray", 48>; -def SPV_C_StorageImageExtendedFormats : I32EnumAttrCase<"StorageImageExtendedFormats", 49>; -def SPV_C_ImageQuery : I32EnumAttrCase<"ImageQuery", 50>; -def SPV_C_DerivativeControl : I32EnumAttrCase<"DerivativeControl", 51>; -def SPV_C_InterpolationFunction : I32EnumAttrCase<"InterpolationFunction", 52>; -def SPV_C_TransformFeedback : I32EnumAttrCase<"TransformFeedback", 53>; -def SPV_C_GeometryStreams : I32EnumAttrCase<"GeometryStreams", 54>; -def SPV_C_StorageImageReadWithoutFormat : I32EnumAttrCase<"StorageImageReadWithoutFormat", 55>; -def SPV_C_StorageImageWriteWithoutFormat : I32EnumAttrCase<"StorageImageWriteWithoutFormat", 56>; -def SPV_C_MultiViewport : I32EnumAttrCase<"MultiViewport", 57>; -def SPV_C_SubgroupDispatch : I32EnumAttrCase<"SubgroupDispatch", 58>; -def SPV_C_NamedBarrier : I32EnumAttrCase<"NamedBarrier", 59>; -def SPV_C_PipeStorage : I32EnumAttrCase<"PipeStorage", 60>; -def SPV_C_GroupNonUniform : I32EnumAttrCase<"GroupNonUniform", 61>; -def SPV_C_GroupNonUniformVote : I32EnumAttrCase<"GroupNonUniformVote", 62>; -def SPV_C_GroupNonUniformArithmetic : I32EnumAttrCase<"GroupNonUniformArithmetic", 63>; -def SPV_C_GroupNonUniformBallot : I32EnumAttrCase<"GroupNonUniformBallot", 64>; -def SPV_C_GroupNonUniformShuffle : I32EnumAttrCase<"GroupNonUniformShuffle", 65>; -def SPV_C_GroupNonUniformShuffleRelative : I32EnumAttrCase<"GroupNonUniformShuffleRelative", 66>; -def SPV_C_GroupNonUniformClustered : I32EnumAttrCase<"GroupNonUniformClustered", 67>; -def SPV_C_GroupNonUniformQuad : I32EnumAttrCase<"GroupNonUniformQuad", 68>; -def SPV_C_SubgroupBallotKHR : I32EnumAttrCase<"SubgroupBallotKHR", 4423>; -def SPV_C_DrawParameters : I32EnumAttrCase<"DrawParameters", 4427>; -def SPV_C_SubgroupVoteKHR : I32EnumAttrCase<"SubgroupVoteKHR", 4431>; -def SPV_C_StorageBuffer16BitAccess : I32EnumAttrCase<"StorageBuffer16BitAccess", 4433>; -def SPV_C_UniformAndStorageBuffer16BitAccess : I32EnumAttrCase<"UniformAndStorageBuffer16BitAccess", 4434>; -def SPV_C_StoragePushConstant16 : I32EnumAttrCase<"StoragePushConstant16", 4435>; -def SPV_C_StorageInputOutput16 : I32EnumAttrCase<"StorageInputOutput16", 4436>; -def SPV_C_DeviceGroup : I32EnumAttrCase<"DeviceGroup", 4437>; -def SPV_C_MultiView : I32EnumAttrCase<"MultiView", 4439>; -def SPV_C_VariablePointersStorageBuffer : I32EnumAttrCase<"VariablePointersStorageBuffer", 4441>; -def SPV_C_VariablePointers : I32EnumAttrCase<"VariablePointers", 4442>; -def SPV_C_AtomicStorageOps : I32EnumAttrCase<"AtomicStorageOps", 4445>; -def SPV_C_SampleMaskPostDepthCoverage : I32EnumAttrCase<"SampleMaskPostDepthCoverage", 4447>; -def SPV_C_StorageBuffer8BitAccess : I32EnumAttrCase<"StorageBuffer8BitAccess", 4448>; -def SPV_C_UniformAndStorageBuffer8BitAccess : I32EnumAttrCase<"UniformAndStorageBuffer8BitAccess", 4449>; -def SPV_C_StoragePushConstant8 : I32EnumAttrCase<"StoragePushConstant8", 4450>; -def SPV_C_DenormPreserve : I32EnumAttrCase<"DenormPreserve", 4464>; -def SPV_C_DenormFlushToZero : I32EnumAttrCase<"DenormFlushToZero", 4465>; -def SPV_C_SignedZeroInfNanPreserve : I32EnumAttrCase<"SignedZeroInfNanPreserve", 4466>; -def SPV_C_RoundingModeRTE : I32EnumAttrCase<"RoundingModeRTE", 4467>; -def SPV_C_RoundingModeRTZ : I32EnumAttrCase<"RoundingModeRTZ", 4468>; -def SPV_C_Float16ImageAMD : I32EnumAttrCase<"Float16ImageAMD", 5008>; -def SPV_C_ImageGatherBiasLodAMD : I32EnumAttrCase<"ImageGatherBiasLodAMD", 5009>; -def SPV_C_FragmentMaskAMD : I32EnumAttrCase<"FragmentMaskAMD", 5010>; -def SPV_C_StencilExportEXT : I32EnumAttrCase<"StencilExportEXT", 5013>; -def SPV_C_ImageReadWriteLodAMD : I32EnumAttrCase<"ImageReadWriteLodAMD", 5015>; -def SPV_C_ShaderClockKHR : I32EnumAttrCase<"ShaderClockKHR", 5055>; -def SPV_C_SampleMaskOverrideCoverageNV : I32EnumAttrCase<"SampleMaskOverrideCoverageNV", 5249>; -def SPV_C_GeometryShaderPassthroughNV : I32EnumAttrCase<"GeometryShaderPassthroughNV", 5251>; -def SPV_C_ShaderViewportIndexLayerEXT : I32EnumAttrCase<"ShaderViewportIndexLayerEXT", 5254>; -def SPV_C_ShaderViewportMaskNV : I32EnumAttrCase<"ShaderViewportMaskNV", 5255>; -def SPV_C_ShaderStereoViewNV : I32EnumAttrCase<"ShaderStereoViewNV", 5259>; -def SPV_C_PerViewAttributesNV : I32EnumAttrCase<"PerViewAttributesNV", 5260>; -def SPV_C_FragmentFullyCoveredEXT : I32EnumAttrCase<"FragmentFullyCoveredEXT", 5265>; -def SPV_C_MeshShadingNV : I32EnumAttrCase<"MeshShadingNV", 5266>; -def SPV_C_ImageFootprintNV : I32EnumAttrCase<"ImageFootprintNV", 5282>; -def SPV_C_FragmentBarycentricNV : I32EnumAttrCase<"FragmentBarycentricNV", 5284>; -def SPV_C_ComputeDerivativeGroupQuadsNV : I32EnumAttrCase<"ComputeDerivativeGroupQuadsNV", 5288>; -def SPV_C_FragmentDensityEXT : I32EnumAttrCase<"FragmentDensityEXT", 5291>; -def SPV_C_GroupNonUniformPartitionedNV : I32EnumAttrCase<"GroupNonUniformPartitionedNV", 5297>; -def SPV_C_ShaderNonUniformEXT : I32EnumAttrCase<"ShaderNonUniformEXT", 5301>; -def SPV_C_RuntimeDescriptorArrayEXT : I32EnumAttrCase<"RuntimeDescriptorArrayEXT", 5302>; -def SPV_C_InputAttachmentArrayDynamicIndexingEXT : I32EnumAttrCase<"InputAttachmentArrayDynamicIndexingEXT", 5303>; -def SPV_C_UniformTexelBufferArrayDynamicIndexingEXT : I32EnumAttrCase<"UniformTexelBufferArrayDynamicIndexingEXT", 5304>; -def SPV_C_StorageTexelBufferArrayDynamicIndexingEXT : I32EnumAttrCase<"StorageTexelBufferArrayDynamicIndexingEXT", 5305>; -def SPV_C_UniformBufferArrayNonUniformIndexingEXT : I32EnumAttrCase<"UniformBufferArrayNonUniformIndexingEXT", 5306>; -def SPV_C_SampledImageArrayNonUniformIndexingEXT : I32EnumAttrCase<"SampledImageArrayNonUniformIndexingEXT", 5307>; -def SPV_C_StorageBufferArrayNonUniformIndexingEXT : I32EnumAttrCase<"StorageBufferArrayNonUniformIndexingEXT", 5308>; -def SPV_C_StorageImageArrayNonUniformIndexingEXT : I32EnumAttrCase<"StorageImageArrayNonUniformIndexingEXT", 5309>; -def SPV_C_InputAttachmentArrayNonUniformIndexingEXT : I32EnumAttrCase<"InputAttachmentArrayNonUniformIndexingEXT", 5310>; -def SPV_C_UniformTexelBufferArrayNonUniformIndexingEXT : I32EnumAttrCase<"UniformTexelBufferArrayNonUniformIndexingEXT", 5311>; -def SPV_C_StorageTexelBufferArrayNonUniformIndexingEXT : I32EnumAttrCase<"StorageTexelBufferArrayNonUniformIndexingEXT", 5312>; -def SPV_C_RayTracingNV : I32EnumAttrCase<"RayTracingNV", 5340>; -def SPV_C_VulkanMemoryModelKHR : I32EnumAttrCase<"VulkanMemoryModelKHR", 5345>; -def SPV_C_VulkanMemoryModelDeviceScopeKHR : I32EnumAttrCase<"VulkanMemoryModelDeviceScopeKHR", 5346>; -def SPV_C_PhysicalStorageBufferAddressesEXT : I32EnumAttrCase<"PhysicalStorageBufferAddressesEXT", 5347>; -def SPV_C_ComputeDerivativeGroupLinearNV : I32EnumAttrCase<"ComputeDerivativeGroupLinearNV", 5350>; -def SPV_C_CooperativeMatrixNV : I32EnumAttrCase<"CooperativeMatrixNV", 5357>; -def SPV_C_FragmentShaderSampleInterlockEXT : I32EnumAttrCase<"FragmentShaderSampleInterlockEXT", 5363>; -def SPV_C_FragmentShaderShadingRateInterlockEXT : I32EnumAttrCase<"FragmentShaderShadingRateInterlockEXT", 5372>; -def SPV_C_ShaderSMBuiltinsNV : I32EnumAttrCase<"ShaderSMBuiltinsNV", 5373>; -def SPV_C_FragmentShaderPixelInterlockEXT : I32EnumAttrCase<"FragmentShaderPixelInterlockEXT", 5378>; -def SPV_C_DemoteToHelperInvocationEXT : I32EnumAttrCase<"DemoteToHelperInvocationEXT", 5379>; -def SPV_C_SubgroupShuffleINTEL : I32EnumAttrCase<"SubgroupShuffleINTEL", 5568>; -def SPV_C_SubgroupBufferBlockIOINTEL : I32EnumAttrCase<"SubgroupBufferBlockIOINTEL", 5569>; -def SPV_C_SubgroupImageBlockIOINTEL : I32EnumAttrCase<"SubgroupImageBlockIOINTEL", 5570>; -def SPV_C_SubgroupImageMediaBlockIOINTEL : I32EnumAttrCase<"SubgroupImageMediaBlockIOINTEL", 5579>; -def SPV_C_IntegerFunctions2INTEL : I32EnumAttrCase<"IntegerFunctions2INTEL", 5584>; -def SPV_C_SubgroupAvcMotionEstimationINTEL : I32EnumAttrCase<"SubgroupAvcMotionEstimationINTEL", 5696>; -def SPV_C_SubgroupAvcMotionEstimationIntraINTEL : I32EnumAttrCase<"SubgroupAvcMotionEstimationIntraINTEL", 5697>; -def SPV_C_SubgroupAvcMotionEstimationChromaINTEL : I32EnumAttrCase<"SubgroupAvcMotionEstimationChromaINTEL", 5698>; +def SPV_C_Matrix : I32EnumAttrCase<"Matrix", 0>; +def SPV_C_Shader : I32EnumAttrCase<"Shader", 1>; +def SPV_C_Geometry : I32EnumAttrCase<"Geometry", 2>; +def SPV_C_Tessellation : I32EnumAttrCase<"Tessellation", 3>; +def SPV_C_Addresses : I32EnumAttrCase<"Addresses", 4>; +def SPV_C_Linkage : I32EnumAttrCase<"Linkage", 5>; +def SPV_C_Kernel : I32EnumAttrCase<"Kernel", 6>; +def SPV_C_Vector16 : I32EnumAttrCase<"Vector16", 7>; +def SPV_C_Float16Buffer : I32EnumAttrCase<"Float16Buffer", 8>; +def SPV_C_Float16 : I32EnumAttrCase<"Float16", 9>; +def SPV_C_Float64 : I32EnumAttrCase<"Float64", 10>; +def SPV_C_Int64 : I32EnumAttrCase<"Int64", 11>; +def SPV_C_Int64Atomics : I32EnumAttrCase<"Int64Atomics", 12>; +def SPV_C_ImageBasic : I32EnumAttrCase<"ImageBasic", 13>; +def SPV_C_ImageReadWrite : I32EnumAttrCase<"ImageReadWrite", 14>; +def SPV_C_ImageMipmap : I32EnumAttrCase<"ImageMipmap", 15>; +def SPV_C_Pipes : I32EnumAttrCase<"Pipes", 17>; +def SPV_C_Groups : I32EnumAttrCase<"Groups", 18>; +def SPV_C_DeviceEnqueue : I32EnumAttrCase<"DeviceEnqueue", 19>; +def SPV_C_LiteralSampler : I32EnumAttrCase<"LiteralSampler", 20>; +def SPV_C_AtomicStorage : I32EnumAttrCase<"AtomicStorage", 21>; +def SPV_C_Int16 : I32EnumAttrCase<"Int16", 22>; +def SPV_C_TessellationPointSize : I32EnumAttrCase<"TessellationPointSize", 23>; +def SPV_C_GeometryPointSize : I32EnumAttrCase<"GeometryPointSize", 24>; +def SPV_C_ImageGatherExtended : I32EnumAttrCase<"ImageGatherExtended", 25>; +def SPV_C_StorageImageMultisample : I32EnumAttrCase<"StorageImageMultisample", 27>; +def SPV_C_UniformBufferArrayDynamicIndexing : I32EnumAttrCase<"UniformBufferArrayDynamicIndexing", 28>; +def SPV_C_SampledImageArrayDynamicIndexing : I32EnumAttrCase<"SampledImageArrayDynamicIndexing", 29>; +def SPV_C_StorageBufferArrayDynamicIndexing : I32EnumAttrCase<"StorageBufferArrayDynamicIndexing", 30>; +def SPV_C_StorageImageArrayDynamicIndexing : I32EnumAttrCase<"StorageImageArrayDynamicIndexing", 31>; +def SPV_C_ClipDistance : I32EnumAttrCase<"ClipDistance", 32>; +def SPV_C_CullDistance : I32EnumAttrCase<"CullDistance", 33>; +def SPV_C_ImageCubeArray : I32EnumAttrCase<"ImageCubeArray", 34>; +def SPV_C_SampleRateShading : I32EnumAttrCase<"SampleRateShading", 35>; +def SPV_C_ImageRect : I32EnumAttrCase<"ImageRect", 36>; +def SPV_C_SampledRect : I32EnumAttrCase<"SampledRect", 37>; +def SPV_C_GenericPointer : I32EnumAttrCase<"GenericPointer", 38>; +def SPV_C_Int8 : I32EnumAttrCase<"Int8", 39>; +def SPV_C_InputAttachment : I32EnumAttrCase<"InputAttachment", 40>; +def SPV_C_SparseResidency : I32EnumAttrCase<"SparseResidency", 41>; +def SPV_C_MinLod : I32EnumAttrCase<"MinLod", 42>; +def SPV_C_Sampled1D : I32EnumAttrCase<"Sampled1D", 43>; +def SPV_C_Image1D : I32EnumAttrCase<"Image1D", 44>; +def SPV_C_SampledCubeArray : I32EnumAttrCase<"SampledCubeArray", 45>; +def SPV_C_SampledBuffer : I32EnumAttrCase<"SampledBuffer", 46>; +def SPV_C_ImageBuffer : I32EnumAttrCase<"ImageBuffer", 47>; +def SPV_C_ImageMSArray : I32EnumAttrCase<"ImageMSArray", 48>; +def SPV_C_StorageImageExtendedFormats : I32EnumAttrCase<"StorageImageExtendedFormats", 49>; +def SPV_C_ImageQuery : I32EnumAttrCase<"ImageQuery", 50>; +def SPV_C_DerivativeControl : I32EnumAttrCase<"DerivativeControl", 51>; +def SPV_C_InterpolationFunction : I32EnumAttrCase<"InterpolationFunction", 52>; +def SPV_C_TransformFeedback : I32EnumAttrCase<"TransformFeedback", 53>; +def SPV_C_GeometryStreams : I32EnumAttrCase<"GeometryStreams", 54>; +def SPV_C_StorageImageReadWithoutFormat : I32EnumAttrCase<"StorageImageReadWithoutFormat", 55>; +def SPV_C_StorageImageWriteWithoutFormat : I32EnumAttrCase<"StorageImageWriteWithoutFormat", 56>; +def SPV_C_MultiViewport : I32EnumAttrCase<"MultiViewport", 57>; +def SPV_C_SubgroupDispatch : I32EnumAttrCase<"SubgroupDispatch", 58>; +def SPV_C_NamedBarrier : I32EnumAttrCase<"NamedBarrier", 59>; +def SPV_C_PipeStorage : I32EnumAttrCase<"PipeStorage", 60>; +def SPV_C_GroupNonUniform : I32EnumAttrCase<"GroupNonUniform", 61>; +def SPV_C_GroupNonUniformVote : I32EnumAttrCase<"GroupNonUniformVote", 62>; +def SPV_C_GroupNonUniformArithmetic : I32EnumAttrCase<"GroupNonUniformArithmetic", 63>; +def SPV_C_GroupNonUniformBallot : I32EnumAttrCase<"GroupNonUniformBallot", 64>; +def SPV_C_GroupNonUniformShuffle : I32EnumAttrCase<"GroupNonUniformShuffle", 65>; +def SPV_C_GroupNonUniformShuffleRelative : I32EnumAttrCase<"GroupNonUniformShuffleRelative", 66>; +def SPV_C_GroupNonUniformClustered : I32EnumAttrCase<"GroupNonUniformClustered", 67>; +def SPV_C_GroupNonUniformQuad : I32EnumAttrCase<"GroupNonUniformQuad", 68>; +def SPV_C_ShaderLayer : I32EnumAttrCase<"ShaderLayer", 69>; +def SPV_C_ShaderViewportIndex : I32EnumAttrCase<"ShaderViewportIndex", 70>; +def SPV_C_SubgroupBallotKHR : I32EnumAttrCase<"SubgroupBallotKHR", 4423>; +def SPV_C_DrawParameters : I32EnumAttrCase<"DrawParameters", 4427>; +def SPV_C_SubgroupVoteKHR : I32EnumAttrCase<"SubgroupVoteKHR", 4431>; +def SPV_C_StorageBuffer16BitAccess : I32EnumAttrCase<"StorageBuffer16BitAccess", 4433>; +def SPV_C_UniformAndStorageBuffer16BitAccess : I32EnumAttrCase<"UniformAndStorageBuffer16BitAccess", 4434>; +def SPV_C_StoragePushConstant16 : I32EnumAttrCase<"StoragePushConstant16", 4435>; +def SPV_C_StorageInputOutput16 : I32EnumAttrCase<"StorageInputOutput16", 4436>; +def SPV_C_DeviceGroup : I32EnumAttrCase<"DeviceGroup", 4437>; +def SPV_C_MultiView : I32EnumAttrCase<"MultiView", 4439>; +def SPV_C_VariablePointersStorageBuffer : I32EnumAttrCase<"VariablePointersStorageBuffer", 4441>; +def SPV_C_VariablePointers : I32EnumAttrCase<"VariablePointers", 4442>; +def SPV_C_AtomicStorageOps : I32EnumAttrCase<"AtomicStorageOps", 4445>; +def SPV_C_SampleMaskPostDepthCoverage : I32EnumAttrCase<"SampleMaskPostDepthCoverage", 4447>; +def SPV_C_StorageBuffer8BitAccess : I32EnumAttrCase<"StorageBuffer8BitAccess", 4448>; +def SPV_C_UniformAndStorageBuffer8BitAccess : I32EnumAttrCase<"UniformAndStorageBuffer8BitAccess", 4449>; +def SPV_C_StoragePushConstant8 : I32EnumAttrCase<"StoragePushConstant8", 4450>; +def SPV_C_DenormPreserve : I32EnumAttrCase<"DenormPreserve", 4464>; +def SPV_C_DenormFlushToZero : I32EnumAttrCase<"DenormFlushToZero", 4465>; +def SPV_C_SignedZeroInfNanPreserve : I32EnumAttrCase<"SignedZeroInfNanPreserve", 4466>; +def SPV_C_RoundingModeRTE : I32EnumAttrCase<"RoundingModeRTE", 4467>; +def SPV_C_RoundingModeRTZ : I32EnumAttrCase<"RoundingModeRTZ", 4468>; +def SPV_C_Float16ImageAMD : I32EnumAttrCase<"Float16ImageAMD", 5008>; +def SPV_C_ImageGatherBiasLodAMD : I32EnumAttrCase<"ImageGatherBiasLodAMD", 5009>; +def SPV_C_FragmentMaskAMD : I32EnumAttrCase<"FragmentMaskAMD", 5010>; +def SPV_C_StencilExportEXT : I32EnumAttrCase<"StencilExportEXT", 5013>; +def SPV_C_ImageReadWriteLodAMD : I32EnumAttrCase<"ImageReadWriteLodAMD", 5015>; +def SPV_C_ShaderClockKHR : I32EnumAttrCase<"ShaderClockKHR", 5055>; +def SPV_C_SampleMaskOverrideCoverageNV : I32EnumAttrCase<"SampleMaskOverrideCoverageNV", 5249>; +def SPV_C_GeometryShaderPassthroughNV : I32EnumAttrCase<"GeometryShaderPassthroughNV", 5251>; +def SPV_C_ShaderViewportIndexLayerEXT : I32EnumAttrCase<"ShaderViewportIndexLayerEXT", 5254>; +def SPV_C_ShaderViewportMaskNV : I32EnumAttrCase<"ShaderViewportMaskNV", 5255>; +def SPV_C_ShaderStereoViewNV : I32EnumAttrCase<"ShaderStereoViewNV", 5259>; +def SPV_C_PerViewAttributesNV : I32EnumAttrCase<"PerViewAttributesNV", 5260>; +def SPV_C_FragmentFullyCoveredEXT : I32EnumAttrCase<"FragmentFullyCoveredEXT", 5265>; +def SPV_C_MeshShadingNV : I32EnumAttrCase<"MeshShadingNV", 5266>; +def SPV_C_ImageFootprintNV : I32EnumAttrCase<"ImageFootprintNV", 5282>; +def SPV_C_FragmentBarycentricNV : I32EnumAttrCase<"FragmentBarycentricNV", 5284>; +def SPV_C_ComputeDerivativeGroupQuadsNV : I32EnumAttrCase<"ComputeDerivativeGroupQuadsNV", 5288>; +def SPV_C_FragmentDensityEXT : I32EnumAttrCase<"FragmentDensityEXT", 5291>; +def SPV_C_GroupNonUniformPartitionedNV : I32EnumAttrCase<"GroupNonUniformPartitionedNV", 5297>; +def SPV_C_ShaderNonUniform : I32EnumAttrCase<"ShaderNonUniform", 5301>; +def SPV_C_RuntimeDescriptorArray : I32EnumAttrCase<"RuntimeDescriptorArray", 5302>; +def SPV_C_InputAttachmentArrayDynamicIndexing : I32EnumAttrCase<"InputAttachmentArrayDynamicIndexing", 5303>; +def SPV_C_UniformTexelBufferArrayDynamicIndexing : I32EnumAttrCase<"UniformTexelBufferArrayDynamicIndexing", 5304>; +def SPV_C_StorageTexelBufferArrayDynamicIndexing : I32EnumAttrCase<"StorageTexelBufferArrayDynamicIndexing", 5305>; +def SPV_C_UniformBufferArrayNonUniformIndexing : I32EnumAttrCase<"UniformBufferArrayNonUniformIndexing", 5306>; +def SPV_C_SampledImageArrayNonUniformIndexing : I32EnumAttrCase<"SampledImageArrayNonUniformIndexing", 5307>; +def SPV_C_StorageBufferArrayNonUniformIndexing : I32EnumAttrCase<"StorageBufferArrayNonUniformIndexing", 5308>; +def SPV_C_StorageImageArrayNonUniformIndexing : I32EnumAttrCase<"StorageImageArrayNonUniformIndexing", 5309>; +def SPV_C_InputAttachmentArrayNonUniformIndexing : I32EnumAttrCase<"InputAttachmentArrayNonUniformIndexing", 5310>; +def SPV_C_UniformTexelBufferArrayNonUniformIndexing : I32EnumAttrCase<"UniformTexelBufferArrayNonUniformIndexing", 5311>; +def SPV_C_StorageTexelBufferArrayNonUniformIndexing : I32EnumAttrCase<"StorageTexelBufferArrayNonUniformIndexing", 5312>; +def SPV_C_RayTracingNV : I32EnumAttrCase<"RayTracingNV", 5340>; +def SPV_C_VulkanMemoryModel : I32EnumAttrCase<"VulkanMemoryModel", 5345>; +def SPV_C_VulkanMemoryModelDeviceScope : I32EnumAttrCase<"VulkanMemoryModelDeviceScope", 5346>; +def SPV_C_PhysicalStorageBufferAddresses : I32EnumAttrCase<"PhysicalStorageBufferAddresses", 5347>; +def SPV_C_ComputeDerivativeGroupLinearNV : I32EnumAttrCase<"ComputeDerivativeGroupLinearNV", 5350>; +def SPV_C_CooperativeMatrixNV : I32EnumAttrCase<"CooperativeMatrixNV", 5357>; +def SPV_C_FragmentShaderSampleInterlockEXT : I32EnumAttrCase<"FragmentShaderSampleInterlockEXT", 5363>; +def SPV_C_FragmentShaderShadingRateInterlockEXT : I32EnumAttrCase<"FragmentShaderShadingRateInterlockEXT", 5372>; +def SPV_C_ShaderSMBuiltinsNV : I32EnumAttrCase<"ShaderSMBuiltinsNV", 5373>; +def SPV_C_FragmentShaderPixelInterlockEXT : I32EnumAttrCase<"FragmentShaderPixelInterlockEXT", 5378>; +def SPV_C_DemoteToHelperInvocationEXT : I32EnumAttrCase<"DemoteToHelperInvocationEXT", 5379>; +def SPV_C_SubgroupShuffleINTEL : I32EnumAttrCase<"SubgroupShuffleINTEL", 5568>; +def SPV_C_SubgroupBufferBlockIOINTEL : I32EnumAttrCase<"SubgroupBufferBlockIOINTEL", 5569>; +def SPV_C_SubgroupImageBlockIOINTEL : I32EnumAttrCase<"SubgroupImageBlockIOINTEL", 5570>; +def SPV_C_SubgroupImageMediaBlockIOINTEL : I32EnumAttrCase<"SubgroupImageMediaBlockIOINTEL", 5579>; +def SPV_C_IntegerFunctions2INTEL : I32EnumAttrCase<"IntegerFunctions2INTEL", 5584>; +def SPV_C_SubgroupAvcMotionEstimationINTEL : I32EnumAttrCase<"SubgroupAvcMotionEstimationINTEL", 5696>; +def SPV_C_SubgroupAvcMotionEstimationIntraINTEL : I32EnumAttrCase<"SubgroupAvcMotionEstimationIntraINTEL", 5697>; +def SPV_C_SubgroupAvcMotionEstimationChromaINTEL : I32EnumAttrCase<"SubgroupAvcMotionEstimationChromaINTEL", 5698>; def SPV_CapabilityAttr : I32EnumAttr<"Capability", "valid SPIR-V Capability", [ @@ -574,11 +576,12 @@ def SPV_CapabilityAttr : SPV_C_GroupNonUniform, SPV_C_GroupNonUniformVote, SPV_C_GroupNonUniformArithmetic, SPV_C_GroupNonUniformBallot, SPV_C_GroupNonUniformShuffle, SPV_C_GroupNonUniformShuffleRelative, - SPV_C_GroupNonUniformClustered, SPV_C_GroupNonUniformQuad, - SPV_C_SubgroupBallotKHR, SPV_C_DrawParameters, SPV_C_SubgroupVoteKHR, - SPV_C_StorageBuffer16BitAccess, SPV_C_UniformAndStorageBuffer16BitAccess, - SPV_C_StoragePushConstant16, SPV_C_StorageInputOutput16, SPV_C_DeviceGroup, - SPV_C_MultiView, SPV_C_VariablePointersStorageBuffer, SPV_C_VariablePointers, + SPV_C_GroupNonUniformClustered, SPV_C_GroupNonUniformQuad, SPV_C_ShaderLayer, + SPV_C_ShaderViewportIndex, SPV_C_SubgroupBallotKHR, SPV_C_DrawParameters, + SPV_C_SubgroupVoteKHR, SPV_C_StorageBuffer16BitAccess, + SPV_C_UniformAndStorageBuffer16BitAccess, SPV_C_StoragePushConstant16, + SPV_C_StorageInputOutput16, SPV_C_DeviceGroup, SPV_C_MultiView, + SPV_C_VariablePointersStorageBuffer, SPV_C_VariablePointers, SPV_C_AtomicStorageOps, SPV_C_SampleMaskPostDepthCoverage, SPV_C_StorageBuffer8BitAccess, SPV_C_UniformAndStorageBuffer8BitAccess, SPV_C_StoragePushConstant8, SPV_C_DenormPreserve, SPV_C_DenormFlushToZero, @@ -591,19 +594,19 @@ def SPV_CapabilityAttr : SPV_C_FragmentFullyCoveredEXT, SPV_C_MeshShadingNV, SPV_C_ImageFootprintNV, SPV_C_FragmentBarycentricNV, SPV_C_ComputeDerivativeGroupQuadsNV, SPV_C_FragmentDensityEXT, SPV_C_GroupNonUniformPartitionedNV, - SPV_C_ShaderNonUniformEXT, SPV_C_RuntimeDescriptorArrayEXT, - SPV_C_InputAttachmentArrayDynamicIndexingEXT, - SPV_C_UniformTexelBufferArrayDynamicIndexingEXT, - SPV_C_StorageTexelBufferArrayDynamicIndexingEXT, - SPV_C_UniformBufferArrayNonUniformIndexingEXT, - SPV_C_SampledImageArrayNonUniformIndexingEXT, - SPV_C_StorageBufferArrayNonUniformIndexingEXT, - SPV_C_StorageImageArrayNonUniformIndexingEXT, - SPV_C_InputAttachmentArrayNonUniformIndexingEXT, - SPV_C_UniformTexelBufferArrayNonUniformIndexingEXT, - SPV_C_StorageTexelBufferArrayNonUniformIndexingEXT, SPV_C_RayTracingNV, - SPV_C_VulkanMemoryModelKHR, SPV_C_VulkanMemoryModelDeviceScopeKHR, - SPV_C_PhysicalStorageBufferAddressesEXT, SPV_C_ComputeDerivativeGroupLinearNV, + SPV_C_ShaderNonUniform, SPV_C_RuntimeDescriptorArray, + SPV_C_InputAttachmentArrayDynamicIndexing, + SPV_C_UniformTexelBufferArrayDynamicIndexing, + SPV_C_StorageTexelBufferArrayDynamicIndexing, + SPV_C_UniformBufferArrayNonUniformIndexing, + SPV_C_SampledImageArrayNonUniformIndexing, + SPV_C_StorageBufferArrayNonUniformIndexing, + SPV_C_StorageImageArrayNonUniformIndexing, + SPV_C_InputAttachmentArrayNonUniformIndexing, + SPV_C_UniformTexelBufferArrayNonUniformIndexing, + SPV_C_StorageTexelBufferArrayNonUniformIndexing, SPV_C_RayTracingNV, + SPV_C_VulkanMemoryModel, SPV_C_VulkanMemoryModelDeviceScope, + SPV_C_PhysicalStorageBufferAddresses, SPV_C_ComputeDerivativeGroupLinearNV, SPV_C_CooperativeMatrixNV, SPV_C_FragmentShaderSampleInterlockEXT, SPV_C_FragmentShaderShadingRateInterlockEXT, SPV_C_ShaderSMBuiltinsNV, SPV_C_FragmentShaderPixelInterlockEXT, SPV_C_DemoteToHelperInvocationEXT, @@ -676,9 +679,9 @@ def SPV_D_PerPrimitiveNV : I32EnumAttrCase<"PerPrimitiveNV", 5271>; def SPV_D_PerViewNV : I32EnumAttrCase<"PerViewNV", 5272>; def SPV_D_PerTaskNV : I32EnumAttrCase<"PerTaskNV", 5273>; def SPV_D_PerVertexNV : I32EnumAttrCase<"PerVertexNV", 5285>; -def SPV_D_NonUniformEXT : I32EnumAttrCase<"NonUniformEXT", 5300>; -def SPV_D_RestrictPointerEXT : I32EnumAttrCase<"RestrictPointerEXT", 5355>; -def SPV_D_AliasedPointerEXT : I32EnumAttrCase<"AliasedPointerEXT", 5356>; +def SPV_D_NonUniform : I32EnumAttrCase<"NonUniform", 5300>; +def SPV_D_RestrictPointer : I32EnumAttrCase<"RestrictPointer", 5355>; +def SPV_D_AliasedPointer : I32EnumAttrCase<"AliasedPointer", 5356>; def SPV_D_CounterBuffer : I32EnumAttrCase<"CounterBuffer", 5634>; def SPV_D_UserSemantic : I32EnumAttrCase<"UserSemantic", 5635>; def SPV_D_UserTypeGOOGLE : I32EnumAttrCase<"UserTypeGOOGLE", 5636>; @@ -700,9 +703,9 @@ def SPV_DecorationAttr : SPV_D_NoUnsignedWrap, SPV_D_ExplicitInterpAMD, SPV_D_OverrideCoverageNV, SPV_D_PassthroughNV, SPV_D_ViewportRelativeNV, SPV_D_SecondaryViewportRelativeNV, SPV_D_PerPrimitiveNV, SPV_D_PerViewNV, - SPV_D_PerTaskNV, SPV_D_PerVertexNV, SPV_D_NonUniformEXT, - SPV_D_RestrictPointerEXT, SPV_D_AliasedPointerEXT, SPV_D_CounterBuffer, - SPV_D_UserSemantic, SPV_D_UserTypeGOOGLE + SPV_D_PerTaskNV, SPV_D_PerVertexNV, SPV_D_NonUniform, SPV_D_RestrictPointer, + SPV_D_AliasedPointer, SPV_D_CounterBuffer, SPV_D_UserSemantic, + SPV_D_UserTypeGOOGLE ]> { let returnType = "::mlir::spirv::Decoration"; let convertFromStorage = "static_cast<::mlir::spirv::Decoration>($_self.getInt())"; @@ -947,59 +950,59 @@ def SPV_LoopControlAttr : let cppNamespace = "::mlir::spirv"; } -def SPV_MA_None : I32EnumAttrCase<"None", 0x0000>; -def SPV_MA_Volatile : I32EnumAttrCase<"Volatile", 0x0001>; -def SPV_MA_Aligned : I32EnumAttrCase<"Aligned", 0x0002>; -def SPV_MA_Nontemporal : I32EnumAttrCase<"Nontemporal", 0x0004>; -def SPV_MA_MakePointerAvailableKHR : I32EnumAttrCase<"MakePointerAvailableKHR", 0x0008>; -def SPV_MA_MakePointerVisibleKHR : I32EnumAttrCase<"MakePointerVisibleKHR", 0x0010>; -def SPV_MA_NonPrivatePointerKHR : I32EnumAttrCase<"NonPrivatePointerKHR", 0x0020>; +def SPV_MA_None : I32EnumAttrCase<"None", 0x0000>; +def SPV_MA_Volatile : I32EnumAttrCase<"Volatile", 0x0001>; +def SPV_MA_Aligned : I32EnumAttrCase<"Aligned", 0x0002>; +def SPV_MA_Nontemporal : I32EnumAttrCase<"Nontemporal", 0x0004>; +def SPV_MA_MakePointerAvailable : I32EnumAttrCase<"MakePointerAvailable", 0x0008>; +def SPV_MA_MakePointerVisible : I32EnumAttrCase<"MakePointerVisible", 0x0010>; +def SPV_MA_NonPrivatePointer : I32EnumAttrCase<"NonPrivatePointer", 0x0020>; def SPV_MemoryAccessAttr : I32EnumAttr<"MemoryAccess", "valid SPIR-V MemoryAccess", [ SPV_MA_None, SPV_MA_Volatile, SPV_MA_Aligned, SPV_MA_Nontemporal, - SPV_MA_MakePointerAvailableKHR, SPV_MA_MakePointerVisibleKHR, - SPV_MA_NonPrivatePointerKHR + SPV_MA_MakePointerAvailable, SPV_MA_MakePointerVisible, + SPV_MA_NonPrivatePointer ]> { let returnType = "::mlir::spirv::MemoryAccess"; let convertFromStorage = "static_cast<::mlir::spirv::MemoryAccess>($_self.getInt())"; let cppNamespace = "::mlir::spirv"; } -def SPV_MM_Simple : I32EnumAttrCase<"Simple", 0>; -def SPV_MM_GLSL450 : I32EnumAttrCase<"GLSL450", 1>; -def SPV_MM_OpenCL : I32EnumAttrCase<"OpenCL", 2>; -def SPV_MM_VulkanKHR : I32EnumAttrCase<"VulkanKHR", 3>; +def SPV_MM_Simple : I32EnumAttrCase<"Simple", 0>; +def SPV_MM_GLSL450 : I32EnumAttrCase<"GLSL450", 1>; +def SPV_MM_OpenCL : I32EnumAttrCase<"OpenCL", 2>; +def SPV_MM_Vulkan : I32EnumAttrCase<"Vulkan", 3>; def SPV_MemoryModelAttr : I32EnumAttr<"MemoryModel", "valid SPIR-V MemoryModel", [ - SPV_MM_Simple, SPV_MM_GLSL450, SPV_MM_OpenCL, SPV_MM_VulkanKHR + SPV_MM_Simple, SPV_MM_GLSL450, SPV_MM_OpenCL, SPV_MM_Vulkan ]> { let returnType = "::mlir::spirv::MemoryModel"; let convertFromStorage = "static_cast<::mlir::spirv::MemoryModel>($_self.getInt())"; let cppNamespace = "::mlir::spirv"; } -def SPV_SC_UniformConstant : I32EnumAttrCase<"UniformConstant", 0>; -def SPV_SC_Input : I32EnumAttrCase<"Input", 1>; -def SPV_SC_Uniform : I32EnumAttrCase<"Uniform", 2>; -def SPV_SC_Output : I32EnumAttrCase<"Output", 3>; -def SPV_SC_Workgroup : I32EnumAttrCase<"Workgroup", 4>; -def SPV_SC_CrossWorkgroup : I32EnumAttrCase<"CrossWorkgroup", 5>; -def SPV_SC_Private : I32EnumAttrCase<"Private", 6>; -def SPV_SC_Function : I32EnumAttrCase<"Function", 7>; -def SPV_SC_Generic : I32EnumAttrCase<"Generic", 8>; -def SPV_SC_PushConstant : I32EnumAttrCase<"PushConstant", 9>; -def SPV_SC_AtomicCounter : I32EnumAttrCase<"AtomicCounter", 10>; -def SPV_SC_Image : I32EnumAttrCase<"Image", 11>; -def SPV_SC_StorageBuffer : I32EnumAttrCase<"StorageBuffer", 12>; -def SPV_SC_CallableDataNV : I32EnumAttrCase<"CallableDataNV", 5328>; -def SPV_SC_IncomingCallableDataNV : I32EnumAttrCase<"IncomingCallableDataNV", 5329>; -def SPV_SC_RayPayloadNV : I32EnumAttrCase<"RayPayloadNV", 5338>; -def SPV_SC_HitAttributeNV : I32EnumAttrCase<"HitAttributeNV", 5339>; -def SPV_SC_IncomingRayPayloadNV : I32EnumAttrCase<"IncomingRayPayloadNV", 5342>; -def SPV_SC_ShaderRecordBufferNV : I32EnumAttrCase<"ShaderRecordBufferNV", 5343>; -def SPV_SC_PhysicalStorageBufferEXT : I32EnumAttrCase<"PhysicalStorageBufferEXT", 5349>; +def SPV_SC_UniformConstant : I32EnumAttrCase<"UniformConstant", 0>; +def SPV_SC_Input : I32EnumAttrCase<"Input", 1>; +def SPV_SC_Uniform : I32EnumAttrCase<"Uniform", 2>; +def SPV_SC_Output : I32EnumAttrCase<"Output", 3>; +def SPV_SC_Workgroup : I32EnumAttrCase<"Workgroup", 4>; +def SPV_SC_CrossWorkgroup : I32EnumAttrCase<"CrossWorkgroup", 5>; +def SPV_SC_Private : I32EnumAttrCase<"Private", 6>; +def SPV_SC_Function : I32EnumAttrCase<"Function", 7>; +def SPV_SC_Generic : I32EnumAttrCase<"Generic", 8>; +def SPV_SC_PushConstant : I32EnumAttrCase<"PushConstant", 9>; +def SPV_SC_AtomicCounter : I32EnumAttrCase<"AtomicCounter", 10>; +def SPV_SC_Image : I32EnumAttrCase<"Image", 11>; +def SPV_SC_StorageBuffer : I32EnumAttrCase<"StorageBuffer", 12>; +def SPV_SC_CallableDataNV : I32EnumAttrCase<"CallableDataNV", 5328>; +def SPV_SC_IncomingCallableDataNV : I32EnumAttrCase<"IncomingCallableDataNV", 5329>; +def SPV_SC_RayPayloadNV : I32EnumAttrCase<"RayPayloadNV", 5338>; +def SPV_SC_HitAttributeNV : I32EnumAttrCase<"HitAttributeNV", 5339>; +def SPV_SC_IncomingRayPayloadNV : I32EnumAttrCase<"IncomingRayPayloadNV", 5342>; +def SPV_SC_ShaderRecordBufferNV : I32EnumAttrCase<"ShaderRecordBufferNV", 5343>; +def SPV_SC_PhysicalStorageBuffer : I32EnumAttrCase<"PhysicalStorageBuffer", 5349>; def SPV_StorageClassAttr : I32EnumAttr<"StorageClass", "valid SPIR-V StorageClass", [ @@ -1008,7 +1011,7 @@ def SPV_StorageClassAttr : SPV_SC_Generic, SPV_SC_PushConstant, SPV_SC_AtomicCounter, SPV_SC_Image, SPV_SC_StorageBuffer, SPV_SC_CallableDataNV, SPV_SC_IncomingCallableDataNV, SPV_SC_RayPayloadNV, SPV_SC_HitAttributeNV, SPV_SC_IncomingRayPayloadNV, - SPV_SC_ShaderRecordBufferNV, SPV_SC_PhysicalStorageBufferEXT + SPV_SC_ShaderRecordBufferNV, SPV_SC_PhysicalStorageBuffer ]> { let returnType = "::mlir::spirv::StorageClass"; let convertFromStorage = "static_cast<::mlir::spirv::StorageClass>($_self.getInt())"; diff --git a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp index 544232e..4760ed0 100644 --- a/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp +++ b/mlir/lib/Conversion/GPUToSPIRV/GPUToSPIRV.cpp @@ -134,7 +134,7 @@ void GPUToSPIRVPass::runOnModule() { builder.getI32IntegerAttr( static_cast(spirv::AddressingModel::Logical)), builder.getI32IntegerAttr( - static_cast(spirv::MemoryModel::VulkanKHR))); + static_cast(spirv::MemoryModel::GLSL450))); OpBuilder moduleBuilder(spvModule.getOperation()->getRegion(0)); moduleBuilder.clone(*funcOp.getOperation()); spirvModules.push_back(spvModule); diff --git a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir index ce9421e..73fb183 100644 --- a/mlir/test/Conversion/GPUToSPIRV/builtins.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/builtins.mlir @@ -6,7 +6,7 @@ func @builtin() { return } -// CHECK-LABEL: spv.module "Logical" "VulkanKHR" +// CHECK-LABEL: spv.module "Logical" "GLSL450" // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") func @builtin_workgroup_id_x() attributes {gpu.kernel} { @@ -25,7 +25,7 @@ func @builtin() { return } -// CHECK-LABEL: spv.module "Logical" "VulkanKHR" +// CHECK-LABEL: spv.module "Logical" "GLSL450" // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") func @builtin_workgroup_id_y() attributes {gpu.kernel} { @@ -44,7 +44,7 @@ func @builtin() { return } -// CHECK-LABEL: spv.module "Logical" "VulkanKHR" +// CHECK-LABEL: spv.module "Logical" "GLSL450" // CHECK: spv.globalVariable [[WORKGROUPID:@.*]] built_in("WorkgroupId") func @builtin_workgroup_id_z() attributes {gpu.kernel} { @@ -63,7 +63,7 @@ func @builtin() { return } -// CHECK-LABEL: spv.module "Logical" "VulkanKHR" +// CHECK-LABEL: spv.module "Logical" "GLSL450" // CHECK: spv.globalVariable [[WORKGROUPSIZE:@.*]] built_in("WorkgroupSize") func @builtin_workgroup_size_x() attributes {gpu.kernel} { @@ -82,7 +82,7 @@ func @builtin() { return } -// CHECK-LABEL: spv.module "Logical" "VulkanKHR" +// CHECK-LABEL: spv.module "Logical" "GLSL450" // CHECK: spv.globalVariable [[LOCALINVOCATIONID:@.*]] built_in("LocalInvocationId") func @builtin_local_id_x() attributes {gpu.kernel} { @@ -101,7 +101,7 @@ func @builtin() { return } -// CHECK-LABEL: spv.module "Logical" "VulkanKHR" +// CHECK-LABEL: spv.module "Logical" "GLSL450" // CHECK: spv.globalVariable [[NUMWORKGROUPS:@.*]] built_in("NumWorkgroups") func @builtin_num_workgroups_x() attributes {gpu.kernel} { diff --git a/mlir/test/Conversion/GPUToSPIRV/load_store.mlir b/mlir/test/Conversion/GPUToSPIRV/load_store.mlir index cc8ed07..e86cc19 100644 --- a/mlir/test/Conversion/GPUToSPIRV/load_store.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/load_store.mlir @@ -14,7 +14,7 @@ func @load_store(%arg0: memref<12x4xf32>, %arg1: memref<12x4xf32>, %arg2: memref return } -// CHECK-LABEL: spv.module "Logical" "VulkanKHR" +// CHECK-LABEL: spv.module "Logical" "GLSL450" // CHECK: spv.globalVariable {{@.*}} bind(0, 0) : [[TYPE1:!spv.ptr>, StorageBuffer>]] // CHECK-NEXT: spv.globalVariable {{@.*}} bind(0, 1) : [[TYPE2:!spv.ptr>, StorageBuffer>]] // CHECK-NEXT: spv.globalVariable {{@.*}} bind(0, 2) : [[TYPE3:!spv.ptr>, StorageBuffer>]] diff --git a/mlir/test/Conversion/GPUToSPIRV/simple.mlir b/mlir/test/Conversion/GPUToSPIRV/simple.mlir index ca522a3..a92ec96 100644 --- a/mlir/test/Conversion/GPUToSPIRV/simple.mlir +++ b/mlir/test/Conversion/GPUToSPIRV/simple.mlir @@ -1,6 +1,6 @@ // RUN: mlir-opt -convert-gpu-to-spirv %s -o - | FileCheck %s -// CHECK: spv.module "Logical" "VulkanKHR" { +// CHECK: spv.module "Logical" "GLSL450" { // CHECK-NEXT: spv.globalVariable [[VAR1:@.*]] bind(0, 0) : !spv.ptr // CHECK-NEXT: spv.globalVariable [[VAR2:@.*]] bind(0, 1) : !spv.ptr, StorageBuffer> // CHECK-NEXT: func @kernel_1 diff --git a/mlir/test/Dialect/SPIRV/Serialization/access_chain.mlir b/mlir/test/Dialect/SPIRV/Serialization/access_chain.mlir index e2172f6..34af0fd 100644 --- a/mlir/test/Dialect/SPIRV/Serialization/access_chain.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/access_chain.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -serialize-spirv %s | mlir-translate -deserialize-spirv | FileCheck %s -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @access_chain(%arg0 : !spv.ptr>, Function>, %arg1 : i32, %arg2 : i32) { // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}] : !spv.ptr>, Function> diff --git a/mlir/test/Dialect/SPIRV/Serialization/array_stride.mlir b/mlir/test/Dialect/SPIRV/Serialization/array_stride.mlir index fa76c8a..9356c7d 100644 --- a/mlir/test/Dialect/SPIRV/Serialization/array_stride.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/array_stride.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -serialize-spirv %s | mlir-translate -deserialize-spirv | FileCheck %s -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @array_stride(%arg0 : !spv.ptr [128]>, StorageBuffer>, %arg1 : i32, %arg2 : i32) { // CHECK: {{%.*}} = spv.AccessChain {{%.*}}[{{%.*}}, {{%.*}}] : !spv.ptr [128]>, StorageBuffer> diff --git a/mlir/test/Dialect/SPIRV/Serialization/bin_ops.mlir b/mlir/test/Dialect/SPIRV/Serialization/bin_ops.mlir index 891d9fd..0a699b2 100644 --- a/mlir/test/Dialect/SPIRV/Serialization/bin_ops.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/bin_ops.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -serialize-spirv %s | mlir-translate -deserialize-spirv | FileCheck %s -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @fmul(%arg0 : f32, %arg1 : f32) { // CHECK: {{%.*}}= spv.FMul {{%.*}}, {{%.*}} : f32 %0 = spv.FMul %arg0, %arg1 : f32 diff --git a/mlir/test/Dialect/SPIRV/Serialization/entry.mlir b/mlir/test/Dialect/SPIRV/Serialization/entry.mlir index cd48e20..c58134a 100644 --- a/mlir/test/Dialect/SPIRV/Serialization/entry.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/entry.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -serialize-spirv %s | mlir-translate -deserialize-spirv | FileCheck %s -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @noop() -> () { spv.Return } diff --git a/mlir/test/Dialect/SPIRV/Serialization/entry_interface.mlir b/mlir/test/Dialect/SPIRV/Serialization/entry_interface.mlir index 1c8488f..924318b 100644 --- a/mlir/test/Dialect/SPIRV/Serialization/entry_interface.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/entry_interface.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -serialize-spirv %s | mlir-translate -deserialize-spirv | FileCheck %s -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { // CHECK: spv.globalVariable @var2 : !spv.ptr // CHECK-NEXT: spv.globalVariable @var3 : !spv.ptr // CHECK-NEXT: func @noop({{%.*}}: !spv.ptr, {{%.*}}: !spv.ptr) diff --git a/mlir/test/Dialect/SPIRV/Serialization/execution_mode.mlir b/mlir/test/Dialect/SPIRV/Serialization/execution_mode.mlir index b8bd230..278fe1b 100644 --- a/mlir/test/Dialect/SPIRV/Serialization/execution_mode.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/execution_mode.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -serialize-spirv %s | mlir-translate -deserialize-spirv | FileCheck %s -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @foo() -> () { spv.Return } diff --git a/mlir/test/Dialect/SPIRV/Serialization/load_store.mlir b/mlir/test/Dialect/SPIRV/Serialization/load_store.mlir index 07ec6da..b6c0263 100644 --- a/mlir/test/Dialect/SPIRV/Serialization/load_store.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/load_store.mlir @@ -4,7 +4,7 @@ // CHECK-NEXT: [[VALUE:%.*]] = spv.Load "Input" [[ARG1]] : f32 // CHECK-NEXT: spv.Store "Output" [[ARG2]], [[VALUE]] : f32 -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @load_store(%arg0 : !spv.ptr, %arg1 : !spv.ptr) { %1 = spv.Load "Input" %arg0 : f32 spv.Store "Output" %arg1, %1 : f32 diff --git a/mlir/test/Dialect/SPIRV/Serialization/minimal-module.mlir b/mlir/test/Dialect/SPIRV/Serialization/minimal-module.mlir index 8754a67..82b7aee 100644 --- a/mlir/test/Dialect/SPIRV/Serialization/minimal-module.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/minimal-module.mlir @@ -1,12 +1,12 @@ // RUN: mlir-translate -serialize-spirv %s | mlir-translate -deserialize-spirv | FileCheck %s -// CHECK: spv.module "Logical" "VulkanKHR" { +// CHECK: spv.module "Logical" "GLSL450" { // CHECK-NEXT: func @foo() { // CHECK-NEXT: spv.Return // CHECK-NEXT: } // CHECK-NEXT: } attributes {major_version = 1 : i32, minor_version = 0 : i32} -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @foo() -> () { spv.Return } diff --git a/mlir/test/Dialect/SPIRV/Serialization/select.mlir b/mlir/test/Dialect/SPIRV/Serialization/select.mlir index aec39e8..614d71a 100644 --- a/mlir/test/Dialect/SPIRV/Serialization/select.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/select.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -serialize-spirv %s | mlir-translate -deserialize-spirv | FileCheck %s -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { spv.specConstant @condition_scalar = true func @select() -> () { %0 = spv.constant 4.0 : f32 diff --git a/mlir/test/Dialect/SPIRV/Serialization/struct.mlir b/mlir/test/Dialect/SPIRV/Serialization/struct.mlir index 98481e1..086e5778 100644 --- a/mlir/test/Dialect/SPIRV/Serialization/struct.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/struct.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -serialize-spirv %s | mlir-translate -deserialize-spirv | FileCheck %s -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { // CHECK: !spv.ptr [0]>, Input> spv.globalVariable @var0 bind(0, 1) : !spv.ptr [0]>, Input> diff --git a/mlir/test/Dialect/SPIRV/Serialization/variable_reference.mlir b/mlir/test/Dialect/SPIRV/Serialization/variable_reference.mlir index 5441976..9e5d9e1 100644 --- a/mlir/test/Dialect/SPIRV/Serialization/variable_reference.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/variable_reference.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -serialize-spirv %s | mlir-translate -deserialize-spirv | FileCheck %s -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { spv.globalVariable @globalInvocationID built_in("GlobalInvocationId") : !spv.ptr, Input> func @foo() { // CHECK: %[[ADDR:.*]] = spv._address_of @globalInvocationID : !spv.ptr, Input> diff --git a/mlir/test/Dialect/SPIRV/Serialization/variables.mlir b/mlir/test/Dialect/SPIRV/Serialization/variables.mlir index d87b944..990a1a6 100644 --- a/mlir/test/Dialect/SPIRV/Serialization/variables.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/variables.mlir @@ -5,7 +5,7 @@ // CHECK-NEXT: spv.globalVariable @var2 built_in("GlobalInvocationId") : !spv.ptr, Input> // CHECK-NEXT: spv.globalVariable @var3 built_in("GlobalInvocationId") : !spv.ptr, Input> -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { spv.globalVariable @var0 bind(1, 0) : !spv.ptr spv.globalVariable @var1 bind(0, 1) : !spv.ptr spv.globalVariable @var2 {built_in = "GlobalInvocationId"} : !spv.ptr, Input> diff --git a/mlir/test/Dialect/SPIRV/Serialization/variables_init.mlir b/mlir/test/Dialect/SPIRV/Serialization/variables_init.mlir index 7432d48..64849ff 100644 --- a/mlir/test/Dialect/SPIRV/Serialization/variables_init.mlir +++ b/mlir/test/Dialect/SPIRV/Serialization/variables_init.mlir @@ -1,6 +1,6 @@ // RUN: mlir-translate -serialize-spirv %s | mlir-translate -deserialize-spirv | FileCheck %s -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { // CHECK: spv.globalVariable @var1 : !spv.ptr // CHECK-NEXT: spv.globalVariable @var2 initializer(@var1) bind(1, 0) : !spv.ptr spv.globalVariable @var1 : !spv.ptr diff --git a/mlir/test/Dialect/SPIRV/control-flow-ops.mlir b/mlir/test/Dialect/SPIRV/control-flow-ops.mlir index 8199d0e..c3638a8 100644 --- a/mlir/test/Dialect/SPIRV/control-flow-ops.mlir +++ b/mlir/test/Dialect/SPIRV/control-flow-ops.mlir @@ -340,7 +340,7 @@ func @only_allowed_in_last_block() -> () { // ----- // Return mismatches function signature -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @work() -> (i32) { // expected-error @+1 {{cannot be used in functions returning value}} spv.Return diff --git a/mlir/test/Dialect/SPIRV/ops.mlir b/mlir/test/Dialect/SPIRV/ops.mlir index 524f1d2..348a685 100644 --- a/mlir/test/Dialect/SPIRV/ops.mlir +++ b/mlir/test/Dialect/SPIRV/ops.mlir @@ -266,7 +266,7 @@ func @composite_extract_result_type_mismatch(%arg0: !spv.array<4xf32>) -> i32 { // spv.ExecutionMode //===----------------------------------------------------------------------===// -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @do_nothing() -> () { spv.Return } @@ -275,7 +275,7 @@ spv.module "Logical" "VulkanKHR" { spv.ExecutionMode @do_nothing "ContractionOff" } -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @do_nothing() -> () { spv.Return } @@ -286,7 +286,7 @@ spv.module "Logical" "VulkanKHR" { // ----- -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @do_nothing() -> () { spv.Return } @@ -417,7 +417,7 @@ func @aligned_load_incorrect_attributes() -> () { // ----- -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { spv.globalVariable @var0 : !spv.ptr // CHECK_LABEL: @simple_load func @simple_load() -> () { @@ -644,7 +644,7 @@ func @aligned_store_incorrect_attributes(%arg0 : f32) -> () { // ----- -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { spv.globalVariable @var0 : !spv.ptr func @simple_store(%arg0 : f32) -> () { %0 = spv._address_of @var0 : !spv.ptr diff --git a/mlir/test/Dialect/SPIRV/structure-ops.mlir b/mlir/test/Dialect/SPIRV/structure-ops.mlir index 356855d..271edc8 100644 --- a/mlir/test/Dialect/SPIRV/structure-ops.mlir +++ b/mlir/test/Dialect/SPIRV/structure-ops.mlir @@ -4,7 +4,7 @@ // spv._address_of //===----------------------------------------------------------------------===// -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { spv.globalVariable @var1 : !spv.ptr>, Input> func @access_chain() -> () { %0 = spv.constant 1: i32 @@ -18,7 +18,7 @@ spv.module "Logical" "VulkanKHR" { // ----- -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { spv.globalVariable @var1 : !spv.ptr>, Input> func @foo() -> () { // expected-error @+1 {{expected spv.globalVariable symbol}} @@ -28,7 +28,7 @@ spv.module "Logical" "VulkanKHR" { // ----- -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { spv.globalVariable @var1 : !spv.ptr>, Input> func @foo() -> () { // expected-error @+1 {{result type mismatch with the referenced global variable's type}} @@ -95,7 +95,7 @@ func @value_result_type_mismatch() -> () { // spv.EntryPoint //===----------------------------------------------------------------------===// -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @do_nothing() -> () { spv.Return } @@ -103,7 +103,7 @@ spv.module "Logical" "VulkanKHR" { spv.EntryPoint "GLCompute" @do_nothing } -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { spv.globalVariable @var2 : !spv.ptr spv.globalVariable @var3 : !spv.ptr func @do_something(%arg0 : !spv.ptr, %arg1 : !spv.ptr) -> () { @@ -117,7 +117,7 @@ spv.module "Logical" "VulkanKHR" { // ----- -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @do_nothing() -> () { spv.Return } @@ -127,7 +127,7 @@ spv.module "Logical" "VulkanKHR" { // ----- -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @do_nothing() -> () { spv.Return } @@ -142,7 +142,7 @@ spv.module "Logical" "VulkanKHR" { // ----- -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @do_nothing() -> () { // expected-error @+1 {{'spv.EntryPoint' op failed to verify that op must appear in a 'spv.module' block}} spv.EntryPoint "GLCompute" @do_something @@ -151,7 +151,7 @@ spv.module "Logical" "VulkanKHR" { // ----- -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @do_nothing() -> () { spv.Return } @@ -162,7 +162,7 @@ spv.module "Logical" "VulkanKHR" { // ----- -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @do_nothing() -> () { spv.Return } @@ -177,13 +177,13 @@ spv.module "Logical" "VulkanKHR" { // spv.globalVariable //===----------------------------------------------------------------------===// -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { // CHECK: spv.globalVariable @var0 : !spv.ptr spv.globalVariable @var0 : !spv.ptr } // TODO: Fix test case after initialization with normal constant is addressed -// spv.module "Logical" "VulkanKHR" { +// spv.module "Logical" "GLSL450" { // %0 = spv.constant 4.0 : f32 // // CHECK1: spv.Variable init(%0) : !spv.ptr // spv.globalVariable @var1 init(%0) : !spv.ptr @@ -191,7 +191,7 @@ spv.module "Logical" "VulkanKHR" { // ----- -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { spv.specConstant @sc = 4.0 : f32 // CHECK: spv.globalVariable @var initializer(@sc) : !spv.ptr spv.globalVariable @var initializer(@sc) : !spv.ptr @@ -199,13 +199,13 @@ spv.module "Logical" "VulkanKHR" { // ----- -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { // CHECK: spv.globalVariable @var0 bind(1, 2) : !spv.ptr spv.globalVariable @var0 bind(1, 2) : !spv.ptr } // TODO: Fix test case after initialization with constant is addressed -// spv.module "Logical" "VulkanKHR" { +// spv.module "Logical" "GLSL450" { // %0 = spv.constant 4.0 : f32 // // CHECK1: spv.globalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr // spv.globalVariable @var1 initializer(%0) {binding = 5 : i32} : !spv.ptr @@ -213,7 +213,7 @@ spv.module "Logical" "VulkanKHR" { // ----- -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { // CHECK: spv.globalVariable @var1 built_in("GlobalInvocationID") : !spv.ptr, Input> spv.globalVariable @var1 built_in("GlobalInvocationID") : !spv.ptr, Input> // CHECK: spv.globalVariable @var2 built_in("GlobalInvocationID") : !spv.ptr, Input> @@ -222,28 +222,28 @@ spv.module "Logical" "VulkanKHR" { // ----- -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { // expected-error @+1 {{expected spv.ptr type}} spv.globalVariable @var0 : f32 } // ----- -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { // expected-error @+1 {{op initializer must be result of a spv.specConstant or spv.globalVariable op}} spv.globalVariable @var0 initializer(@var1) : !spv.ptr } // ----- -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { // expected-error @+1 {{storage class cannot be 'Generic'}} spv.globalVariable @var0 : !spv.ptr } // ----- -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @foo() { // expected-error @+1 {{op failed to verify that op must appear in a 'spv.module' block}} spv.globalVariable @var0 : !spv.ptr @@ -258,25 +258,25 @@ spv.module "Logical" "VulkanKHR" { //===----------------------------------------------------------------------===// // Module without capability and extension -// CHECK: spv.module "Logical" "VulkanKHR" -spv.module "Logical" "VulkanKHR" { } +// CHECK: spv.module "Logical" "GLSL450" +spv.module "Logical" "GLSL450" { } // Module with capability and extension // CHECK: attributes {capability = ["Shader"], extension = ["SPV_KHR_16bit_storage"]} -spv.module "Logical" "VulkanKHR" { } attributes { +spv.module "Logical" "GLSL450" { } attributes { capability = ["Shader"], extension = ["SPV_KHR_16bit_storage"] } // Module with explict spv._module_end // CHECK: spv.module -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { spv._module_end } // Module with function // CHECK: spv.module -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @do_nothing() -> () { spv.Return } @@ -310,7 +310,7 @@ spv.module "Logical" "Bla" { } // Module with multiple blocks // expected-error @+1 {{expects region #0 to have 0 or 1 blocks}} -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { ^first: spv.Return ^second: @@ -329,7 +329,7 @@ spv.module "Logical" "VulkanKHR" { // ----- // Use non SPIR-V op inside.module -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { // expected-error @+1 {{'spv.module' can only contain func and spv.* ops}} "dialect.op"() : () -> () } @@ -337,7 +337,7 @@ spv.module "Logical" "VulkanKHR" { // ----- // Use non SPIR-V op inside function -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @do_nothing() -> () { // expected-error @+1 {{functions in 'spv.module' can only contain spv.* ops}} "dialect.op"() : () -> () @@ -347,7 +347,7 @@ spv.module "Logical" "VulkanKHR" { // ----- // Use external function -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { // expected-error @+1 {{'spv.module' cannot contain external functions}} func @extern() -> () } @@ -355,7 +355,7 @@ spv.module "Logical" "VulkanKHR" { // ----- // Module with nested function -spv.module "Logical" "VulkanKHR" { +spv.module "Logical" "GLSL450" { func @outer_func() -> () { // expected-error @+1 {{'spv.module' cannot contain nested functions}} func @inner_func() -> () { diff --git a/mlir/utils/spirv/gen_spirv_dialect.py b/mlir/utils/spirv/gen_spirv_dialect.py index e74a40e..cca152f 100755 --- a/mlir/utils/spirv/gen_spirv_dialect.py +++ b/mlir/utils/spirv/gen_spirv_dialect.py @@ -379,7 +379,7 @@ def get_op_definition(instruction, doc, existing_info, inst_category): 'form\n{assembly}}}];\n') if inst_category == 'Op': fmt_str +='\n let arguments = (ins{args});\n\n'\ - ' let results = (outs{results});\n\n' + ' let results = (outs{results});\n' fmt_str +='{extras}'\ '}}\n' -- 2.7.4