[SPIRV] Add Lifetime intrinsics/instructions

llvm/include/llvm/IR/IntrinsicsSPIRV.td

@@ -40,6 +40,18 @@ let TargetPrefix = "spv" in {
   def int_spv_assume : Intrinsic<[], [llvm_i1_ty]>;
   def int_spv_expect : Intrinsic<[llvm_anyint_ty], [LLVMMatchType<0>, LLVMMatchType<0>]>;
 
+  // Memory Use Markers
+  def int_spv_lifetime_start : Intrinsic<[],
+                                [llvm_i64_ty, llvm_anyptr_ty],
+                                [IntrArgMemOnly, IntrWillReturn,
+                                NoCapture<ArgIndex<1>>,
+                                ImmArg<ArgIndex<0>>]>;
+  def int_spv_lifetime_end   : Intrinsic<[],
+                                [llvm_i64_ty, llvm_anyptr_ty],
+                                [IntrArgMemOnly, IntrWillReturn,
+                                NoCapture<ArgIndex<1>>,
+                                ImmArg<ArgIndex<0>>]>;
+
   // The following intrinsic(s) are mirrored from IntrinsicsDirectX.td for HLSL support.
   def int_spv_thread_id : Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrWillReturn]>;
   def int_spv_create_handle : ClangBuiltin<"__builtin_hlsl_create_handle">,

llvm/lib/MC/SPIRVObjectWriter.cpp

