KhronosGroup · AlexeySotkin · Dec 21, 2021 · Dec 10, 2021 · Dec 13, 2021 · Dec 13, 2021
diff --git a/lib/SPIRV/SPIRVWriter.cpp b/lib/SPIRV/SPIRVWriter.cpp
@@ -606,9 +606,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
@@ -3567,12 +3565,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;
 
@@ -3584,9 +3585,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(
@@ -3689,10 +3696,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/lib/SPIRV/SPIRVWriter.h b/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/lib/SPIRV/libSPIRV/SPIRVDecorate.h b/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/lib/SPIRV/libSPIRV/SPIRVEntry.cpp b/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/lib/SPIRV/libSPIRV/SPIRVEntry.h b/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/lib/SPIRV/libSPIRV/SPIRVModule.cpp b/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/lib/SPIRV/libSPIRV/SPIRVModule.h b/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/lib/SPIRV/libSPIRV/SPIRVStream.cpp b/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/test/ExecutionMode.ll b/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/test/copy_object.spt b/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/test/entry-point-interfaces.ll b/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/test/negative/unimplemented.spt b/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/test/right_shift.spt b/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/test/transcoding/SPV_INTEL_inline_assembly/inline_asm_clobbers.cl b/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) {