From d934838ef91ff00ff62c1dbfbf1dd0a9fb65f59e Mon Sep 17 00:00:00 2001
From: Dmitry Sidorov <dmitry.sidorov@intel.com>
Date: Fri, 10 Sep 2021 13:46:26 +0300
Subject: [PATCH 1/3] Add SPIR-V 1.4 checks

Signed-off-by: Dmitry Sidorov <dmitry.sidorov@intel.com>

Original commit:
https://github.com/KhronosGroup/SPIRV-LLVM-Translator/commit/c5b3c8e3283b
---
 llvm-spirv/lib/SPIRV/SPIRVWriter.cpp          | 83 ++++++++++---------
 llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h |  5 +-
 llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h    |  6 --
 llvm-spirv/lib/SPIRV/libSPIRV/SPIRVValue.cpp  | 69 ++++++++-------
 llvm-spirv/lib/SPIRV/libSPIRV/SPIRVValue.h    |  6 +-
 .../test/exec_mode_float_control_khr.ll       | 14 +++-
 llvm-spirv/test/transcoding/LoopUnroll.ll     |  8 ++
 .../test/transcoding/NoSignedUnsignedWrap.ll  | 43 +++++-----
 .../exec_mode_float_control.ll                | 11 ++-
 .../test/transcoding/annotate_attribute.ll    |  3 +
 .../test/transcoding/block_w_struct_return.cl | 14 +++-
 llvm-spirv/test/transcoding/global_block.cl   | 14 +++-
 12 files changed, 161 insertions(+), 115 deletions(-)

