diff --git a/llvm/include/llvm/IR/IntrinsicsSPIRV.td b/llvm/include/llvm/IR/IntrinsicsSPIRV.td index 766b3b542a9e0..0eb09b1699aff 100644 --- a/llvm/include/llvm/IR/IntrinsicsSPIRV.td +++ b/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>, + ImmArg>]>; + def int_spv_lifetime_end : Intrinsic<[], + [llvm_i64_ty, llvm_anyptr_ty], + [IntrArgMemOnly, IntrWillReturn, + NoCapture>, + ImmArg>]>; + // 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">, diff --git a/llvm/lib/MC/SPIRVObjectWriter.cpp b/llvm/lib/MC/SPIRVObjectWriter.cpp index 39856e96e9be5..d72d6e07f2e6f 100644 --- a/llvm/lib/MC/SPIRVObjectWriter.cpp +++ b/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(MagicNumber); W.write(VersionNumber); W.write(GeneratorMagicNumber); diff --git a/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp b/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp index 1fbf3c3e11aed..30c67d3fde633 100644 --- a/llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp +++ b/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(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(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(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 = diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp index bda9c57e534c3..42f8397a3023b 100644 --- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp +++ b/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) diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h index 25d82ebf9bc79..da480b22a525f 100644 --- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h +++ b/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`. diff --git a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp b/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp index fd19b7412c4c9..0fef19c2d5341 100644 --- a/llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp +++ b/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(I).getIntrinsicID()) { + Intrinsic::ID IID = cast(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; } diff --git a/llvm/lib/Target/SPIRV/SPIRVMCInstLower.cpp b/llvm/lib/Target/SPIRV/SPIRVMCInstLower.cpp index 8c6649bf62826..afa550d6dd424 100644 --- a/llvm/lib/Target/SPIRV/SPIRVMCInstLower.cpp +++ b/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(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; } diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp index 2b4cb5ccc7b1e..00d0cbd763736 100644 --- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp +++ b/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; } diff --git a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.h b/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.h index 708384fc55f52..6e86eed30c5dc 100644 --- a/llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.h +++ b/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]; } diff --git a/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp b/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp index c376497469ce3..a8a0577f60564 100644 --- a/llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp +++ b/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 OpNos) { + Function *F = nullptr; + if (OpNos.empty()) { + F = Intrinsic::getDeclaration(II->getModule(), NewID); + } else { + SmallVector 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(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(*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; } } } diff --git a/llvm/lib/Target/SPIRV/SPIRVSubtarget.h b/llvm/lib/Target/SPIRV/SPIRVSubtarget.h index 62524ebfc9bf8..3b486226a9393 100644 --- a/llvm/lib/Target/SPIRV/SPIRVSubtarget.h +++ b/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 diff --git a/llvm/test/CodeGen/SPIRV/ComparePointers.ll b/llvm/test/CodeGen/SPIRV/ComparePointers.ll index 9be05944789b6..6777fc38024b3 100644 --- a/llvm/test/CodeGen/SPIRV/ComparePointers.ll +++ b/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) diff --git a/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll b/llvm/test/CodeGen/SPIRV/llvm-intrinsics/lifetime.ll new file mode 100644 index 0000000000000..710a1581f760c --- /dev/null +++ b/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) diff --git a/llvm/test/CodeGen/SPIRV/transcoding/AtomicCompareExchangeExplicit_cl20.ll b/llvm/test/CodeGen/SPIRV/transcoding/AtomicCompareExchangeExplicit_cl20.ll index 55cfcea999d84..e0c47798cc6d0 100644 --- a/llvm/test/CodeGen/SPIRV/transcoding/AtomicCompareExchangeExplicit_cl20.ll +++ b/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, diff --git a/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars.ll b/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars.ll index f18f27a6de51d..5074893163565 100644 --- a/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars.ll +++ b/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 %[[#]] diff --git a/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_arithmetics.ll b/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_arithmetics.ll index d39ca3c39383c..d0c4dff43121c 100644 --- a/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_arithmetics.ll +++ b/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 diff --git a/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_opt.ll b/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_opt.ll index 03456aef6b6b2..3885f07023144 100644 --- a/llvm/test/CodeGen/SPIRV/transcoding/builtin_vars_opt.ll +++ b/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