Skip to content

Commit 59f34e8

Browse files
VyacheslavLevytskyyVyacheslavLevytskyy
authored
[SPIRV] Add Lifetime intrinsics/instructions (#85391)
This PR: * adds Lifetime intrinsics/instructions * fixes how the binary header is emitted (correct version and better approximation of Bound) * add validation into more test cases
1 parent 38a44bd commit 59f34e8

17 files changed

+160
-29
lines changed

‎llvm/include/llvm/IR/IntrinsicsSPIRV.td

Lines changed: 12 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -40,6 +40,18 @@ let TargetPrefix = "spv" in {
4040
def int_spv_assume : Intrinsic<[], [llvm_i1_ty]>;
4141
def int_spv_expect : Intrinsic<[llvm_anyint_ty], [LLVMMatchType<0>, LLVMMatchType<0>]>;
4242

43+
// Memory Use Markers
44+
def int_spv_lifetime_start : Intrinsic<[],
45+
[llvm_i64_ty, llvm_anyptr_ty],
46+
[IntrArgMemOnly, IntrWillReturn,
47+
NoCapture<ArgIndex<1>>,
48+
ImmArg<ArgIndex<0>>]>;
49+
def int_spv_lifetime_end : Intrinsic<[],
50+
[llvm_i64_ty, llvm_anyptr_ty],
51+
[IntrArgMemOnly, IntrWillReturn,
52+
NoCapture<ArgIndex<1>>,
53+
ImmArg<ArgIndex<0>>]>;
54+
4355
// The following intrinsic(s) are mirrored from IntrinsicsDirectX.td for HLSL support.
4456
def int_spv_thread_id : Intrinsic<[llvm_i32_ty], [llvm_i32_ty], [IntrNoMem, IntrWillReturn]>;
4557
def int_spv_create_handle : ClangBuiltin<"__builtin_hlsl_create_handle">,

‎llvm/lib/MC/SPIRVObjectWriter.cpp

Lines changed: 5 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -43,18 +43,14 @@ class SPIRVObjectWriter : public MCObjectWriter {
4343

4444
void SPIRVObjectWriter::writeHeader(const MCAssembler &Asm) {
4545
constexpr uint32_t MagicNumber = 0x07230203;
46-
47-
// TODO: set the version on a min-necessary basis (just like the translator
48-
// does) requires some refactoring of MCAssembler::VersionInfoType.
49-
constexpr uint32_t Major = 1;
50-
constexpr uint32_t Minor = 0;
51-
constexpr uint32_t VersionNumber = 0 | (Major << 16) | (Minor << 8);
52-
// TODO: check if we could use anything other than 0 (spec allows).
5346
constexpr uint32_t GeneratorMagicNumber = 0;
54-
// TODO: do not hardcode this as well.
55-
constexpr uint32_t Bound = 900;
5647
constexpr uint32_t Schema = 0;
5748

49+
// Construct SPIR-V version and Bound
50+
const MCAssembler::VersionInfoType &VIT = Asm.getVersionInfo();
51+
uint32_t VersionNumber = 0 | (VIT.Major << 16) | (VIT.Minor << 8);
52+
uint32_t Bound = VIT.Update;
53+
5854
W.write<uint32_t>(MagicNumber);
5955
W.write<uint32_t>(VersionNumber);
6056
W.write<uint32_t>(GeneratorMagicNumber);

‎llvm/lib/Target/SPIRV/SPIRVAsmPrinter.cpp

Lines changed: 24 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -29,7 +29,9 @@
2929
#include "llvm/CodeGen/MachineModuleInfo.h"
3030
#include "llvm/CodeGen/TargetLoweringObjectFileImpl.h"
3131
#include "llvm/MC/MCAsmInfo.h"
32+
#include "llvm/MC/MCAssembler.h"
3233
#include "llvm/MC/MCInst.h"
34+
#include "llvm/MC/MCObjectStreamer.h"
3335
#include "llvm/MC/MCStreamer.h"
3436
#include "llvm/MC/MCSymbol.h"
3537
#include "llvm/MC/TargetRegistry.h"
@@ -101,6 +103,21 @@ void SPIRVAsmPrinter::emitEndOfAsmFile(Module &M) {
101103
if (ModuleSectionsEmitted == false) {
102104
outputModuleSections();
103105
ModuleSectionsEmitted = true;
106+
} else {
107+
ST = static_cast<const SPIRVTargetMachine &>(TM).getSubtargetImpl();
108+
uint32_t DecSPIRVVersion = ST->getSPIRVVersion();
109+
uint32_t Major = DecSPIRVVersion / 10;
110+
uint32_t Minor = DecSPIRVVersion - Major * 10;
111+
// TODO: calculate Bound more carefully from maximum used register number,
112+
// accounting for generated OpLabels and other related instructions if
113+
// needed.
114+
unsigned Bound = 2 * (ST->getBound() + 1);
115+
bool FlagToRestore = OutStreamer->getUseAssemblerInfoForParsing();
116+
OutStreamer->setUseAssemblerInfoForParsing(true);
117+
if (MCAssembler *Asm = OutStreamer->getAssemblerPtr())
118+
Asm->setBuildVersion(static_cast<MachO::PlatformType>(0), Major, Minor,
119+
Bound, VersionTuple(Major, Minor, 0, Bound));
120+
OutStreamer->setUseAssemblerInfoForParsing(FlagToRestore);
104121
}
105122
}
106123

@@ -507,6 +524,13 @@ void SPIRVAsmPrinter::outputAnnotations(const Module &M) {
507524
report_fatal_error("Unsupported value in llvm.global.annotations");
508525
Function *Func = cast<Function>(AnnotatedVar);
509526
Register Reg = MAI->getFuncReg(Func);
527+
if (!Reg.isValid()) {
528+
std::string DiagMsg;
529+
raw_string_ostream OS(DiagMsg);
530+
AnnotatedVar->print(OS);
531+
DiagMsg = "Unknown function in llvm.global.annotations: " + DiagMsg;
532+
report_fatal_error(DiagMsg.c_str());
533+
}
510534

511535
// The second field contains a pointer to a global annotation string.
512536
GlobalVariable *GV =

‎llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.cpp

Lines changed: 10 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -24,7 +24,7 @@
2424

2525
using namespace llvm;
2626
SPIRVGlobalRegistry::SPIRVGlobalRegistry(unsigned PointerSize)
27-
: PointerSize(PointerSize) {}
27+
: PointerSize(PointerSize), Bound(0) {}
2828

2929
SPIRVType *SPIRVGlobalRegistry::assignIntTypeToVReg(unsigned BitWidth,
3030
Register VReg,
@@ -896,6 +896,15 @@ bool SPIRVGlobalRegistry::isScalarOrVectorSigned(const SPIRVType *Type) const {
896896
return IntType && IntType->getOperand(2).getImm() != 0;
897897
}
898898

899+
unsigned SPIRVGlobalRegistry::getPointeeTypeOp(Register PtrReg) {
900+
SPIRVType *PtrType = getSPIRVTypeForVReg(PtrReg);
901+
SPIRVType *ElemType =
902+
PtrType && PtrType->getOpcode() == SPIRV::OpTypePointer
903+
? getSPIRVTypeForVReg(PtrType->getOperand(2).getReg())
904+
: nullptr;
905+
return ElemType ? ElemType->getOpcode() : 0;
906+
}
907+
899908
bool SPIRVGlobalRegistry::isBitcastCompatible(const SPIRVType *Type1,
900909
const SPIRVType *Type2) const {
901910
if (!Type1 || !Type2)

‎llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h

Lines changed: 9 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -56,6 +56,9 @@ class SPIRVGlobalRegistry {
5656
// Number of bits pointers and size_t integers require.
5757
const unsigned PointerSize;
5858

59+
// Holds the maximum ID we have in the module.
60+
unsigned Bound;
61+
5962
// Add a new OpTypeXXX instruction without checking for duplicates.
6063
SPIRVType *createSPIRVType(const Type *Type, MachineIRBuilder &MIRBuilder,
6164
SPIRV::AccessQualifier::AccessQualifier AQ =
@@ -108,6 +111,9 @@ class SPIRVGlobalRegistry {
108111
DT.buildDepsGraph(Graph, MMI);
109112
}
110113

114+
void setBound(unsigned V) { Bound = V; }
115+
unsigned getBound() { return Bound; }
116+
111117
// Map a machine operand that represents a use of a function via function
112118
// pointer to a machine operand that represents the function definition.
113119
// Return either the register or invalid value, because we have no context for
@@ -166,6 +172,9 @@ class SPIRVGlobalRegistry {
166172
return Res->second;
167173
}
168174

175+
// Return a pointee's type op code, or 0 otherwise.
176+
unsigned getPointeeTypeOp(Register PtrReg);
177+
169178
// Either generate a new OpTypeXXX instruction or return an existing one
170179
// corresponding to the given string containing the name of the builtin type.
171180
// Return nullptr if unable to recognize SPIRV type name from `TypeStr`.

‎llvm/lib/Target/SPIRV/SPIRVInstructionSelector.cpp

Lines changed: 21 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1567,7 +1567,8 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
15671567
const SPIRVType *ResType,
15681568
MachineInstr &I) const {
15691569
MachineBasicBlock &BB = *I.getParent();
1570-
switch (cast<GIntrinsic>(I).getIntrinsicID()) {
1570+
Intrinsic::ID IID = cast<GIntrinsic>(I).getIntrinsicID();
1571+
switch (IID) {
15711572
case Intrinsic::spv_load:
15721573
return selectLoad(ResVReg, ResType, I);
15731574
case Intrinsic::spv_store:
@@ -1661,8 +1662,25 @@ bool SPIRVInstructionSelector::selectIntrinsic(Register ResVReg,
16611662
break;
16621663
case Intrinsic::spv_thread_id:
16631664
return selectSpvThreadId(ResVReg, ResType, I);
1664-
default:
1665-
llvm_unreachable("Intrinsic selection not implemented");
1665+
case Intrinsic::spv_lifetime_start:
1666+
case Intrinsic::spv_lifetime_end: {
1667+
unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
1668+
: SPIRV::OpLifetimeStop;
1669+
int64_t Size = I.getOperand(I.getNumExplicitDefs() + 1).getImm();
1670+
Register PtrReg = I.getOperand(I.getNumExplicitDefs() + 2).getReg();
1671+
unsigned PonteeOpType = GR.getPointeeTypeOp(PtrReg);
1672+
bool IsNonvoidPtr = PonteeOpType != 0 && PonteeOpType != SPIRV::OpTypeVoid;
1673+
if (Size == -1 || IsNonvoidPtr)
1674+
Size = 0;
1675+
BuildMI(BB, I, I.getDebugLoc(), TII.get(Op)).addUse(PtrReg).addImm(Size);
1676+
} break;
1677+
default: {
1678+
std::string DiagMsg;
1679+
raw_string_ostream OS(DiagMsg);
1680+
I.print(OS);
1681+
DiagMsg = "Intrinsic selection not implemented: " + DiagMsg;
1682+
report_fatal_error(DiagMsg.c_str(), false);
1683+
}
16661684
}
16671685
return true;
16681686
}

‎llvm/lib/Target/SPIRV/SPIRVMCInstLower.cpp

Lines changed: 7 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -34,7 +34,13 @@ void SPIRVMCInstLower::lower(const MachineInstr *MI, MCInst &OutMI,
3434
llvm_unreachable("unknown operand type");
3535
case MachineOperand::MO_GlobalAddress: {
3636
Register FuncReg = MAI->getFuncReg(dyn_cast<Function>(MO.getGlobal()));
37-
assert(FuncReg.isValid() && "Cannot find function Id");
37+
if (!FuncReg.isValid()) {
38+
std::string DiagMsg;
39+
raw_string_ostream OS(DiagMsg);
40+
MI->print(OS);
41+
DiagMsg = "Unknown function in:" + DiagMsg;
42+
report_fatal_error(DiagMsg.c_str());
43+
}
3844
MCOp = MCOperand::createReg(FuncReg);
3945
break;
4046
}

‎llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.cpp

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1309,5 +1309,8 @@ bool SPIRVModuleAnalysis::runOnModule(Module &M) {
13091309
if (MAI.MS[SPIRV::MB_EntryPoints].empty())
13101310
MAI.Reqs.addCapability(SPIRV::Capability::Linkage);
13111311

1312+
// Set maximum ID used.
1313+
GR->setBound(MAI.MaxID);
1314+
13121315
return false;
13131316
}

‎llvm/lib/Target/SPIRV/SPIRVModuleAnalysis.h

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -163,8 +163,8 @@ struct ModuleAnalysisInfo {
163163
Register getFuncReg(const Function *F) {
164164
assert(F && "Function is null");
165165
auto FuncPtrRegPair = FuncMap.find(F);
166-
assert(FuncPtrRegPair != FuncMap.end() && "Cannot find function ID");
167-
return FuncPtrRegPair->second;
166+
return FuncPtrRegPair == FuncMap.end() ? Register(0)
167+
: FuncPtrRegPair->second;
168168
}
169169
Register getExtInstSetReg(unsigned SetNum) { return ExtInstSetMap[SetNum]; }
170170
InstrList &getMSInstrs(unsigned MSType) { return MS[MSType]; }

‎llvm/lib/Target/SPIRV/SPIRVPrepareFunctions.cpp

Lines changed: 35 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -263,6 +263,21 @@ static void lowerExpectAssume(IntrinsicInst *II) {
263263
return;
264264
}
265265

266+
static bool toSpvOverloadedIntrinsic(IntrinsicInst *II, Intrinsic::ID NewID,
267+
ArrayRef<unsigned> OpNos) {
268+
Function *F = nullptr;
269+
if (OpNos.empty()) {
270+
F = Intrinsic::getDeclaration(II->getModule(), NewID);
271+
} else {
272+
SmallVector<Type *, 4> Tys;
273+
for (unsigned OpNo : OpNos)
274+
Tys.push_back(II->getOperand(OpNo)->getType());
275+
F = Intrinsic::getDeclaration(II->getModule(), NewID, Tys);
276+
}
277+
II->setCalledFunction(F);
278+
return true;
279+
}
280+
266281
static void lowerUMulWithOverflow(IntrinsicInst *UMulIntrinsic) {
267282
// Get a separate function - otherwise, we'd have to rework the CFG of the
268283
// current one. Then simply replace the intrinsic uses with a call to the new
@@ -290,22 +305,35 @@ bool SPIRVPrepareFunctions::substituteIntrinsicCalls(Function *F) {
290305
if (!CF || !CF->isIntrinsic())
291306
continue;
292307
auto *II = cast<IntrinsicInst>(Call);
293-
if (II->getIntrinsicID() == Intrinsic::memset ||
294-
II->getIntrinsicID() == Intrinsic::bswap)
308+
switch (II->getIntrinsicID()) {
309+
case Intrinsic::memset:
310+
case Intrinsic::bswap:
295311
Changed |= lowerIntrinsicToFunction(II);
296-
else if (II->getIntrinsicID() == Intrinsic::fshl ||
297-
II->getIntrinsicID() == Intrinsic::fshr) {
312+
break;
313+
case Intrinsic::fshl:
314+
case Intrinsic::fshr:
298315
lowerFunnelShifts(II);
299316
Changed = true;
300-
} else if (II->getIntrinsicID() == Intrinsic::umul_with_overflow) {
317+
break;
318+
case Intrinsic::umul_with_overflow:
301319
lowerUMulWithOverflow(II);
302320
Changed = true;
303-
} else if (II->getIntrinsicID() == Intrinsic::assume ||
304-
II->getIntrinsicID() == Intrinsic::expect) {
321+
break;
322+
case Intrinsic::assume:
323+
case Intrinsic::expect: {
305324
const SPIRVSubtarget &STI = TM.getSubtarget<SPIRVSubtarget>(*F);
306325
if (STI.canUseExtension(SPIRV::Extension::SPV_KHR_expect_assume))
307326
lowerExpectAssume(II);
308327
Changed = true;
328+
} break;
329+
case Intrinsic::lifetime_start:
330+
Changed |= toSpvOverloadedIntrinsic(
331+
II, Intrinsic::SPVIntrinsics::spv_lifetime_start, {1});
332+
break;
333+
case Intrinsic::lifetime_end:
334+
Changed |= toSpvOverloadedIntrinsic(
335+
II, Intrinsic::SPVIntrinsics::spv_lifetime_end, {1});
336+
break;
309337
}
310338
}
311339
}

‎llvm/lib/Target/SPIRV/SPIRVSubtarget.h

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -71,6 +71,7 @@ class SPIRVSubtarget : public SPIRVGenSubtargetInfo {
7171
// The definition of this function is auto generated by tblgen.
7272
void ParseSubtargetFeatures(StringRef CPU, StringRef TuneCPU, StringRef FS);
7373
unsigned getPointerSize() const { return PointerSize; }
74+
unsigned getBound() const { return GR->getBound(); }
7475
bool canDirectlyComparePointers() const;
7576
// TODO: this environment is not implemented in Triple, we need to decide
7677
// how to standardize its support. For now, let's assume SPIR-V with physical

‎llvm/test/CodeGen/SPIRV/ComparePointers.ll

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
; RUN: llc -O0 -mtriple=spirv64-unknown-unknown --mattr=+spirv1.3 %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
2-
; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
2+
; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
33

44
;; kernel void test(int global *in, int global *in2) {
55
;; if (!in)

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

Lines changed: 25 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,25 @@
1+
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s
2+
3+
; CHECK: OpFunction
4+
; CHECK: %[[FooArg:.*]] = OpVariable
5+
; CHECK: OpLifetimeStart %[[FooArg]], 0
6+
; CHECK: OpCopyMemorySized
7+
; CHECK: OpBitcast
8+
; CHECK: OpInBoundsPtrAccessChain
9+
; CHECK: OpLifetimeStop %[[FooArg]], 0
10+
11+
%tprange = type { %tparray }
12+
%tparray = type { [2 x i64] }
13+
14+
define spir_func void @foo(ptr noundef byval(%tprange) align 8 %_arg_UserRange) {
15+
%RoundedRangeKernel = alloca %tprange, align 8
16+
call void @llvm.lifetime.start.p0(i64 72, ptr nonnull %RoundedRangeKernel) #7
17+
call void @llvm.memcpy.p0.p0.i64(ptr align 8 %RoundedRangeKernel, ptr align 8 %_arg_UserRange, i64 16, i1 false)
18+
%KernelFunc = getelementptr inbounds i8, ptr %RoundedRangeKernel, i64 16
19+
call void @llvm.lifetime.end.p0(i64 72, ptr nonnull %RoundedRangeKernel) #7
20+
ret void
21+
}
22+
23+
declare void @llvm.lifetime.start.p0(i64 immarg, ptr nocapture)
24+
declare void @llvm.memcpy.p0.p0.i64(ptr noalias nocapture writeonly, ptr noalias nocapture readonly, i64, i1 immarg)
25+
declare void @llvm.lifetime.end.p0(i64 immarg, ptr nocapture)

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
1-
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
2-
; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
1+
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown --mattr=+spirv1.3 %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
2+
; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown --mattr=+spirv1.3 %s -o - -filetype=obj | spirv-val %}
33

44
;; __kernel void testAtomicCompareExchangeExplicit_cl20(
55
;; volatile global atomic_int* object,

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
2-
; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
2+
; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %}
33

44
; CHECK-SPIRV: OpDecorate %[[#Id:]] BuiltIn GlobalLinearId
55
; CHECK-SPIRV: %[[#Id:]] = OpVariable %[[#]]

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
; RUN: llc -O0 -mtriple=spirv64-unknown-linux %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
2-
; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
2+
; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
33

44
;; The IR was generated from the following source:
55
;; #include <CL/sycl.hpp>

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,5 @@
11
; RUN: llc -O0 -mtriple=spirv64-unknown-linux %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV
2-
; TODO: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
2+
; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %}
33

44
;; The IR was generated from the following source:
55
;; #include <CL/sycl.hpp>

0 commit comments

Comments
 (0)