@@ -43,18 +43,14 @@ class SPIRVObjectWriter : public MCObjectWriter {
 
 void SPIRVObjectWriter::writeHeader(const MCAssembler &Asm) {
   constexpr uint32_t MagicNumber = 0x07230203;
-
-  // TODO: set the version on a min-necessary basis (just like the translator
-  // does) requires some refactoring of MCAssembler::VersionInfoType.
-  constexpr uint32_t Major = 1;
-  constexpr uint32_t Minor = 0;
-  constexpr uint32_t VersionNumber = 0 | (Major << 16) | (Minor << 8);
-  // TODO: check if we could use anything other than 0 (spec allows).
   constexpr uint32_t GeneratorMagicNumber = 0;
-  // TODO: do not hardcode this as well.
-  constexpr uint32_t Bound = 900;
   constexpr uint32_t Schema = 0;
 
+  // Construct SPIR-V version and Bound
+  const MCAssembler::VersionInfoType &VIT = Asm.getVersionInfo();
+  uint32_t VersionNumber = 0 | (VIT.Major << 16) | (VIT.Minor << 8);
+  uint32_t Bound = VIT.Update;
+
   W.write<uint32_t>(MagicNumber);
   W.write<uint32_t>(VersionNumber);
   W.write<uint32_t>(GeneratorMagicNumber);

llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp

@@ -29,7 +29,9 @@
 #include "llvm/CodeGen/MachineModuleInfo.h"
 #include "llvm/CodeGen/TargetLoweringObjectFileImpl.h"
 #include "llvm/MC/MCAsmInfo.h"
+#include "llvm/MC/MCAssembler.h"
 #include "llvm/MC/MCInst.h"
+#include "llvm/MC/MCObjectStreamer.h"
 #include "llvm/MC/MCStreamer.h"
 #include "llvm/MC/MCSymbol.h"
 #include "llvm/MC/TargetRegistry.h"
@@ -101,6 +103,21 @@ void SPIRVAsmPrinter::emitEndOfAsmFile(Module &M) {
   if (ModuleSectionsEmitted == false) {
     outputModuleSections();
     ModuleSectionsEmitted = true;
+  } else {
+    ST = static_cast<const SPIRVTargetMachine &>(TM).getSubtargetImpl();
+    uint32_t DecSPIRVVersion = ST->getSPIRVVersion();
+    uint32_t Major = DecSPIRVVersion / 10;
+    uint32_t Minor = DecSPIRVVersion - Major * 10;
+    // TODO: calculate Bound more carefully from maximum used register number,
+    // accounting for generated OpLabels and other related instructions if
+    // needed.
+    unsigned Bound = 2 * (ST->getBound() + 1);
+    bool FlagToRestore = OutStreamer->getUseAssemblerInfoForParsing();
+    OutStreamer->setUseAssemblerInfoForParsing(true);
+    if (MCAssembler *Asm = OutStreamer->getAssemblerPtr())
+      Asm->setBuildVersion(static_cast<MachO::PlatformType>(0), Major, Minor,
+                           Bound, VersionTuple(Major, Minor, 0, Bound));
+    OutStreamer->setUseAssemblerInfoForParsing(FlagToRestore);
   }
 }
 
@@ -507,6 +524,13 @@ void SPIRVAsmPrinter::outputAnnotations(const Module &M) {
         report_fatal_error("Unsupported value in llvm.global.annotations");
       Function *Func = cast<Function>(AnnotatedVar);
       Register Reg = MAI->getFuncReg(Func);
+      if (!Reg.isValid()) {
+        std::string DiagMsg;
+        raw_string_ostream OS(DiagMsg);
+        AnnotatedVar->print(OS);
+        DiagMsg = "Unknown function in llvm.global.annotations: " + DiagMsg;
+        report_fatal_error(DiagMsg.c_str());
+      }
 
       // The second field contains a pointer to a global annotation string.
       GlobalVariable *GV =

llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp

@@ -24,7 +24,7 @@
 
 using namespace llvm;
 SPIRVGlobalRegistry::SPIRVGlobalRegistry(unsigned PointerSize)
-    : PointerSize(PointerSize) {}
+    : PointerSize(PointerSize), Bound(0) {}
 
 SPIRVType *SPIRVGlobalRegistry::assignIntTypeToVReg(unsigned BitWidth,
                                                     Register VReg,
@@ -896,6 +896,15 @@ bool SPIRVGlobalRegistry::isScalarOrVectorSigned(const SPIRVType *Type) const {
   return IntType && IntType->getOperand(2).getImm() != 0;
 }
 
+unsigned SPIRVGlobalRegistry::getPointeeTypeOp(Register PtrReg) {
+  SPIRVType *PtrType = getSPIRVTypeForVReg(PtrReg);
+  SPIRVType *ElemType =
+      PtrType && PtrType->getOpcode() == SPIRV::OpTypePointer
+          ? getSPIRVTypeForVReg(PtrType->getOperand(2).getReg())
+          : nullptr;
+  return ElemType ? ElemType->getOpcode() : 0;
+}
+
 bool SPIRVGlobalRegistry::isBitcastCompatible(const SPIRVType *Type1,
                                               const SPIRVType *Type2) const {
   if (!Type1 || !Type2)

llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h

@@ -56,6 +56,9 @@ class SPIRVGlobalRegistry {
   // Number of bits pointers and size_t integers require.
   const unsigned PointerSize;
 
+  // Holds the maximum ID we have in the module.
+  unsigned Bound;
+
   // Add a new OpTypeXXX instruction without checking for duplicates.
   SPIRVType *createSPIRVType(const Type *Type, MachineIRBuilder &MIRBuilder,
                              SPIRV::AccessQualifier::AccessQualifier AQ =
@@ -108,6 +111,9 @@ class SPIRVGlobalRegistry {
     DT.buildDepsGraph(Graph, MMI);
   }
 
+  void setBound(unsigned V) { Bound = V; }
+  unsigned getBound() { return Bound; }
+
   // Map a machine operand that represents a use of a function via function
   // pointer to a machine operand that represents the function definition.
   // Return either the register or invalid value, because we have no context for
@@ -166,6 +172,9 @@ class SPIRVGlobalRegistry {
     return Res->second;
   }
 
+  // Return a pointee's type op code, or 0 otherwise.
+  unsigned getPointeeTypeOp(Register PtrReg);
+
   // Either generate a new OpTypeXXX instruction or return an existing one
   // corresponding to the given string containing the name of the builtin type.
   // Return nullptr if unable to recognize SPIRV type name from `TypeStr`.

llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp

@@ -1567,7 +1567,8 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
                                                const SPIRVType *ResType,
                                                MachineInstr &I) const {
   MachineBasicBlock &BB = *I.getParent();
-  switch (cast<GIntrinsic>(I).getIntrinsicID()) {
+  Intrinsic::ID IID = cast<GIntrinsic>(I).getIntrinsicID();
+  switch (IID) {
   case Intrinsic::spv_load:
     return selectLoad(ResVReg, ResType, I);
   case Intrinsic::spv_store:
@@ -1661,8 +1662,25 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
     break;
   case Intrinsic::spv_thread_id:
     return selectSpvThreadId(ResVReg, ResType, I);
-  default:
-    llvm_unreachable("Intrinsic selection not implemented");
+  case Intrinsic::spv_lifetime_start:
+  case Intrinsic::spv_lifetime_end: {
+    unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
+                                                       : SPIRV::OpLifetimeStop;
+    int64_t Size = I.getOperand(I.getNumExplicitDefs() + 1).getImm();
+    Register PtrReg = I.getOperand(I.getNumExplicitDefs() + 2).getReg();
+    unsigned PonteeOpType = GR.getPointeeTypeOp(PtrReg);
+    bool IsNonvoidPtr = PonteeOpType != 0 && PonteeOpType != SPIRV::OpTypeVoid;
+    if (Size == -1 || IsNonvoidPtr)
+      Size = 0;
+    BuildMI(BB, I, I.getDebugLoc(), TII.get(Op)).addUse(PtrReg).addImm(Size);
+  } break;
+  default: {
+    std::string DiagMsg;
+    raw_string_ostream OS(DiagMsg);
+    I.print(OS);
+    DiagMsg = "Intrinsic selection not implemented: " + DiagMsg;
+    report_fatal_error(DiagMsg.c_str(), false);
+  }
   }
   return true;
 }

llvm/lib/Target/SPIRV/SPIRVMCInstLower.cpp

@@ -34,7 +34,13 @@ void SPIRVMCInstLower::lower(const MachineInstr *MI, MCInst &OutMI,
       llvm_unreachable("unknown operand type");
     case MachineOperand::MO_GlobalAddress: {
       Register FuncReg = MAI->getFuncReg(dyn_cast<Function>(MO.getGlobal()));
-      assert(FuncReg.isValid() && "Cannot find function Id");
+      if (!FuncReg.isValid()) {
+        std::string DiagMsg;
+        raw_string_ostream OS(DiagMsg);
+        MI->print(OS);
+        DiagMsg = "Unknown function in:" + DiagMsg;
+        report_fatal_error(DiagMsg.c_str());
+      }
       MCOp = MCOperand::createReg(FuncReg);
       break;
     }

llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp

@@ -1309,5 +1309,8 @@ bool SPIRVModuleAnalysis::runOnModule(Module &M) {
   if (MAI.MS[SPIRV::MB_EntryPoints].empty())
     MAI.Reqs.addCapability(SPIRV::Capability::Linkage);
 
+  // Set maximum ID used.
+  GR->setBound(MAI.MaxID);
+
   return false;
 }

llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.h

@@ -163,8 +163,8 @@ struct ModuleAnalysisInfo {
   Register getFuncReg(const Function *F) {
     assert(F && "Function is null");
     auto FuncPtrRegPair = FuncMap.find(F);
-    assert(FuncPtrRegPair != FuncMap.end() && "Cannot find function ID");
-    return FuncPtrRegPair->second;
+    return FuncPtrRegPair == FuncMap.end() ? Register(0)
+                                           : FuncPtrRegPair->second;
   }
   Register getExtInstSetReg(unsigned SetNum) { return ExtInstSetMap[SetNum]; }
   InstrList &getMSInstrs(unsigned MSType) { return MS[MSType]; }

llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp

@@ -263,6 +263,21 @@ static void lowerExpectAssume(IntrinsicInst *II) {
   return;
 }
 
+static bool toSpvOverloadedIntrinsic(IntrinsicInst *II, Intrinsic::ID NewID,
+                                     ArrayRef<unsigned> OpNos) {
+  Function *F = nullptr;
+  if (OpNos.empty()) {
+    F = Intrinsic::getDeclaration(II->getModule(), NewID);
+  } else {
+    SmallVector<Type *, 4> Tys;
+    for (unsigned OpNo : OpNos)
+      Tys.push_back(II->getOperand(OpNo)->getType());
+    F = Intrinsic::getDeclaration(II->getModule(), NewID, Tys);
+  }
+  II->setCalledFunction(F);
+  return true;
+}
+
 static void lowerUMulWithOverflow(IntrinsicInst *UMulIntrinsic) {
   // Get a separate function - otherwise, we'd have to rework the CFG of the
   // current one. Then simply replace the intrinsic uses with a call to the new
@@ -290,22 +305,35 @@ bool SPIRVPrepareFunctions::substituteIntrinsicCalls(Function *F) {
       if (!CF || !CF->isIntrinsic())
         continue;
       auto *II = cast<IntrinsicInst>(Call);
-      if (II->getIntrinsicID() == Intrinsic::memset ||
-          II->getIntrinsicID() == Intrinsic::bswap)
+      switch (II->getIntrinsicID()) {
+      case Intrinsic::memset:
+      case Intrinsic::bswap:
         Changed |= lowerIntrinsicToFunction(II);
-      else if (II->getIntrinsicID() == Intrinsic::fshl ||
-               II->getIntrinsicID() == Intrinsic::fshr) {
+        break;
+      case Intrinsic::fshl:
+      case Intrinsic::fshr:
         lowerFunnelShifts(II);
         Changed = true;
-      } else if (II->getIntrinsicID() == Intrinsic::umul_with_overflow) {
+        break;
+      case Intrinsic::umul_with_overflow:
         lowerUMulWithOverflow(II);
         Changed = true;
-      } else if (II->getIntrinsicID() == Intrinsic::assume ||
-                 II->getIntrinsicID() == Intrinsic::expect) {
+        break;
+      case Intrinsic::assume:
+      case Intrinsic::expect: {
         const SPIRVSubtarget &STI = TM.getSubtarget<SPIRVSubtarget>(*F);
         if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_expect_assume))
           lowerExpectAssume(II);
         Changed = true;
+      } break;
+      case Intrinsic::lifetime_start:
+        Changed |= toSpvOverloadedIntrinsic(
+            II, Intrinsic::SPVIntrinsics::spv_lifetime_start, {1});
+        break;
+      case Intrinsic::lifetime_end:
+        Changed |= toSpvOverloadedIntrinsic(
+            II, Intrinsic::SPVIntrinsics::spv_lifetime_end, {1});
+        break;
       }
     }
   }

llvm/lib/Target/SPIRV/SPIRVSubtarget.h

@@ -71,6 +71,7 @@ class SPIRVSubtarget : public SPIRVGenSubtargetInfo {
   // The definition of this function is auto generated by tblgen.
   void ParseSubtargetFeatures(StringRef CPU, StringRef TuneCPU, StringRef FS);
   unsigned getPointerSize() const { return PointerSize; }
+  unsigned getBound() const { return GR->getBound(); }
   bool canDirectlyComparePointers() const;
   // TODO: this environment is not implemented in Triple, we need to decide
   // how to standardize its support. For now, let's assume SPIR-V with physical

llvm/test/CodeGen/SPIRV/ComparePointers.ll

@@ -1,5 +1,5 @@
 ; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --mattr=+spirv1.3  %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
 
 ;; kernel void test(int global *in, int global *in2) {
 ;;   if (!in)

llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll

@@ -0,0 +1,25 @@
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
+
+; CHECK: OpFunction
+; CHECK: %[[FooArg:.*]] = OpVariable
+; CHECK: OpLifetimeStart %[[FooArg]], 0
+; CHECK: OpCopyMemorySized
+; CHECK: OpBitcast
+; CHECK: OpInBoundsPtrAccessChain
+; CHECK: OpLifetimeStop %[[FooArg]], 0
+
+%tprange = type { %tparray }
+%tparray = type { [2 x i64] }
+
+define spir_func void @foo(ptr noundef byval(%tprange) align 8 %_arg_UserRange) {
+  %RoundedRangeKernel = alloca %tprange, align 8
+  call void @llvm.lifetime.start.p0(i64 72, ptr nonnull %RoundedRangeKernel) #7
+  call void @llvm.memcpy.p0.p0.i64(ptr align 8 %RoundedRangeKernel, ptr align 8 %_arg_UserRange, i64 16, i1 false)
+  %KernelFunc = getelementptr inbounds i8, ptr %RoundedRangeKernel, i64 16
+  call void @llvm.lifetime.end.p0(i64 72, ptr nonnull %RoundedRangeKernel) #7
+  ret void
+}
+
+declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture)
+declare void @llvm.memcpy.p0.p0.i64(ptr noalias nocapture writeonly, ptr noalias nocapture readonly, i64, i1 immarg)
+declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture)

llvm/test/CodeGen/SPIRV/transcoding/AtomicCompareExchangeExplicit_cl20.ll

@@ -1,5 +1,5 @@
-; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --mattr=+spirv1.3 %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown --mattr=+spirv1.3 %s -o - -filetype=obj | spirv-val %}
 
 ;; __kernel void testAtomicCompareExchangeExplicit_cl20(
 ;;     volatile global atomic_int* object,

llvm/test/CodeGen/SPIRV/transcoding/builtin_vars.ll

@@ -1,5 +1,5 @@
 ; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
 
 ; CHECK-SPIRV: OpDecorate %[[#Id:]] BuiltIn GlobalLinearId
 ; CHECK-SPIRV: %[[#Id:]] = OpVariable %[[#]]

llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_arithmetics.ll

@@ -1,5 +1,5 @@
 ; RUN: llc -O0 -mtriple=spirv64-unknown-linux %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
 
 ;; The IR was generated from the following source:
 ;; #include <CL/sycl.hpp>

llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_opt.ll

@@ -1,5 +1,5 @@
 ; RUN: llc -O0 -mtriple=spirv64-unknown-linux %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
-; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
+; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
 
 ;; The IR was generated from the following source:
 ;; #include <CL/sycl.hpp>