diff --git a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp
index 4f7fa2f06102d..b8479c64f65e6 100644
--- a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp
+++ b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp
@@ -1195,9 +1195,13 @@ LLVMToSPIRVBase::getLoopControl(const BranchInst *Branch,
       // PartialCount must not be used with the DontUnroll bit
       else if (S == "llvm.loop.unroll.count" &&
                !(LoopControl & LoopControlDontUnrollMask)) {
-        size_t I = getMDOperandAsInt(Node, 1);
-        ParametersToSort.emplace_back(spv::LoopControlPartialCountMask, I);
-        LoopControl |= spv::LoopControlPartialCountMask;
+        if (BM->isAllowedToUseVersion(VersionNumber::SPIRV_1_4)) {
+          BM->setMinSPIRVVersion(
+              static_cast<SPIRVWord>(VersionNumber::SPIRV_1_4));
+          size_t I = getMDOperandAsInt(Node, 1);
+          ParametersToSort.emplace_back(spv::LoopControlPartialCountMask, I);
+          LoopControl |= spv::LoopControlPartialCountMask;
+        }
       } else if (S == "llvm.loop.ivdep.enable")
         LoopControl |= spv::LoopControlDependencyInfiniteMask;
       else if (S == "llvm.loop.ivdep.safelen") {
@@ -2008,10 +2012,10 @@ bool LLVMToSPIRVBase::transDecoration(Value *V, SPIRVValue *BV) {
 
   if (auto BVO = dyn_cast_or_null<OverflowingBinaryOperator>(V)) {
     if (BVO->hasNoSignedWrap()) {
-      BV->setNoSignedWrap(true);
+      BV->setNoIntegerDecorationWrap<DecorationNoSignedWrap>(true);
     }
     if (BVO->hasNoUnsignedWrap()) {
-      BV->setNoUnsignedWrap(true);
+      BV->setNoIntegerDecorationWrap<DecorationNoUnsignedWrap>(true);
     }
   }
 
@@ -3853,6 +3857,12 @@ bool LLVMToSPIRVBase::transExecutionMode() {
       if (!BF)
         return false;
 
+      auto AddSingleArgExecutionMode = [&](ExecutionMode EMode) {
+        uint32_t Arg;
+        N.get(Arg);
+        BF->addExecutionMode(BM->add(new SPIRVExecutionMode(BF, EMode, Arg)));
+      };
+
       switch (EMode) {
       case spv::ExecutionModeContractionOff:
         BF->addExecutionMode(BM->add(
@@ -3888,43 +3898,34 @@ bool LLVMToSPIRVBase::transExecutionMode() {
         }
       } break;
       case spv::ExecutionModeNoGlobalOffsetINTEL: {
-        if (BM->isAllowedToUseExtension(
-                ExtensionID::SPV_INTEL_kernel_attributes)) {
-          BF->addExecutionMode(BM->add(
-              new SPIRVExecutionMode(BF, static_cast<ExecutionMode>(EMode))));
-          BM->addExtension(ExtensionID::SPV_INTEL_kernel_attributes);
-          BM->addCapability(CapabilityKernelAttributesINTEL);
-        }
+        if (!BM->isAllowedToUseExtension(
+                ExtensionID::SPV_INTEL_kernel_attributes))
+          break;
+        BF->addExecutionMode(BM->add(
+            new SPIRVExecutionMode(BF, static_cast<ExecutionMode>(EMode))));
+        BM->addExtension(ExtensionID::SPV_INTEL_kernel_attributes);
+        BM->addCapability(CapabilityKernelAttributesINTEL);
       } break;
       case spv::ExecutionModeVecTypeHint:
       case spv::ExecutionModeSubgroupSize:
-      case spv::ExecutionModeSubgroupsPerWorkgroup: {
-        unsigned X;
-        N.get(X);
-        BF->addExecutionMode(BM->add(
-            new SPIRVExecutionMode(BF, static_cast<ExecutionMode>(EMode), X)));
-      } break;
+      case spv::ExecutionModeSubgroupsPerWorkgroup:
+        AddSingleArgExecutionMode(static_cast<ExecutionMode>(EMode));
+        break;
       case spv::ExecutionModeNumSIMDWorkitemsINTEL:
       case spv::ExecutionModeSchedulerTargetFmaxMhzINTEL:
       case spv::ExecutionModeMaxWorkDimINTEL:
       case spv::internal::ExecutionModeStreamingInterfaceINTEL: {
-        if (BM->isAllowedToUseExtension(
-                ExtensionID::SPV_INTEL_kernel_attributes)) {
-          unsigned X;
-          N.get(X);
-          BF->addExecutionMode(BM->add(new SPIRVExecutionMode(
-              BF, static_cast<ExecutionMode>(EMode), X)));
-          BM->addExtension(ExtensionID::SPV_INTEL_kernel_attributes);
-          BM->addCapability(CapabilityFPGAKernelAttributesINTEL);
-        }
+        if (!BM->isAllowedToUseExtension(
+                ExtensionID::SPV_INTEL_kernel_attributes))
+          break;
+        AddSingleArgExecutionMode(static_cast<ExecutionMode>(EMode));
+        BM->addExtension(ExtensionID::SPV_INTEL_kernel_attributes);
+        BM->addCapability(CapabilityFPGAKernelAttributesINTEL);
       } break;
       case spv::ExecutionModeSharedLocalMemorySizeINTEL: {
         if (!BM->isAllowedToUseExtension(ExtensionID::SPV_INTEL_vector_compute))
           break;
-        unsigned SLMSize;
-        N.get(SLMSize);
-        BF->addExecutionMode(BM->add(new SPIRVExecutionMode(
-            BF, static_cast<ExecutionMode>(EMode), SLMSize)));
+        AddSingleArgExecutionMode(static_cast<ExecutionMode>(EMode));
       } break;
 
       case spv::ExecutionModeDenormPreserve:
@@ -3932,12 +3933,15 @@ bool LLVMToSPIRVBase::transExecutionMode() {
       case spv::ExecutionModeSignedZeroInfNanPreserve:
       case spv::ExecutionModeRoundingModeRTE:
       case spv::ExecutionModeRoundingModeRTZ: {
-        if (!BM->isAllowedToUseExtension(ExtensionID::SPV_KHR_float_controls))
-          break;
-        unsigned TargetWidth;
-        N.get(TargetWidth);
-        BF->addExecutionMode(BM->add(new SPIRVExecutionMode(
-            BF, static_cast<ExecutionMode>(EMode), TargetWidth)));
+        if (BM->isAllowedToUseVersion(VersionNumber::SPIRV_1_4)) {
+          BM->setMinSPIRVVersion(
+              static_cast<SPIRVWord>(VersionNumber::SPIRV_1_4));
+          AddSingleArgExecutionMode(static_cast<ExecutionMode>(EMode));
+        } else if (BM->isAllowedToUseExtension(
+                       ExtensionID::SPV_KHR_float_controls)) {
+          BM->addExtension(ExtensionID::SPV_KHR_float_controls);
+          AddSingleArgExecutionMode(static_cast<ExecutionMode>(EMode));
+        }
       } break;
       case spv::ExecutionModeRoundingModeRTPINTEL:
       case spv::ExecutionModeRoundingModeRTNINTEL:
@@ -3946,10 +3950,7 @@ bool LLVMToSPIRVBase::transExecutionMode() {
         if (!BM->isAllowedToUseExtension(
                 ExtensionID::SPV_INTEL_float_controls2))
           break;
-        unsigned TargetWidth;
-        N.get(TargetWidth);
-        BF->addExecutionMode(BM->add(new SPIRVExecutionMode(
-            BF, static_cast<ExecutionMode>(EMode), TargetWidth)));
+        AddSingleArgExecutionMode(static_cast<ExecutionMode>(EMode));
       } break;
       case spv::internal::ExecutionModeFastCompositeKernelINTEL: {
         if (BM->isAllowedToUseExtension(ExtensionID::SPV_INTEL_fast_composite))
diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h
index b200015821d6b..3b8e0065d371b 100644
--- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h
+++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h
@@ -94,6 +94,8 @@ class SPIRVDecorateGeneric : public SPIRVAnnotationGeneric {
 
     case DecorationMaxByteOffset:
       return static_cast<SPIRVWord>(VersionNumber::SPIRV_1_1);
+    case DecorationUserSemantic:
+      return static_cast<SPIRVWord>(VersionNumber::SPIRV_1_4);
 
     default:
       return static_cast<SPIRVWord>(VersionNumber::SPIRV_1_0);
@@ -127,9 +129,6 @@ class SPIRVDecorate : public SPIRVDecorateGeneric {
 
   llvm::Optional<ExtensionID> getRequiredExtension() const override {
     switch (static_cast<size_t>(Dec)) {
-    case DecorationNoSignedWrap:
-    case DecorationNoUnsignedWrap:
-      return ExtensionID::SPV_KHR_no_integer_wrap_decoration;
     case DecorationRegisterINTEL:
     case DecorationMemoryINTEL:
     case DecorationNumbanksINTEL:
diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h
index acb2aab1f4ab4..2c6f5cacc2c90 100644
--- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h
+++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h
@@ -841,12 +841,6 @@ class SPIRVCapability : public SPIRVEntryNoId<OpCapability> {
 
   llvm::Optional<ExtensionID> getRequiredExtension() const override {
     switch (static_cast<unsigned>(Kind)) {
-    case CapabilityDenormPreserve:
-    case CapabilityDenormFlushToZero:
-    case CapabilitySignedZeroInfNanPreserve:
-    case CapabilityRoundingModeRTE:
-    case CapabilityRoundingModeRTZ:
-      return ExtensionID::SPV_KHR_float_controls;
     case CapabilityRoundToInfinityINTEL:
     case CapabilityFloatingPointModeINTEL:
     case CapabilityFunctionFloatControlINTEL:
diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVValue.cpp b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVValue.cpp
index 3b466f05fabbd..2356c05ddb39e 100644
--- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVValue.cpp
+++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVValue.cpp
@@ -75,45 +75,10 @@ bool SPIRVValue::hasNoSignedWrap() const {
   return hasDecorate(DecorationNoSignedWrap);
 }
 
-void SPIRVValue::setNoSignedWrap(bool HasNoSignedWrap) {
-  if (!HasNoSignedWrap) {
-    eraseDecorate(DecorationNoSignedWrap);
-  }
-  if (Module->isAllowedToUseExtension(
-          ExtensionID::SPV_KHR_no_integer_wrap_decoration)) {
-    // NoSignedWrap decoration is available only if it is allowed to use SPIR-V
-    // 1.4 or if SPV_KHR_no_integer_wrap_decoration extension is allowed
-    // FIXME: update this 'if' to include check for SPIR-V 1.4 once translator
-    // support this version
-    addDecorate(new SPIRVDecorate(DecorationNoSignedWrap, this));
-    SPIRVDBG(spvdbgs() << "Set nsw for obj " << Id << "\n")
-  } else {
-    SPIRVDBG(spvdbgs() << "Skip setting nsw for obj " << Id << "\n")
-  }
-}
-
 bool SPIRVValue::hasNoUnsignedWrap() const {
   return hasDecorate(DecorationNoUnsignedWrap);
 }
 
-void SPIRVValue::setNoUnsignedWrap(bool HasNoUnsignedWrap) {
-  if (!HasNoUnsignedWrap) {
-    eraseDecorate(DecorationNoUnsignedWrap);
-    return;
-  }
-  if (Module->isAllowedToUseExtension(
-          ExtensionID::SPV_KHR_no_integer_wrap_decoration)) {
-    // NoUnsignedWrap decoration is available only if it is allowed to use
-    // SPIR-V 1.4 or if SPV_KHR_no_integer_wrap_decoration extension is allowed
-    // FIXME: update this 'if' to include check for SPIR-V 1.4 once translator
-    // support this version
-    addDecorate(new SPIRVDecorate(DecorationNoUnsignedWrap, this));
-    SPIRVDBG(spvdbgs() << "Set nuw for obj " << Id << "\n")
-  } else {
-    SPIRVDBG(spvdbgs() << "Skip setting nuw for obj " << Id << "\n")
-  }
-}
-
 void SPIRVValue::setFPFastMathMode(SPIRVWord M) {
   if (M == 0) {
     eraseDecorate(DecorationFPFastMathMode);
@@ -124,6 +89,40 @@ void SPIRVValue::setFPFastMathMode(SPIRVWord M) {
                      << "\n")
 }
 
+template <spv::Decoration NoIntegerWrapDecoration>
+void SPIRVValue::setNoIntegerDecorationWrap(bool HasNoIntegerWrap) {
+  if (!HasNoIntegerWrap) {
+    eraseDecorate(NoIntegerWrapDecoration);
+    return;
+  }
+  // NoSignedWrap and NoUnsignedWrap decorations are available only if it is
+  // allowed to use SPIR-V 1.4 or if SPV_KHR_no_integer_wrap_decoration
+  // extension is enabled
+#ifdef _SPIRVDBG
+  const std::string InstStr =
+      NoIntegerWrapDecoration == DecorationNoSignedWrap ? "nsw" : "nuw";
+#endif // _SPIRVDBG
+  if (Module->isAllowedToUseVersion(VersionNumber::SPIRV_1_4)) {
+    Module->setMinSPIRVVersion(
+        static_cast<SPIRVWord>(VersionNumber::SPIRV_1_4));
+    addDecorate(new SPIRVDecorate(NoIntegerWrapDecoration, this));
+    SPIRVDBG(spvdbgs() << "Set " << InstStr << " for obj " << Id << "\n")
+  } else if (Module->isAllowedToUseExtension(
+                 ExtensionID::SPV_KHR_no_integer_wrap_decoration)) {
+    Module->addExtension(ExtensionID::SPV_KHR_no_integer_wrap_decoration);
+    addDecorate(new SPIRVDecorate(NoIntegerWrapDecoration, this));
+    SPIRVDBG(spvdbgs() << "Set " << InstStr << " for obj " << Id << "\n")
+  } else {
+    SPIRVDBG(spvdbgs() << "Skip setting " << InstStr << " for obj " << Id
+                       << "\n")
+  }
+}
+
+template void
+SPIRVValue::setNoIntegerDecorationWrap<DecorationNoSignedWrap>(bool);
+template void
+SPIRVValue::setNoIntegerDecorationWrap<DecorationNoUnsignedWrap>(bool);
+
 template <spv::Op OC>
 void SPIRVConstantBase<OC>::setWords(const uint64_t *TheValue) {
   assert(TheValue && "Nullptr value");
diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVValue.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVValue.h
index c8ba404d2d315..42fb4dd79f9a5 100644
--- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVValue.h
+++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVValue.h
@@ -98,8 +98,10 @@ class SPIRVValue : public SPIRVEntry {
 
   void setAlignment(SPIRVWord);
   void setVolatile(bool IsVolatile);
-  void setNoSignedWrap(bool HasNoSignedWrap);
-  void setNoUnsignedWrap(bool HasNoUnsignedWrap);
+
+  template <spv::Decoration NoIntegerWrapDecoration>
+  void setNoIntegerDecorationWrap(bool HasNoIntegerWrap);
+
   void setFPFastMathMode(SPIRVWord FPFastMathMode);
 
   void validate() const override {
diff --git a/llvm-spirv/test/exec_mode_float_control_khr.ll b/llvm-spirv/test/exec_mode_float_control_khr.ll
index e1b1768995e4e..11430a8822dac 100644
--- a/llvm-spirv/test/exec_mode_float_control_khr.ll
+++ b/llvm-spirv/test/exec_mode_float_control_khr.ll
@@ -1,7 +1,13 @@
 ; RUN: llvm-as %s -o %t.bc
-; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_KHR_float_controls
+; RUN: llvm-spirv %t.bc -o %t.spv --spirv-max-version=1.1 --spirv-ext=+SPV_KHR_float_controls
 ; RUN: llvm-spirv %t.spv -o %t.spt --to-text
-; RUN: FileCheck %s --input-file %t.spt -check-prefix=SPV
+; RUN: FileCheck %s --input-file %t.spt -check-prefixes=SPV,SPVEXT
+; RUN: llvm-spirv %t.bc -o %t.spv --spirv-max-version=1.4
+; RUN: llvm-spirv %t.spv -o %t.spt --to-text
+; RUN: FileCheck %s --input-file %t.spt -check-prefixes=SPV,SPV14
+; RUN: llvm-spirv %t.bc -o %t.spv --spirv-max-version=1.1
+; RUN: llvm-spirv %t.spv -o %t.spt --to-text
+; RUN: FileCheck %s --input-file %t.spt -check-prefix=SPV-NEGATIVE
 
 ; ModuleID = 'float_control.bc'
 source_filename = "float_control.cpp"
@@ -44,6 +50,10 @@ entry:
 !spirv.EntryPoint = !{}
 !spirv.ExecutionMode = !{!15, !16, !17, !18, !19, !20, !21, !22, !23, !24, !25, !26, !27, !28, !29}
 
+; SPVEXT-DAG: Extension "SPV_KHR_float_controls"
+; SPV14-NOT: Extension "SPV_KHR_float_controls"
+; SPV-NEGATIVE-NOT: Extension "SPV_KHR_float_controls"
+
 ; SPV-DAG: EntryPoint {{[0-9]+}} [[KERNEL0:[0-9]+]] "k_float_controls_0"
 ; SPV-DAG: EntryPoint {{[0-9]+}} [[KERNEL1:[0-9]+]] "k_float_controls_1"
 ; SPV-DAG: EntryPoint {{[0-9]+}} [[KERNEL2:[0-9]+]] "k_float_controls_2"
diff --git a/llvm-spirv/test/transcoding/LoopUnroll.ll b/llvm-spirv/test/transcoding/LoopUnroll.ll
index af1712021396f..0bd924d81fd43 100644
--- a/llvm-spirv/test/transcoding/LoopUnroll.ll
+++ b/llvm-spirv/test/transcoding/LoopUnroll.ll
@@ -42,6 +42,13 @@
 ; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
 ; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
 ; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM
+; RUN: llvm-spirv %t.bc -o %t.spv --spirv-max-version=1.1
+; RUN: llvm-spirv -to-text %t.spv -o %t.spt
+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV-NEGATIVE
+
+; Check SPIR-V versions in a format magic number + version
+; CHECK-SPIRV: 119734787 66560
+; CHECK-SPIRV-NEGATIVE: 119734787 65536
 
 target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024"
 target triple = "spir64"
@@ -115,6 +122,7 @@ while.cond:                                       ; preds = %if.end, %if.then, %
 ; Per SPIRV spec p3.23 "Unroll" loop control = 0x1
 ; CHECK-SPIRV: LoopMerge [[#MERGEBLOCK:]] [[#CONTINUE:]] 256 8
 ; CHECK-SPIRV: BranchConditional [[#]] [[#]] [[#MERGEBLOCK]]
+; CHECK-SPIRV-NEGATIVE-NOT: LoopMerge {{.*}} 256
   br i1 %cmp, label %while.body, label %while.end
 
 while.body:                                       ; preds = %while.cond
diff --git a/llvm-spirv/test/transcoding/NoSignedUnsignedWrap.ll b/llvm-spirv/test/transcoding/NoSignedUnsignedWrap.ll
index 4e15d495f8ba0..a67c729840545 100644
--- a/llvm-spirv/test/transcoding/NoSignedUnsignedWrap.ll
+++ b/llvm-spirv/test/transcoding/NoSignedUnsignedWrap.ll
@@ -8,10 +8,15 @@
 ; Positive tests:
 ;
 ; RUN: llvm-as %s -o %t.bc
-; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_KHR_no_integer_wrap_decoration -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_KHR_no_integer_wrap_decoration -o %t.spv
+; RUN: llvm-spirv %t.bc --spirv-max-version=1.1 --spirv-ext=+SPV_KHR_no_integer_wrap_decoration -spirv-text -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-EXT
+; RUN: llvm-spirv %t.bc --spirv-max-version=1.1 --spirv-ext=+SPV_KHR_no_integer_wrap_decoration -o %t.spv
 ; RUN: spirv-val %t.spv
-; RUN: llvm-spirv -r %t.spv --spirv-ext=+SPV_KHR_no_integer_wrap_decoration -o %t.rev.bc
+; RUN: llvm-spirv -r %t.spv --spirv-max-version=1.1 --spirv-ext=+SPV_KHR_no_integer_wrap_decoration -o %t.rev.bc
+; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
+; RUN: llvm-spirv %t.bc --spirv-max-version=1.4 -spirv-text -o - | FileCheck %s --check-prefixes=CHECK-SPIRV,CHECK-SPIRV-NOEXT
+; RUN: llvm-spirv %t.bc --spirv-max-version=1.4 -o %t.spv
+; RUN: spirv-val %t.spv
+; RUN: llvm-spirv -r %t.spv --spirv-max-version=1.4 -o %t.rev.bc
 ; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM
 ;
 ; During consumption, any SPIR-V extension must be accepted by default
@@ -21,29 +26,29 @@
 ;
 ; Negative tests:
 ;
-; Check that translator is able to reject SPIR-V if extension is disallowed
-;
-; RUN: not llvm-spirv -r %t.spv --spirv-ext=-SPV_KHR_no_integer_wrap_decoration -o - 2>&1 | FileCheck %s --check-prefix=CHECK-INVALID-SPIRV
-;
-; Check that translator is able to skip nsw/nuw attributes if extension is disabled implicitly or explicitly
+; Check that translator is able to skip nsw/nuw attributes if extension is
+; disabled implicitly or explicitly and if max SPIR-V version is lower then 1.4
 ;
-; RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV-NOEXT
-; RUN: llvm-spirv %t.bc -o %t.spv
+; RUN: llvm-spirv %t.bc --spirv-max-version=1.1 -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV-NEGATIVE
+; RUN: llvm-spirv --spirv-max-version=1.1 %t.bc -o %t.spv
 ; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
-; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM-NOEXT
+; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM-NEGATIVE
 ;
-; RUN: llvm-spirv %t.bc --spirv-ext=-SPV_KHR_no_integer_wrap_decoration -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV-NOEXT
-; RUN: llvm-spirv %t.bc --spirv-ext=-SPV_KHR_no_integer_wrap_decoration -o %t.spv
+; RUN: llvm-spirv %t.bc --spirv-max-version=1.1 --spirv-ext=-SPV_KHR_no_integer_wrap_decoration -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV-NEGATIVE
+; RUN: llvm-spirv %t.bc --spirv-max-version=1.1 --spirv-ext=-SPV_KHR_no_integer_wrap_decoration -o %t.spv
 ; RUN: llvm-spirv -r %t.spv -o %t.rev.bc
-; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM-NOEXT
+; RUN: llvm-dis < %t.rev.bc | FileCheck %s --check-prefix=CHECK-LLVM-NEGATIVE
 
-; CHECK-SPIRV: Extension "SPV_KHR_no_integer_wrap_decoration"
+; Check SPIR-V versions in a format magic number + version
+; CHECK-SPIRV-EXT: 119734787 65536
+; CHECK-SPIRV-EXT: Extension "SPV_KHR_no_integer_wrap_decoration"
+; CHECK-SPIRV-NOEXT: 119734787 66560
 ; CHECK-SPIRV-DAG: Decorate {{[0-9]+}} NoSignedWrap
 ; CHECK-SPIRV-DAG: Decorate {{[0-9]+}} NoUnsignedWrap
 ;
-; CHECK-SPIRV-NOEXT-NOT: Extension "SPV_KHR_no_integer_wrap_decoration"
-; CHECK-SPIRV-NOEXT-NOT: Decorate {{[0-9]+}} NoSignedWrap
-; CHECK-SPIRV-NOEXT-NOT: Decorate {{[0-9]+}} NoUnsignedWrap
+; CHECK-SPIRV-NEGATIVE-NOT: Extension "SPV_KHR_no_integer_wrap_decoration"
+; CHECK-SPIRV-NEGATIVE-NOT: Decorate {{[0-9]+}} NoSignedWrap
+; CHECK-SPIRV-NEGATIVE-NOT: Decorate {{[0-9]+}} NoUnsignedWrap
 ;
 ; CHECK-INVALID-SPIRV: input SPIR-V module uses extension 'SPV_KHR_no_integer_wrap_decoration' which were disabled
 
@@ -55,7 +60,7 @@ define spir_func i32 @square(i16 zeroext %a) local_unnamed_addr #0 {
 entry:
   %conv = zext i16 %a to i32
   ; CHECK-LLVM: mul nuw nsw
-  ; CHECK-LLVM-NOEXT: mul i32
+  ; CHECK-LLVM-NEGATIVE: mul i32
   %mul = mul nuw nsw i32 %conv, %conv
   ret i32 %mul
 }
diff --git a/llvm-spirv/test/transcoding/SPV_INTEL_vector_compute/exec_mode_float_control.ll b/llvm-spirv/test/transcoding/SPV_INTEL_vector_compute/exec_mode_float_control.ll
index d14f10c18aba6..ddec4135cc8e8 100644
--- a/llvm-spirv/test/transcoding/SPV_INTEL_vector_compute/exec_mode_float_control.ll
+++ b/llvm-spirv/test/transcoding/SPV_INTEL_vector_compute/exec_mode_float_control.ll
@@ -1,11 +1,15 @@
 ; RUN: llvm-as %s -o %t.bc
-; RUN: llvm-spirv %t.bc -o %t.spv --spirv-ext=+SPV_INTEL_vector_compute,+SPV_KHR_float_controls,+SPV_INTEL_float_controls2
+; RUN: llvm-spirv %t.bc -o %t.spv --spirv-max-version=1.1 --spirv-ext=+SPV_INTEL_vector_compute,+SPV_KHR_float_controls,+SPV_INTEL_float_controls2
 ; RUN: llvm-spirv %t.spv -o %t.spt --to-text
 ; RUN: llvm-spirv -r %t.spv -o %t.bc
 ; RUN: llvm-dis %t.bc -o %t.ll
-; RUN: FileCheck %s --input-file %t.spt -check-prefix=SPV
+; RUN: FileCheck %s --input-file %t.spt -check-prefixes=SPV,SPVEXT
 ; RUN: FileCheck %s --input-file %t.ll  -check-prefix=LLVM
 
+; RUN: llvm-spirv %t.bc -o %t.spv --spirv-max-version=1.4 --spirv-ext=+SPV_INTEL_vector_compute,+SPV_INTEL_float_controls2
+; RUN: llvm-spirv %t.spv -o %t.spt --to-text
+; RUN: FileCheck %s --input-file %t.spt -check-prefixes=SPV,SPV14
+
 
 ; ModuleID = 'float_control.bc'
 source_filename = "float_control.cpp"
@@ -13,7 +17,8 @@ target datalayout = "e-p:64:64-i64:64-n8:16:32"
 target triple = "spir"
 
 
-; SPV-DAG: Extension "SPV_KHR_float_controls"
+; SPVEXT-DAG: Extension "SPV_KHR_float_controls"
+; SPV14-NOT: Extension "SPV_KHR_float_controls"
 ; SPV-DAG: Extension "SPV_INTEL_float_controls2"
 
 ; LLVM-DAG: @k_rte{{[^a-zA-Z0-9_][^#]*}}#[[K_RTE:[0-9]+]]
diff --git a/llvm-spirv/test/transcoding/annotate_attribute.ll b/llvm-spirv/test/transcoding/annotate_attribute.ll
index c9055c83f6501..2bdc6c6d04502 100644
--- a/llvm-spirv/test/transcoding/annotate_attribute.ll
+++ b/llvm-spirv/test/transcoding/annotate_attribute.ll
@@ -10,6 +10,9 @@
 ; RUN: llvm-as %s -o %t.bc
 ; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_INTEL_fpga_memory_accesses,+SPV_INTEL_fpga_memory_attributes -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
 
+; Check SPIR-V versions in a format magic number + version
+; CHECK-SPIRV: 119734787 66560
+
 ; CHECK-SPIRV-DAG: Decorate {{[0-9]+}} UserSemantic "42"
 ; CHECK-SPIRV-DAG: Decorate {{[0-9]+}} UserSemantic "bar"
 ; CHECK-SPIRV-DAG: Decorate {{[0-9]+}} UserSemantic "{FOO}"
diff --git a/llvm-spirv/test/transcoding/block_w_struct_return.cl b/llvm-spirv/test/transcoding/block_w_struct_return.cl
index d23aeb331a330..dbc478b54195b 100644
--- a/llvm-spirv/test/transcoding/block_w_struct_return.cl
+++ b/llvm-spirv/test/transcoding/block_w_struct_return.cl
@@ -1,7 +1,17 @@
 // RUN: %clang_cc1 -triple spir -cl-std=cl2.0 -disable-llvm-passes -fdeclare-opencl-builtins -finclude-default-header %s -emit-llvm-bc -o %t.bc
-// RUN: llvm-spirv %t.bc -spirv-text -o %t.spv.txt
+// TODO: currently max version is limited to 1.1 for this test. Issues here
+// that the SPIR-V module generated for blocks is invalid for versions starting
+// from 1.4, spirv-val is failing with:
+//   error: line 63: Interface variable id <13> is used by entry point
+//                   'block_kernel' id <24>, but is not listed as an interface
+//   %__block_literal_global = OpVariable %_ptr_CrossWorkgroup__struct_10
+//                         CrossWorkgroup %11
+// details can be found in:
+// – Public issue #35: OpEntryPoint must list all global variables in the
+//   interface. Additionally, duplication in the list is not allowed.
+// RUN: llvm-spirv --spirv-max-version=1.1 %t.bc -spirv-text -o %t.spv.txt
 // RUN: FileCheck < %t.spv.txt %s --check-prefix=CHECK-SPIRV
-// RUN: llvm-spirv %t.bc -o %t.spv
+// RUN: llvm-spirv --spirv-max-version=1.1 %t.bc -o %t.spv
 // RUN: spirv-val %t.spv
 // RUN: llvm-spirv -r %t.spv -o %t.rev.bc
 // RUN: llvm-dis %t.rev.bc
diff --git a/llvm-spirv/test/transcoding/global_block.cl b/llvm-spirv/test/transcoding/global_block.cl
index 4a8be9528ece4..25e0e8a0ae8e6 100644
--- a/llvm-spirv/test/transcoding/global_block.cl
+++ b/llvm-spirv/test/transcoding/global_block.cl
@@ -4,8 +4,18 @@
 // removed
 
 // RUN: %clang_cc1 -O0 -triple spir-unknown-unknown -cl-std=CL2.0 -x cl %s -emit-llvm-bc -o %t.bc
-// RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-// RUN: llvm-spirv %t.bc -o %t.spv
+// TODO: currently max version is limited to 1.1 for this test. Issues here
+// that the SPIR-V module generated for blocks is invalid for versions starting
+// from 1.4, spirv-val is failing with:
+//   error: line 63: Interface variable id <13> is used by entry point
+//                   'block_kernel' id <24>, but is not listed as an interface
+//   %__block_literal_global = OpVariable %_ptr_CrossWorkgroup__struct_10
+//                         CrossWorkgroup %11
+// details can be found in:
+// – Public issue #35: OpEntryPoint must list all global variables in the
+//   interface. Additionally, duplication in the list is not allowed.
+// RUN: llvm-spirv --spirv-max-version=1.1 %t.bc -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+// RUN: llvm-spirv --spirv-max-version=1.1 %t.bc -o %t.spv
 // RUN: spirv-val %t.spv
 // RUN: llvm-spirv -r %t.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM
 

From 014151f28469dbf7344630c19b4642606b2d7c38 Mon Sep 17 00:00:00 2001
From: KornevNikita <nikita.kornev@intel.com>
Date: Tue, 21 Dec 2021 16:07:34 +0300
Subject: [PATCH 2/3] Fix the collection of entry point interfaces

This is a patch to expand the collection of entry point interfaces.
In SPIR-V 1.4 and later OpEntryPoint must list all global variables in the
interface. Also fix quoted string output in SPIRV text format.

Co-authored-by: Alexey Sotkin <alexey.sotkin@intel.com>

Original commit:
https://github.com/KhronosGroup/SPIRV-LLVM-Translator/commit/352ea14
---
 llvm-spirv/lib/SPIRV/SPIRVWriter.cpp          | 31 ++++++-----
 llvm-spirv/lib/SPIRV/SPIRVWriter.h            |  3 +-
 llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h |  1 -
 llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.cpp  |  6 ++-
 llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h    |  1 +
 llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.cpp | 38 +++++---------
 llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.h   |  7 ++-
 llvm-spirv/lib/SPIRV/libSPIRV/SPIRVStream.cpp |  1 +
 llvm-spirv/test/ExecutionMode.ll              |  3 --
 llvm-spirv/test/copy_object.spt               |  2 +-
 llvm-spirv/test/entry-point-interfaces.ll     | 52 +++++++++++++++++++
 llvm-spirv/test/negative/unimplemented.spt    |  2 +-
 llvm-spirv/test/right_shift.spt               |  2 +-
 .../inline_asm_clobbers.cl                    |  8 +--
 .../inline_asm_constraints.cl                 | 14 ++---
 .../test/transcoding/block_w_struct_return.cl | 34 ++++++------
 llvm-spirv/test/transcoding/global_block.cl   | 24 ++++-----
 17 files changed, 138 insertions(+), 91 deletions(-)
 create mode 100644 llvm-spirv/test/entry-point-interfaces.ll

diff --git a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp
index cf317502f662c..85810cffd694e 100644
--- a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp
+++ b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp
@@ -610,9 +610,7 @@ SPIRVFunction *LLVMToSPIRVBase::transFunctionDecl(Function *F) {
   BF->setFunctionControlMask(transFunctionControlMask(F));
   if (F->hasName())
     BM->setName(BF, F->getName().str());
-  if (isKernel(F))
-    BM->addEntryPoint(ExecutionModelKernel, BF->getId());
-  else if (F->getLinkage() != GlobalValue::InternalLinkage)
+  if (!isKernel(F) && F->getLinkage() != GlobalValue::InternalLinkage)
     BF->setLinkageType(transLinkageType(F));
 
   // Translate OpenCL/SYCL buffer_location metadata if it's attached to the
@@ -3581,12 +3579,15 @@ bool LLVMToSPIRVBase::isAnyFunctionReachableFromFunction(
   return false;
 }
 
-void LLVMToSPIRVBase::collectInputOutputVariables(SPIRVFunction *SF,
-                                                  Function *F) {
+std::vector<SPIRVId>
+LLVMToSPIRVBase::collectEntryPointInterfaces(SPIRVFunction *SF, Function *F) {
+  std::vector<SPIRVId> Interface;
   for (auto &GV : M->globals()) {
     const auto AS = GV.getAddressSpace();
-    if (AS != SPIRAS_Input && AS != SPIRAS_Output)
-      continue;
+    SPIRVModule *BM = SF->getModule();
+    if (!BM->isAllowedToUseVersion(VersionNumber::SPIRV_1_4))
+      if (AS != SPIRAS_Input && AS != SPIRAS_Output)
+        continue;
 
     std::unordered_set<const Function *> Funcs;
 
@@ -3598,9 +3599,15 @@ void LLVMToSPIRVBase::collectInputOutputVariables(SPIRVFunction *SF,
     }
 
     if (isAnyFunctionReachableFromFunction(F, Funcs)) {
-      SF->addVariable(ValueMap[&GV]);
+      SPIRVWord ModuleVersion = static_cast<SPIRVWord>(BM->getSPIRVVersion());
+      if (AS != SPIRAS_Input && AS != SPIRAS_Output &&
+          ModuleVersion < static_cast<SPIRVWord>(VersionNumber::SPIRV_1_4))
+        BM->setMinSPIRVVersion(
+            static_cast<SPIRVWord>(VersionNumber::SPIRV_1_4));
+      Interface.push_back(ValueMap[&GV]->getId());
     }
   }
+  return Interface;
 }
 
 void LLVMToSPIRVBase::mutateFuncArgType(
@@ -3703,10 +3710,10 @@ void LLVMToSPIRVBase::transFunction(Function *I) {
   joinFPContract(I, FPContract::ENABLED);
   fpContractUpdateRecursive(I, getFPContract(I));
 
-  bool IsKernelEntryPoint = isKernel(I);
-
-  if (IsKernelEntryPoint) {
-    collectInputOutputVariables(BF, I);
+  if (isKernel(I)) {
+    auto Interface = collectEntryPointInterfaces(BF, I);
+    BM->addEntryPoint(ExecutionModelKernel, BF->getId(), I->getName().str(),
+                      Interface);
   }
 }
 
diff --git a/llvm-spirv/lib/SPIRV/SPIRVWriter.h b/llvm-spirv/lib/SPIRV/SPIRVWriter.h
index 78ba39e36c261..ea128709cf081 100644
--- a/llvm-spirv/lib/SPIRV/SPIRVWriter.h
+++ b/llvm-spirv/lib/SPIRV/SPIRVWriter.h
@@ -216,7 +216,8 @@ class LLVMToSPIRVBase {
   bool isAnyFunctionReachableFromFunction(
       const Function *FS,
       const std::unordered_set<const Function *> Funcs) const;
-  void collectInputOutputVariables(SPIRVFunction *SF, Function *F);
+  std::vector<SPIRVId> collectEntryPointInterfaces(SPIRVFunction *BF,
+                                                   Function *F);
 };
 
 class LLVMToSPIRVPass : public PassInfoMixin<LLVMToSPIRVPass>,
diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h
index b5a91474cee1f..e499ee2acbc33 100644
--- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h
+++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVDecorate.h
@@ -241,7 +241,6 @@ class SPIRVDecorateLinkageAttr : public SPIRVDecorate {
 #ifdef _SPIRV_SUPPORT_TEXT_FMT
     if (SPIRVUseTextFormat) {
       Encoder << getString(Literals.cbegin(), Literals.cend() - 1);
-      Encoder.OS << " ";
       Encoder << (SPIRVLinkageTypeKind)Literals.back();
     } else
 #endif
diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.cpp b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.cpp
index 443b5a8a40f5d..ad63f8fa1b925 100644
--- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.cpp
+++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.cpp
@@ -541,9 +541,11 @@ void SPIRVEntryPoint::encode(spv_ostream &O) const {
 }
 
 void SPIRVEntryPoint::decode(std::istream &I) {
-  getDecoder(I) >> ExecModel >> Target >> Name >> Variables;
+  getDecoder(I) >> ExecModel >> Target >> Name;
+  Variables.resize(WordCount - FixedWC - getSizeInWords(Name) + 1);
+  getDecoder(I) >> Variables;
   Module->setName(getOrCreateTarget(), Name);
-  Module->addEntryPoint(ExecModel, Target);
+  Module->addEntryPoint(ExecModel, Target, Name, Variables);
 }
 
 void SPIRVExecutionMode::encode(spv_ostream &O) const {
diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h
index 2c6f5cacc2c90..975b9c2044ce2 100644
--- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h
+++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVEntry.h
@@ -524,6 +524,7 @@ template <Op OC> class SPIRVAnnotation : public SPIRVAnnotationGeneric {
 
 class SPIRVEntryPoint : public SPIRVAnnotation<OpEntryPoint> {
 public:
+  static const SPIRVWord FixedWC = 4;
   SPIRVEntryPoint(SPIRVModule *TheModule, SPIRVExecutionModelKind,
                   SPIRVId TheId, const std::string &TheName,
                   std::vector<SPIRVId> Variables);
diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.cpp b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.cpp
index 6d8c1a9a5aabf..95112c2add449 100644
--- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.cpp
+++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.cpp
@@ -128,20 +128,6 @@ class SPIRVModuleImpl : public SPIRVModule {
   getValueTypes(const std::vector<SPIRVId> &) const override;
   SPIRVMemoryModelKind getMemoryModel() const override { return MemoryModel; }
   SPIRVConstant *getLiteralAsConstant(unsigned Literal) override;
-  unsigned getNumEntryPoints(SPIRVExecutionModelKind EM) const override {
-    auto Loc = EntryPointVec.find(EM);
-    if (Loc == EntryPointVec.end())
-      return 0;
-    return Loc->second.size();
-  }
-  SPIRVFunction *getEntryPoint(SPIRVExecutionModelKind EM,
-                               unsigned I) const override {
-    auto Loc = EntryPointVec.find(EM);
-    if (Loc == EntryPointVec.end())
-      return nullptr;
-    assert(I < Loc->second.size());
-    return get<SPIRVFunction>(Loc->second[I]);
-  }
   unsigned getNumFunctions() const override { return FuncVec.size(); }
   unsigned getNumVariables() const override { return VariableVec.size(); }
   SourceLanguage getSourceLanguage(SPIRVWord *Ver = nullptr) const override {
@@ -215,8 +201,9 @@ class SPIRVModuleImpl : public SPIRVModule {
   SPIRVGroupMemberDecorate *
   addGroupMemberDecorate(SPIRVDecorationGroup *Group,
                          const std::vector<SPIRVEntry *> &Targets) override;
-  void addEntryPoint(SPIRVExecutionModelKind ExecModel,
-                     SPIRVId EntryPoint) override;
+  void addEntryPoint(SPIRVExecutionModelKind ExecModel, SPIRVId EntryPoint,
+                     const std::string &Name,
+                     const std::vector<SPIRVId> &Variables) override;
   SPIRVForward *addForward(SPIRVType *Ty) override;
   SPIRVForward *addForward(SPIRVId, SPIRVType *Ty) override;
   SPIRVFunction *addFunction(SPIRVFunction *) override;
@@ -495,11 +482,11 @@ class SPIRVModuleImpl : public SPIRVModule {
   typedef std::vector<SPIRVGroupDecorateGeneric *> SPIRVGroupDecVec;
   typedef std::vector<SPIRVAsmTargetINTEL *> SPIRVAsmTargetVector;
   typedef std::vector<SPIRVAsmINTEL *> SPIRVAsmVector;
+  typedef std::vector<SPIRVEntryPoint *> SPIRVEntryPointVec;
   typedef std::map<SPIRVId, SPIRVExtInstSetKind> SPIRVIdToInstructionSetMap;
   std::map<SPIRVExtInstSetKind, SPIRVId> ExtInstSetIds;
   typedef std::map<SPIRVId, SPIRVExtInstSetKind> SPIRVIdToBuiltinSetMap;
   typedef std::map<SPIRVExecutionModelKind, SPIRVIdSet> SPIRVExecModelIdSetMap;
-  typedef std::map<SPIRVExecutionModelKind, SPIRVIdVec> SPIRVExecModelIdVecMap;
   typedef std::unordered_map<std::string, SPIRVString *> SPIRVStringMap;
   typedef std::map<SPIRVTypeStruct *, std::vector<std::pair<unsigned, SPIRVId>>>
       SPIRVUnknownStructFieldMap;
@@ -525,7 +512,7 @@ class SPIRVModuleImpl : public SPIRVModule {
   SPIRVAsmTargetVector AsmTargetVec;
   SPIRVAsmVector AsmVec;
   SPIRVExecModelIdSetMap EntryPointSet;
-  SPIRVExecModelIdVecMap EntryPointVec;
+  SPIRVEntryPointVec EntryPointVec;
   SPIRVStringMap StrMap;
   SPIRVCapMap CapMap;
   SPIRVUnknownStructFieldMap UnknownStructFieldMap;
@@ -1000,11 +987,14 @@ SPIRVModuleImpl::addDecorate(SPIRVDecorateGeneric *Dec) {
 }
 
 void SPIRVModuleImpl::addEntryPoint(SPIRVExecutionModelKind ExecModel,
-                                    SPIRVId EntryPoint) {
+                                    SPIRVId EntryPoint, const std::string &Name,
+                                    const std::vector<SPIRVId> &Variables) {
   assert(isValid(ExecModel) && "Invalid execution model");
   assert(EntryPoint != SPIRVID_INVALID && "Invalid entry point");
+  auto *EP =
+      add(new SPIRVEntryPoint(this, ExecModel, EntryPoint, Name, Variables));
+  EntryPointVec.push_back(EP);
   EntryPointSet[ExecModel].insert(EntryPoint);
-  EntryPointVec[ExecModel].push_back(EntryPoint);
   addCapabilities(SPIRV::getCapability(ExecModel));
 }
 
@@ -1833,14 +1823,10 @@ spv_ostream &operator<<(spv_ostream &O, SPIRVModule &M) {
 
   O << SPIRVMemoryModel(&M);
 
-  for (auto &I : MI.EntryPointVec)
-    for (auto &II : I.second)
-      O << SPIRVEntryPoint(&M, I.first, II, M.get<SPIRVFunction>(II)->getName(),
-                           M.get<SPIRVFunction>(II)->getVariables());
+  O << MI.EntryPointVec;
 
   for (auto &I : MI.EntryPointVec)
-    for (auto &II : I.second)
-      MI.get<SPIRVFunction>(II)->encodeExecutionModes(O);
+    MI.get<SPIRVFunction>(I->getTargetId())->encodeExecutionModes(O);
 
   O << MI.StringVec;
 
diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.h b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.h
index 07bbfe0c60e05..f45eead3fee6e 100644
--- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.h
+++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVModule.h
@@ -133,14 +133,11 @@ class SPIRVModule {
   virtual const SPIRVCapMap &getCapability() const = 0;
   virtual bool hasCapability(SPIRVCapabilityKind) const = 0;
   virtual SPIRVExtInstSetKind getBuiltinSet(SPIRVId) const = 0;
-  virtual SPIRVFunction *getEntryPoint(SPIRVExecutionModelKind,
-                                       unsigned) const = 0;
   virtual std::set<std::string> &getExtension() = 0;
   virtual SPIRVFunction *getFunction(unsigned) const = 0;
   virtual SPIRVVariable *getVariable(unsigned) const = 0;
   virtual SPIRVMemoryModelKind getMemoryModel() const = 0;
   virtual unsigned getNumFunctions() const = 0;
-  virtual unsigned getNumEntryPoints(SPIRVExecutionModelKind) const = 0;
   virtual unsigned getNumVariables() const = 0;
   virtual SourceLanguage getSourceLanguage(SPIRVWord *) const = 0;
   virtual std::set<std::string> &getSourceExtension() = 0;
@@ -213,7 +210,9 @@ class SPIRVModule {
                          const std::vector<SPIRVEntry *> &Targets) = 0;
   virtual SPIRVGroupDecorateGeneric *
   addGroupDecorateGeneric(SPIRVGroupDecorateGeneric *GDec) = 0;
-  virtual void addEntryPoint(SPIRVExecutionModelKind, SPIRVId) = 0;
+  virtual void addEntryPoint(SPIRVExecutionModelKind, SPIRVId,
+                             const std::string &,
+                             const std::vector<SPIRVId> &) = 0;
   virtual SPIRVForward *addForward(SPIRVType *Ty) = 0;
   virtual SPIRVForward *addForward(SPIRVId, SPIRVType *Ty) = 0;
   virtual SPIRVFunction *addFunction(SPIRVFunction *) = 0;
diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVStream.cpp b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVStream.cpp
index 555b7a274f9f6..d7b43c50b81ea 100644
--- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVStream.cpp
+++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVStream.cpp
@@ -169,6 +169,7 @@ const SPIRVEncoder &operator<<(const SPIRVEncoder &O, const std::string &Str) {
 #ifdef _SPIRV_SUPPORT_TEXT_FMT
   if (SPIRVUseTextFormat) {
     writeQuotedString(O.OS, Str);
+    O.OS << " ";
     return O;
   }
 #endif
diff --git a/llvm-spirv/test/ExecutionMode.ll b/llvm-spirv/test/ExecutionMode.ll
index 9858342eb21c3..5f6aef54e69e2 100644
--- a/llvm-spirv/test/ExecutionMode.ll
+++ b/llvm-spirv/test/ExecutionMode.ll
@@ -1,9 +1,6 @@
 ; RUN: llvm-as < %s | llvm-spirv -spirv-text -o %t
 ; RUN: FileCheck < %t %s
 
-; check for magic number followed by version 1.1
-; CHECK: 119734787 65792
-
 ; CHECK-DAG: TypeVoid [[VOID:[0-9]+]]
 
 ; CHECK-DAG: EntryPoint 6 [[WORKER:[0-9]+]] "worker"
diff --git a/llvm-spirv/test/copy_object.spt b/llvm-spirv/test/copy_object.spt
index 3587a4c966173..44d621431c905 100644
--- a/llvm-spirv/test/copy_object.spt
+++ b/llvm-spirv/test/copy_object.spt
@@ -5,7 +5,7 @@
 2 Capability Int64
 2 Capability Int8
 3 MemoryModel 2 2
-8 EntryPoint 6 1 "copy_object"
+6 EntryPoint 6 1 "copy_object" 
 3 Source 3 102000
 3 Name 2 "in"
 4 Decorate 3 BuiltIn 28
diff --git a/llvm-spirv/test/entry-point-interfaces.ll b/llvm-spirv/test/entry-point-interfaces.ll
new file mode 100644
index 0000000000000..99934b397a72f
--- /dev/null
+++ b/llvm-spirv/test/entry-point-interfaces.ll
@@ -0,0 +1,52 @@
+; RUN: llvm-as %s -o %t.bc
+
+; RUN: llvm-spirv %t.bc -o %t.spv
+; RUN: spirv-val --target-env spv1.4 %t.spv
+; RUN: llvm-spirv -to-text %t.spv -o %t.from.spv.spt
+; RUN: FileCheck < %t.from.spv.spt %s --check-prefix=CHECK-SPIRV
+
+; RUN: llvm-spirv -spirv-text %t.bc -o %t.from.bc.spt
+; RUN: FileCheck < %t.from.bc.spt %s --check-prefix=CHECK-SPIRV
+
+; CHECK-SPIRV: 7 EntryPoint 6 [[#]] "test" [[#Interface1:]] [[#Interface2:]]
+; CHECK-SPIRV: TypeInt [[#TypeInt:]] 32 0
+; CHECK-SPIRV: Constant [[#TypeInt]] [[#Constant1:]] 1
+; CHECK-SPIRV: Constant [[#TypeInt]] [[#Constant2:]] 3
+; CHECK-SPIRV: Variable [[#]] [[#Interface1]] 0 [[#Constant1]]
+; CHECK-SPIRV: Variable [[#]] [[#Interface2]] 0 [[#Constant2]]
+
+; ModuleID = 'source.cpp'
+source_filename = "source.cpp"
+target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64"
+target triple = "spir"
+
+@var = dso_local addrspace(2) constant i32 1, align 4
+@var2 = dso_local addrspace(2) constant i32 3, align 4
+@var.const = private unnamed_addr addrspace(2) constant i32 1, align 4
+@var2.const = private unnamed_addr addrspace(2) constant i32 3, align 4
+
+; Function Attrs: convergent noinline norecurse nounwind optnone
+define dso_local spir_kernel void @test() #0 !kernel_arg_addr_space !2 !kernel_arg_access_qual !2 !kernel_arg_type !2 !kernel_arg_base_type !2 !kernel_arg_type_qual !2 !kernel_arg_host_accessible !2 !kernel_arg_pipe_depth !2 !kernel_arg_pipe_io !2 !kernel_arg_buffer_location !2 {
+entry:
+  %0 = load i32, i32 addrspace(2)* @var.const, align 4
+  %1 = load i32, i32 addrspace(2)* @var2.const, align 4
+  %mul = mul nsw i32 %0, %1
+  %mul1 = mul nsw i32 %mul, 2
+  ret void
+}
+
+attributes #0 = { convergent noinline norecurse nounwind optnone "frame-pointer"="all" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="false" }
+
+!opencl.enable.FP_CONTRACT = !{}
+!opencl.ocl.version = !{!0}
+!opencl.spir.version = !{!0}
+!llvm.module.flags = !{!1}
+!opencl.used.extensions = !{!2}
+!opencl.used.optional.core.features = !{!2}
+!opencl.compiler.options = !{!2}
+!llvm.ident = !{!3}
+
+!0 = !{i32 2, i32 0}
+!1 = !{i32 7, !"frame-pointer", i32 2}
+!2 = !{}
+!3 = !{!"Compiler"}
diff --git a/llvm-spirv/test/negative/unimplemented.spt b/llvm-spirv/test/negative/unimplemented.spt
index fab5fe6105159..0f902155fcf39 100644
--- a/llvm-spirv/test/negative/unimplemented.spt
+++ b/llvm-spirv/test/negative/unimplemented.spt
@@ -2,7 +2,7 @@
 2 Capability Addresses
 2 Capability Shader
 3 MemoryModel 2 2
-6 EntryPoint 6 2 "foo"
+4 EntryPoint 6 2 "foo"
 3 Name 3 "res"
 2 TypeVoid 12
 3 TypeFloat 13 32
diff --git a/llvm-spirv/test/right_shift.spt b/llvm-spirv/test/right_shift.spt
index f31dc10b24b3e..a0bac4b69b584 100644
--- a/llvm-spirv/test/right_shift.spt
+++ b/llvm-spirv/test/right_shift.spt
@@ -4,7 +4,7 @@
 2 Capability Kernel
 2 Capability Int64
 3 MemoryModel 2 2
-10 EntryPoint 6 1 "shift_right_arithmetic"
+9 EntryPoint 6 1 "shift_right_arithmetic" 
 3 Source 3 102000
 3 Name 2 "in"
 4 Decorate 3 BuiltIn 28
diff --git a/llvm-spirv/test/transcoding/SPV_INTEL_inline_assembly/inline_asm_clobbers.cl b/llvm-spirv/test/transcoding/SPV_INTEL_inline_assembly/inline_asm_clobbers.cl
index a640430ec9f50..3f9237495b76e 100644
--- a/llvm-spirv/test/transcoding/SPV_INTEL_inline_assembly/inline_asm_clobbers.cl
+++ b/llvm-spirv/test/transcoding/SPV_INTEL_inline_assembly/inline_asm_clobbers.cl
@@ -20,7 +20,7 @@ size_t __ovld __cnfn get_global_id(unsigned int dimindx);
 // XCHECK-LLVM: [[STRUCTYPE:%[a-z0-9]+]] = type { i32, i32 }
 
 // CHECK-LLVM-LABEL: define spir_kernel void @mem_clobber
-// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} """~{cc},~{memory}"
+// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "" "~{cc},~{memory}"
 // CHECK-LLVM: [[VALUE:%[0-9]+]] = load i32 addrspace(1)*, i32 addrspace(1)**
 // CHECK-LLVM-NEXT: getelementptr inbounds i32, i32 addrspace(1)* [[VALUE]], i64 0
 // CHECK-LLVM-NEXT: store i32 1, i32 addrspace(1)*
@@ -34,7 +34,7 @@ kernel void mem_clobber(global int *x) {
 }
 
 // CHECK-LLVM-LABEL: define spir_kernel void @out_clobber
-// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "earlyclobber_instruction_out $0""=&r"
+// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "earlyclobber_instruction_out $0" "=&r"
 // CHECK-LLVM: barrier
 // CHECK-LLVM: store i32 %{{[a-z0-9]+}}, i32* [[VALUE:%[a-z0-9]+]], align 4
 // CHECK-LLVM-NEXT: [[STOREVAL:%[a-z0-9]+]] = call i32 asm "earlyclobber_instruction_out $0", "=&r"()
@@ -54,7 +54,7 @@ kernel void out_clobber(global int *x) {
 //       Or bug in clang FE. To investigate later, change xchecks to checks and enable
 
 // XCHECK-LLVM-LABEL: define spir_kernel void @in_clobber
-// XCHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "earlyclobber_instruction_in $0""&r"
+// XCHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "earlyclobber_instruction_in $0" "&r"
 // XCHECK-LLVM: barrier
 // XCHECK-LLVM: getelementptr
 // XCHECK-LLVM: store i32  %{{[a-z0-9]+}}, i32* [[LOADVAL:%[a-z0-9]+]], align 4
@@ -74,7 +74,7 @@ kernel void in_clobber(global int *x) {
 #endif
 
 // XCHECK-LLVM-LABEL: define spir_kernel void @mixed_clobber
-// XCHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixedclobber_instruction $0 $1 $2""=&r,=&r,&r,1,~{cc},~{memory}"
+// XCHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixedclobber_instruction $0 $1 $2" "=&r,=&r,&r,1,~{cc},~{memory}"
 
 #if 0
 kernel void mixed_clobber(global int *x, global int *y, global int *z) {
diff --git a/llvm-spirv/test/transcoding/SPV_INTEL_inline_assembly/inline_asm_constraints.cl b/llvm-spirv/test/transcoding/SPV_INTEL_inline_assembly/inline_asm_constraints.cl
index 98bf490642550..d4417f7d1ee12 100644
--- a/llvm-spirv/test/transcoding/SPV_INTEL_inline_assembly/inline_asm_constraints.cl
+++ b/llvm-spirv/test/transcoding/SPV_INTEL_inline_assembly/inline_asm_constraints.cl
@@ -24,7 +24,7 @@ size_t __ovld __cnfn get_global_id(unsigned int dimindx);
 // CHECK-LLVM: [[STRUCTYPE:%[a-z]+]] = type { i32, i8, float }
 
 // CHECK-LLVM-LABEL: define spir_kernel void @test_int
-// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "intcommand $0 $1""=r,r"
+// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "intcommand $0 $1" "=r,r"
 // CHECK-LLVM: [[VALUE:%[0-9]+]] = call i32 asm sideeffect "intcommand $0 $1", "=r,r"(i32 %{{[0-9]+}})
 // CHECK-LLVM-NEXT: store i32 [[VALUE]], i32 addrspace(1)*
 
@@ -34,7 +34,7 @@ kernel void test_int(global int *in, global int *out) {
 }
 
 // CHECK-LLVM-LABEL: define spir_kernel void @test_float
-// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "floatcommand $0 $1""=r,r"
+// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "floatcommand $0 $1" "=r,r"
 // CHECK-LLVM: [[VALUE:%[0-9]+]] = call float asm sideeffect "floatcommand $0 $1", "=r,r"(float %{{[0-9]+}})
 // CHECK-LLVM-NEXT: store float [[VALUE]], float addrspace(1)*
 
@@ -44,7 +44,7 @@ kernel void test_float(global float *in, global float *out) {
 }
 
 // CHECK-LLVM-LABEL: define spir_kernel void @test_mixed_integral
-// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_integral_command $0 $3 $1 $2""=r,r,r,r"
+// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_integral_command $0 $3 $1 $2" "=r,r,r,r"
 // CHECK-LLVM: [[VALUE:%[0-9]+]] = call i64 asm sideeffect "mixed_integral_command $0 $3 $1 $2", "=r,r,r,r"(i16 %{{[0-9]+}}, i32 %{{[0-9]+}}, i8 %{{[0-9]+}})
 // CHECK-LLVM-NEXT: store i64 [[VALUE]], i64 addrspace(1)*
 
@@ -55,7 +55,7 @@ kernel void test_mixed_integral(global uchar *A, global ushort *B, global uint *
 }
 
 // CHECK-LLVM-LABEL: define spir_kernel void @test_mixed_floating
-// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_floating_command $0 $1 $2""=r,r,r"
+// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_floating_command $0 $1 $2" "=r,r,r"
 // CHECK-LLVM: [[VALUE:%[0-9]+]] = call half asm sideeffect "mixed_floating_command $0 $1 $2", "=r,r,r"(double %{{[0-9]+}}, float %{{[0-9]+}})
 // CHECK-LLVM-NEXT: store half [[VALUE]], half addrspace(1)*
 
@@ -66,7 +66,7 @@ kernel void test_mixed_floating(global float *A, global half *B, global double *
 }
 
 // CHECK-LLVM-LABEL: define spir_kernel void @test_mixed_all
-// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_all_command $0 $3 $1 $2""=r,r,r,r"
+// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "mixed_all_command $0 $3 $1 $2" "=r,r,r,r"
 // CHECK-LLVM: [[VALUE:%[0-9]+]] = call i8 asm sideeffect "mixed_all_command $0 $3 $1 $2", "=r,r,r,r"(float %{{[0-9]+}}, i32 %{{[0-9]+}}, i8 %{{[0-9]+}})
 // CHECK-LLVM-NEXT: store i8 [[VALUE]], i8 addrspace(1)*
 
@@ -77,7 +77,7 @@ kernel void test_mixed_all(global uchar *A, global float *B, global uint *C, glo
 }
 
 // CHECK-LLVM-LABEL: define spir_kernel void @test_multiple
-// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "multiple_command $0 $0 $1 $1 $2 $2""=r,=r,=r,0,1,2"
+// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "multiple_command $0 $0 $1 $1 $2 $2" "=r,=r,=r,0,1,2"
 // CHECK-LLVM: [[VALUE:%[0-9]+]] = call [[STRUCTYPE]] asm sideeffect "multiple_command $0 $0 $1 $1 $2 $2", "=r,=r,=r,0,1,2"(i32 %{{[0-9]+}}, i8 %{{[0-9]+}}, float %{{[0-9]+}})
 // CHECK-LLVM-NEXT: extractvalue [[STRUCTYPE]] [[VALUE]], 0
 // CHECK-LLVM-NEXT: extractvalue [[STRUCTYPE]] [[VALUE]], 1
@@ -90,7 +90,7 @@ kernel void test_multiple(global uchar *A, global float *B, global uint *C) {
 }
 
 // CHECK-LLVM-LABEL: define spir_kernel void @test_constants
-// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "constcommand $0 $1""i,i"
+// CHECK-SPIRV: {{[0-9]+}} AsmINTEL {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} {{[0-9]+}} "constcommand $0 $1" "i,i"
 // CHECK-LLVM: call void asm sideeffect "constcommand $0 $1", "i,i"(i32 1, double 2.000000e+00)
 
 kernel void test_constants() {
diff --git a/llvm-spirv/test/transcoding/block_w_struct_return.cl b/llvm-spirv/test/transcoding/block_w_struct_return.cl
index dbc478b54195b..638ae8ffcf5b1 100644
--- a/llvm-spirv/test/transcoding/block_w_struct_return.cl
+++ b/llvm-spirv/test/transcoding/block_w_struct_return.cl
@@ -1,19 +1,16 @@
 // RUN: %clang_cc1 -triple spir -cl-std=cl2.0 -disable-llvm-passes -fdeclare-opencl-builtins -finclude-default-header %s -emit-llvm-bc -o %t.bc
-// TODO: currently max version is limited to 1.1 for this test. Issues here
-// that the SPIR-V module generated for blocks is invalid for versions starting
-// from 1.4, spirv-val is failing with:
-//   error: line 63: Interface variable id <13> is used by entry point
-//                   'block_kernel' id <24>, but is not listed as an interface
-//   %__block_literal_global = OpVariable %_ptr_CrossWorkgroup__struct_10
-//                         CrossWorkgroup %11
-// details can be found in:
-// – Public issue #35: OpEntryPoint must list all global variables in the
-//   interface. Additionally, duplication in the list is not allowed.
-// RUN: llvm-spirv --spirv-max-version=1.1 %t.bc -spirv-text -o %t.spv.txt
-// RUN: FileCheck < %t.spv.txt %s --check-prefix=CHECK-SPIRV
-// RUN: llvm-spirv --spirv-max-version=1.1 %t.bc -o %t.spv
-// RUN: spirv-val %t.spv
-// RUN: llvm-spirv -r %t.spv -o %t.rev.bc
+
+// RUN: llvm-spirv --spirv-max-version=1.1 %t.bc -spirv-text -o - | FileCheck %s --check-prefixes=CHECK-SPIRV1_1,CHECK-SPIRV
+// RUN: llvm-spirv --spirv-max-version=1.1 %t.bc -o %t.spirv1.1.spv
+// RUN: spirv-val --target-env spv1.1 %t.spirv1.1.spv
+// RUN: llvm-spirv -r %t.spirv1.1.spv -o %t.rev.bc
+// RUN: llvm-dis %t.rev.bc
+// RUN: FileCheck < %t.rev.ll %s --check-prefix=CHECK-LLVM
+
+// RUN: llvm-spirv --spirv-max-version=1.4 %t.bc -spirv-text -o - | FileCheck %s --check-prefixes=CHECK-SPIRV1_4,CHECK-SPIRV
+// RUN: llvm-spirv --spirv-max-version=1.4 %t.bc -o %t.spirv1.4.spv
+// RUN: spirv-val --target-env spv1.4 %t.spirv1.4.spv
+// RUN: llvm-spirv -r %t.spirv1.4.spv -o %t.rev.bc
 // RUN: llvm-dis %t.rev.bc
 // RUN: FileCheck < %t.rev.ll %s --check-prefix=CHECK-LLVM
 
@@ -34,6 +31,13 @@ kernel void block_ret_struct(__global int* res)
   res[tid] = kernelBlock(aa).a - 6;
 }
 
+// CHECK-SPIRV1_4: EntryPoint 6 [[#]] "block_ret_struct" [[#InterdaceId1:]] [[#InterdaceId2:]]
+// CHECK-SPIRV1_4: Name [[#InterdaceId1]] "__block_literal_global"
+// CHECK-SPIRV1_4: Name [[#InterdaceId2]] "__spirv_BuiltInGlobalInvocationId"
+
+// CHECK-SPIRV1_1: EntryPoint 6 [[#]] "block_ret_struct" [[#InterdaceId1:]]
+// CHECK-SPIRV1_1: Name [[#InterdaceId1]] "__spirv_BuiltInGlobalInvocationId"
+
 // CHECK-SPIRV: Name [[BlockInv:[0-9]+]] "__block_ret_struct_block_invoke"
 
 // CHECK-SPIRV: 4 TypeInt [[IntTy:[0-9]+]] 32
diff --git a/llvm-spirv/test/transcoding/global_block.cl b/llvm-spirv/test/transcoding/global_block.cl
index 25e0e8a0ae8e6..50cb4cb09da39 100644
--- a/llvm-spirv/test/transcoding/global_block.cl
+++ b/llvm-spirv/test/transcoding/global_block.cl
@@ -4,20 +4,16 @@
 // removed
 
 // RUN: %clang_cc1 -O0 -triple spir-unknown-unknown -cl-std=CL2.0 -x cl %s -emit-llvm-bc -o %t.bc
-// TODO: currently max version is limited to 1.1 for this test. Issues here
-// that the SPIR-V module generated for blocks is invalid for versions starting
-// from 1.4, spirv-val is failing with:
-//   error: line 63: Interface variable id <13> is used by entry point
-//                   'block_kernel' id <24>, but is not listed as an interface
-//   %__block_literal_global = OpVariable %_ptr_CrossWorkgroup__struct_10
-//                         CrossWorkgroup %11
-// details can be found in:
-// – Public issue #35: OpEntryPoint must list all global variables in the
-//   interface. Additionally, duplication in the list is not allowed.
+
 // RUN: llvm-spirv --spirv-max-version=1.1 %t.bc -spirv-text -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-// RUN: llvm-spirv --spirv-max-version=1.1 %t.bc -o %t.spv
-// RUN: spirv-val %t.spv
-// RUN: llvm-spirv -r %t.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM
+// RUN: llvm-spirv --spirv-max-version=1.1 %t.bc -o %t.spirv1.1.spv
+// RUN: spirv-val --target-env spv1.1 %t.spirv1.1.spv
+// RUN: llvm-spirv -r %t.spirv1.1.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM
+
+// RUN: llvm-spirv --spirv-max-version=1.4 %t.bc -spirv-text -o - | FileCheck %s --check-prefixes=CHECK-SPIRV1_4,CHECK-SPIRV
+// RUN: llvm-spirv --spirv-max-version=1.4 %t.bc -o %t.spirv1.4.spv
+// RUN: spirv-val --target-env spv1.4 %t.spirv1.4.spv
+// RUN: llvm-spirv -r %t.spirv1.4.spv -o - | llvm-dis | FileCheck %s --check-prefix=CHECK-LLVM
 
 kernel void block_kernel(__global int* res) {
   typedef int (^block_t)(int);
@@ -25,6 +21,8 @@ kernel void block_kernel(__global int* res) {
   *res = b1(5);
 }
 
+// CHECK-SPIRV1_4: EntryPoint 6 [[#]] "block_kernel" [[#InterfaceId:]]
+// CHECK-SPIRV1_4: Name [[#InterfaceId]] "__block_literal_global"
 // CHECK-SPIRV: Name [[block_invoke:[0-9]+]] "_block_invoke"
 // CHECK-SPIRV: TypeInt [[int:[0-9]+]] 32
 // CHECK-SPIRV: TypeInt [[int8:[0-9]+]] 8

From ce8b30cf6d95324726d004ea4d9bc066634c578f Mon Sep 17 00:00:00 2001
From: "Maksimova, Viktoria" <viktoria.maksimova@intel.com>
Date: Fri, 2 Sep 2022 10:05:56 -0700
Subject: [PATCH 3/3] Fix build error

---
 llvm-spirv/lib/SPIRV/SPIRVWriter.cpp         | 9 +++------
 llvm-spirv/lib/SPIRV/libSPIRV/SPIRVValue.cpp | 3 +--
 2 files changed, 4 insertions(+), 8 deletions(-)

diff --git a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp
index a07c1325cfc4a..3714c73773c23 100644
--- a/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp
+++ b/llvm-spirv/lib/SPIRV/SPIRVWriter.cpp
@@ -1393,8 +1393,7 @@ LLVMToSPIRVBase::getLoopControl(const BranchInst *Branch,
       else if (S == "llvm.loop.unroll.count" &&
                !(LoopControl & LoopControlDontUnrollMask)) {
         if (BM->isAllowedToUseVersion(VersionNumber::SPIRV_1_4)) {
-          BM->setMinSPIRVVersion(
-              static_cast<SPIRVWord>(VersionNumber::SPIRV_1_4));
+          BM->setMinSPIRVVersion(VersionNumber::SPIRV_1_4);
           size_t I = getMDOperandAsInt(Node, 1);
           ParametersToSort.emplace_back(spv::LoopControlPartialCountMask, I);
           LoopControl |= spv::LoopControlPartialCountMask;
@@ -4225,8 +4224,7 @@ LLVMToSPIRVBase::collectEntryPointInterfaces(SPIRVFunction *SF, Function *F) {
       SPIRVWord ModuleVersion = static_cast<SPIRVWord>(BM->getSPIRVVersion());
       if (AS != SPIRAS_Input && AS != SPIRAS_Output &&
           ModuleVersion < static_cast<SPIRVWord>(VersionNumber::SPIRV_1_4))
-        BM->setMinSPIRVVersion(
-            static_cast<SPIRVWord>(VersionNumber::SPIRV_1_4));
+        BM->setMinSPIRVVersion(VersionNumber::SPIRV_1_4);
       Interface.push_back(ValueMap[&GV]->getId());
     }
   }
@@ -4581,8 +4579,7 @@ bool LLVMToSPIRVBase::transExecutionMode() {
       case spv::ExecutionModeRoundingModeRTE:
       case spv::ExecutionModeRoundingModeRTZ: {
         if (BM->isAllowedToUseVersion(VersionNumber::SPIRV_1_4)) {
-          BM->setMinSPIRVVersion(
-              static_cast<SPIRVWord>(VersionNumber::SPIRV_1_4));
+          BM->setMinSPIRVVersion(VersionNumber::SPIRV_1_4);
           AddSingleArgExecutionMode(static_cast<ExecutionMode>(EMode));
         } else if (BM->isAllowedToUseExtension(
                        ExtensionID::SPV_KHR_float_controls)) {
diff --git a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVValue.cpp b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVValue.cpp
index 2356c05ddb39e..6879b2e2552ae 100644
--- a/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVValue.cpp
+++ b/llvm-spirv/lib/SPIRV/libSPIRV/SPIRVValue.cpp
@@ -103,8 +103,7 @@ void SPIRVValue::setNoIntegerDecorationWrap(bool HasNoIntegerWrap) {
       NoIntegerWrapDecoration == DecorationNoSignedWrap ? "nsw" : "nuw";
 #endif // _SPIRVDBG
   if (Module->isAllowedToUseVersion(VersionNumber::SPIRV_1_4)) {
-    Module->setMinSPIRVVersion(
-        static_cast<SPIRVWord>(VersionNumber::SPIRV_1_4));
+    Module->setMinSPIRVVersion(VersionNumber::SPIRV_1_4);
     addDecorate(new SPIRVDecorate(NoIntegerWrapDecoration, this));
     SPIRVDBG(spvdbgs() << "Set " << InstStr << " for obj " << Id << "\n")
   } else if (Module->isAllowedToUseExtension(