Minor changes and cleanups (#553)

* Refined condition for CapabilityVector16
* Removed 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.

* Added required capabilities for ExecutionModeSubgroupSize
* Added missing float control capabilities to SPIRVIsValidEnum.h

Author: NikitaRudenkoIntel

llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEnum.h

@@ -237,6 +237,7 @@ template <> inline void SPIRVMap<SPIRVExecutionModeKind, SPIRVCapVec>::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,

llvm-spirv/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;

llvm-spirv/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:

llvm-spirv/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h

@@ -495,6 +495,7 @@ template <> inline void SPIRVMap<Capability, std::string>::init() {
   add(CapabilityStorageImageWriteWithoutFormat,
       "StorageImageWriteWithoutFormat");
   add(CapabilityMultiViewport, "MultiViewport");
+  add(CapabilitySubgroupDispatch, "CapabilitySubgroupDispatch");
   add(CapabilityDenormPreserve, "DenormPreserve");
   add(CapabilityDenormFlushToZero, "DenormFlushToZero");
   add(CapabilitySignedZeroInfNanPreserve, "SignedZeroInfNanPreserve");

llvm-spirv/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;
   }

llvm-spirv/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