From 8360e8a748fad483a6a1ee621fbdbad6aa5b0c33 Mon Sep 17 00:00:00 2001 From: nrudenko Date: Wed, 27 May 2020 15:30:55 +0300 Subject: [PATCH 1/4] Refine condition for CapabilityVector16 --- lib/SPIRV/libSPIRV/SPIRVType.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/lib/SPIRV/libSPIRV/SPIRVType.h b/lib/SPIRV/libSPIRV/SPIRVType.h index c717bc005f..1ce1d3058f 100644 --- a/lib/SPIRV/libSPIRV/SPIRVType.h +++ b/lib/SPIRV/libSPIRV/SPIRVType.h @@ -305,7 +305,7 @@ class SPIRVTypeVector : public SPIRVType { SPIRVCapVec V(getComponentType()->getRequiredCapability()); // Even though the capability name is "Vector16", it describes // usage of 8-component or 16-component vectors. - if (CompCount >= 8) + if (CompCount == 8 || CompCount == 16) V.push_back(CapabilityVector16); return V; } From f2e321407b295e96722b6b006634278441b1220b Mon Sep 17 00:00:00 2001 From: nrudenko Date: Wed, 27 May 2020 15:39:32 +0300 Subject: [PATCH 2/4] Remove redundant assert from SPIRVVectorShuffle This assert is redundant because the allowed size of Components is checked by Type. Moreover it can interfere with capabilities allowing non-standart sizes. --- lib/SPIRV/libSPIRV/SPIRVInstruction.h | 1 - 1 file changed, 1 deletion(-) diff --git a/lib/SPIRV/libSPIRV/SPIRVInstruction.h b/lib/SPIRV/libSPIRV/SPIRVInstruction.h index 375d08b83f..c20ad86fd3 100644 --- a/lib/SPIRV/libSPIRV/SPIRVInstruction.h +++ b/lib/SPIRV/libSPIRV/SPIRVInstruction.h @@ -2275,7 +2275,6 @@ class SPIRVVectorShuffle : public SPIRVInstruction { return; assert(getValueType(Vector1) == getValueType(Vector2)); assert(Components.size() == Type->getVectorComponentCount()); - assert(Components.size() > 1); } SPIRVId Vector1; SPIRVId Vector2; From 946a1b232b51724b1b514f2d2cc105f901eb6b0d Mon Sep 17 00:00:00 2001 From: nrudenko Date: Wed, 27 May 2020 16:03:20 +0300 Subject: [PATCH 3/4] Add required capabilities for ExecutionModeSubgroupSize --- lib/SPIRV/libSPIRV/SPIRVEnum.h | 1 + lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h | 1 + test/transcoding/ReqdSubgroupSize.ll | 1 + 3 files changed, 3 insertions(+) diff --git a/lib/SPIRV/libSPIRV/SPIRVEnum.h b/lib/SPIRV/libSPIRV/SPIRVEnum.h index 95555240d9..6d24877a02 100644 --- a/lib/SPIRV/libSPIRV/SPIRVEnum.h +++ b/lib/SPIRV/libSPIRV/SPIRVEnum.h @@ -237,6 +237,7 @@ template <> inline void SPIRVMap::init() { ADD_VEC_INIT(ExecutionModeOutputTriangleStrip, {CapabilityGeometry}); ADD_VEC_INIT(ExecutionModeVecTypeHint, {CapabilityKernel}); ADD_VEC_INIT(ExecutionModeContractionOff, {CapabilityKernel}); + ADD_VEC_INIT(ExecutionModeSubgroupSize, {CapabilitySubgroupDispatch}); ADD_VEC_INIT(ExecutionModeDenormPreserve, {CapabilityDenormPreserve}); ADD_VEC_INIT(ExecutionModeDenormFlushToZero, {CapabilityDenormFlushToZero}); ADD_VEC_INIT(ExecutionModeSignedZeroInfNanPreserve, diff --git a/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h b/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h index 9ef434659a..0a62838dff 100644 --- a/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h +++ b/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h @@ -495,6 +495,7 @@ template <> inline void SPIRVMap::init() { add(CapabilityStorageImageWriteWithoutFormat, "StorageImageWriteWithoutFormat"); add(CapabilityMultiViewport, "MultiViewport"); + add(CapabilitySubgroupDispatch, "CapabilitySubgroupDispatch"); add(CapabilityDenormPreserve, "DenormPreserve"); add(CapabilityDenormFlushToZero, "DenormFlushToZero"); add(CapabilitySignedZeroInfNanPreserve, "SignedZeroInfNanPreserve"); diff --git a/test/transcoding/ReqdSubgroupSize.ll b/test/transcoding/ReqdSubgroupSize.ll index 703a87f5e6..6978c3d424 100644 --- a/test/transcoding/ReqdSubgroupSize.ll +++ b/test/transcoding/ReqdSubgroupSize.ll @@ -7,6 +7,7 @@ ; RUN: llvm-spirv %t.spv -to-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV ; RUN: llvm-spirv %t.spv -r -o - | llvm-dis -o - | FileCheck %s --check-prefix=CHECK-LLVM +; CHECK-SPIRV: Capability CapabilitySubgroupDispatch ; CHECK-SPIRV: EntryPoint 6 [[kernel:[0-9]+]] "foo" ; CHECK-SPIRV: ExecutionMode [[kernel]] 35 8 From aaacb91a192b769a17e1a8371381d477d641602c Mon Sep 17 00:00:00 2001 From: nrudenko Date: Wed, 27 May 2020 17:30:08 +0300 Subject: [PATCH 4/4] Add missing float control capabilities to SPIRVIsValidEnum.h --- lib/SPIRV/libSPIRV/SPIRVIsValidEnum.h | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/lib/SPIRV/libSPIRV/SPIRVIsValidEnum.h b/lib/SPIRV/libSPIRV/SPIRVIsValidEnum.h index a6000908b2..3db6bb9f8a 100644 --- a/lib/SPIRV/libSPIRV/SPIRVIsValidEnum.h +++ b/lib/SPIRV/libSPIRV/SPIRVIsValidEnum.h @@ -581,6 +581,11 @@ inline bool isValid(spv::Capability V) { case CapabilityGroupNonUniformShuffleRelative: case CapabilityGroupNonUniformClustered: case CapabilityGroupNonUniformQuad: + case CapabilityDenormPreserve: + case CapabilityDenormFlushToZero: + case CapabilitySignedZeroInfNanPreserve: + case CapabilityRoundingModeRTE: + case CapabilityRoundingModeRTZ: case CapabilityFPGAMemoryAttributesINTEL: case CapabilityArbitraryPrecisionIntegersINTEL: case CapabilityFPGALoopControlsINTEL: