Skip to content

Commit 1eda567

Browse files
MrSidimssys-ce-bb
MrSidims
authored and
sys-ce-bb
committed
Adjust Lifetime intrinsic translation after 92c55a3 (#3278)
This patch removes bitcast introduced in TypeScavenger that was previously served as W/A to solve a discrepancy between LLVM's and SPIR-V's lifetime inst definitions. Signed-off-by: Sidorov, Dmitry <[email protected]> Original commit: KhronosGroup/SPIRV-LLVM-Translator@8e829705b660a4c
1 parent 6cece96 commit 1eda567

File tree

13 files changed

+74
-153
lines changed

13 files changed

+74
-153
lines changed

‎llvm-spirv/lib/SPIRV/SPIRVReader.cpp

Lines changed: 18 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -1878,9 +1878,17 @@ Value *SPIRVToLLVM::transValueWithoutDecoration(SPIRVValue *BV, Function *F,
18781878
IRBuilder<> Builder(BB);
18791879
SPIRVWord Size = LTStart->getSize();
18801880
ConstantInt *S = nullptr;
1881+
auto *Var = transValue(LTStart->getObject(), F, BB);
1882+
Var = Var->stripPointerCasts();
18811883
if (Size)
18821884
S = Builder.getInt64(Size);
1883-
Value *Var = transValue(LTStart->getObject(), F, BB);
1885+
if (Size == 0) {
1886+
auto *Alloca = cast<AllocaInst>(Var);
1887+
if (Alloca->getAllocatedType()->isSized())
1888+
Size = M->getDataLayout().getTypeAllocSize(Alloca->getAllocatedType());
1889+
else
1890+
Size = static_cast<SPIRVWord>(-1);
1891+
}
18841892
CallInst *Start = Builder.CreateLifetimeStart(Var, S);
18851893
return mapValue(BV, Start);
18861894
}
@@ -1890,9 +1898,17 @@ Value *SPIRVToLLVM::transValueWithoutDecoration(SPIRVValue *BV, Function *F,
18901898
IRBuilder<> Builder(BB);
18911899
SPIRVWord Size = LTStop->getSize();
18921900
ConstantInt *S = nullptr;
1901+
auto *Var = transValue(LTStop->getObject(), F, BB);
1902+
Var = Var->stripPointerCasts();
18931903
if (Size)
18941904
S = Builder.getInt64(Size);
1895-
auto *Var = transValue(LTStop->getObject(), F, BB);
1905+
if (Size == 0) {
1906+
auto *Alloca = cast<AllocaInst>(Var);
1907+
if (Alloca->getAllocatedType()->isSized())
1908+
Size = M->getDataLayout().getTypeAllocSize(Alloca->getAllocatedType());
1909+
else
1910+
Size = static_cast<SPIRVWord>(-1);
1911+
}
18961912
for (const auto &I : Var->users())
18971913
if (auto *II = getLifetimeStartIntrinsic(dyn_cast<Instruction>(I)))
18981914
return mapValue(BV, Builder.CreateLifetimeEnd(II->getOperand(1), S));

‎llvm-spirv/lib/SPIRV/SPIRVTypeScavenger.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -499,6 +499,8 @@ bool SPIRVTypeScavenger::typeIntrinsicCall(
499499
break;
500500
case Intrinsic::lifetime_start:
501501
case Intrinsic::lifetime_end:
502+
// Translate the types properly.
503+
break;
502504
case Intrinsic::invariant_start:
503505
// These intrinsics were stored as i8* as typed pointers, and the SPIR-V
504506
// writer will expect these to be i8*, even if they can be any pointer

‎llvm-spirv/lib/SPIRV/SPIRVWriter.cpp

Lines changed: 5 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -4835,19 +4835,14 @@ SPIRVValue *LLVMToSPIRVBase::transIntrinsicInst(IntrinsicInst *II,
48354835
auto *PtrOp = transValue(LLVMPtrOp, BB);
48364836
if (PtrAS == SPIRAS_Private)
48374837
return BM->addLifetimeInst(OC, PtrOp, Size, BB);
4838-
// If pointer address space is Generic - cast to private first
4838+
// If pointer address space is Generic - use original allocation.
48394839
BM->getErrorLog().checkError(
48404840
PtrAS == SPIRAS_Generic, SPIRVEC_InvalidInstruction, II,
48414841
"lifetime intrinsic pointer operand must be in private or generic AS");
4842-
auto *SrcTy = PtrOp->getType();
4843-
SPIRVType *DstTy = nullptr;
4844-
if (SrcTy->isTypeUntypedPointerKHR())
4845-
DstTy = BM->addPointerType(StorageClassFunction, nullptr);
4846-
else
4847-
DstTy = BM->addPointerType(StorageClassFunction,
4848-
SrcTy->getPointerElementType());
4849-
PtrOp = BM->addUnaryInst(OpGenericCastToPtr, DstTy, PtrOp, BB);
4850-
ValueMap[LLVMPtrOp] = PtrOp;
4842+
if (PtrOp->getOpCode() == OpPtrCastToGeneric) {
4843+
auto *UI = static_cast<SPIRVUnary *>(PtrOp);
4844+
PtrOp = UI->getOperand(0);
4845+
}
48514846
return BM->addLifetimeInst(OC, PtrOp, Size, BB);
48524847
}
48534848
// We don't want to mix translation of regular code and debug info, because

‎llvm-spirv/lib/SPIRV/libSPIRV/SPIRVInstruction.h

Lines changed: 13 additions & 12 deletions
Original file line numberDiff line numberDiff line change
@@ -2504,6 +2504,18 @@ template <Op OC> class SPIRVLifetime : public SPIRVInstruction {
25042504
// Complete constructor
25052505
SPIRVLifetime(SPIRVId TheObject, SPIRVWord TheSize, SPIRVBasicBlock *TheBB)
25062506
: SPIRVInstruction(3, OC, TheBB), Object(TheObject), Size(TheSize) {
2507+
auto ObjType = getValue(Object)->getType();
2508+
// Size must be 0 if Pointer is a pointer to a non-void type or the
2509+
// Addresses capability is not being used. If Size is non-zero, it is the
2510+
// number of bytes of memory whose lifetime is starting. Its type must be an
2511+
// integer type scalar. It is treated as unsigned; if its type has
2512+
// Signedness of 1, its sign bit cannot be set.
2513+
if (!(ObjType->getPointerElementType()->isTypeVoid() ||
2514+
// (void *) is i8* in LLVM IR
2515+
ObjType->getPointerElementType()->isTypeInt(8) ||
2516+
ObjType->getPointerElementType()->isTypeUntypedPointerKHR()) ||
2517+
!Module->hasCapability(CapabilityAddresses))
2518+
Size = 0;
25072519
validate();
25082520
assert(TheBB && "Invalid BB");
25092521
}
@@ -2521,23 +2533,12 @@ template <Op OC> class SPIRVLifetime : public SPIRVInstruction {
25212533

25222534
protected:
25232535
void validate() const override {
2524-
auto ObjType = getValue(Object)->getType();
2536+
[[maybe_unused]] auto ObjType = getValue(Object)->getType();
25252537
// Type must be an OpTypePointer with Storage Class Function.
25262538
assert(ObjType->isTypePointer() && "Objects type must be a pointer");
25272539
assert(static_cast<SPIRVTypePointer *>(ObjType)->getStorageClass() ==
25282540
StorageClassFunction &&
25292541
"Invalid storage class");
2530-
// Size must be 0 if Pointer is a pointer to a non-void type or the
2531-
// Addresses capability is not being used. If Size is non-zero, it is the
2532-
// number of bytes of memory whose lifetime is starting. Its type must be an
2533-
// integer type scalar. It is treated as unsigned; if its type has
2534-
// Signedness of 1, its sign bit cannot be set.
2535-
if (!(ObjType->getPointerElementType()->isTypeVoid() ||
2536-
// (void *) is i8* in LLVM IR
2537-
ObjType->getPointerElementType()->isTypeInt(8) ||
2538-
ObjType->getPointerElementType()->isTypeUntypedPointerKHR()) ||
2539-
!Module->hasCapability(CapabilityAddresses))
2540-
assert(Size == 0 && "Size must be 0");
25412542
}
25422543
_SPIRV_DEF_ENCDEC2(Object, Size)
25432544
SPIRVId Object;

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_joint_matrix/cooperative_matrix_apply.ll

Lines changed: 0 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -58,7 +58,6 @@ entry:
5858
%ref.tmp6.i = alloca float, align 4
5959
%__SYCLKernel = alloca %class.anon.0, align 8
6060
%__SYCLKernel.ascast = addrspacecast ptr %__SYCLKernel to ptr addrspace(4)
61-
call void @llvm.lifetime.start.p0(i64 64, ptr nonnull %__SYCLKernel)
6261
%agg.tmp.sroa.0.sroa.0.0.copyload = load i64, ptr %_arg_accA1, align 8
6362
%agg.tmp.sroa.0.sroa.2.0._arg_accA1.ascast.sroa_idx = getelementptr inbounds i8, ptr %_arg_accA1, i64 8
6463
%agg.tmp.sroa.0.sroa.2.0.copyload = load i64, ptr %agg.tmp.sroa.0.sroa.2.0._arg_accA1.ascast.sroa_idx, align 8
@@ -95,20 +94,14 @@ entry:
9594
%sub.i = sub nsw i64 %2, %4
9695
%cmp.i12 = icmp ult i64 %5, 2147483648
9796
%sub5.i = sub nsw i64 %3, %5
98-
call void @llvm.lifetime.start.p0(i64 4, ptr nonnull %ref.tmp6.i)
9997
store float 5.000000e+00, ptr %ref.tmp6.i, align 4
10098
%call.i.i = call spir_func noundef zeroext i16 @__devicelib_ConvertFToBF16INTEL(ptr addrspace(4) noundef align 4 dereferenceable(4) %ref.tmp6.ascast.i)
101-
call void @llvm.lifetime.start.p0(i64 2, ptr nonnull %agg.tmp.i17)
10299
store i16 %call.i.i, ptr %agg.tmp.i17, align 2
103100
%call.i18 = call spir_func noundef target("spirv.CooperativeMatrixKHR", i16, 3, 8, 16, 0) @_Z26__spirv_CompositeConstruct(ptr noundef nonnull byval(%"class.sycl::_V1::ext::oneapi::bfloat16") align 2 %agg.tmp.i17)
104-
call void @llvm.lifetime.end.p0(i64 2, ptr nonnull %agg.tmp.i17)
105-
call void @llvm.lifetime.end.p0(i64 4, ptr nonnull %ref.tmp6.i)
106101
%lambda.i = getelementptr inbounds %class.anon.0, ptr addrspace(4) %__SYCLKernel.ascast, i64 0, i32 1
107102
%ref.tmp.ascast.i21 = addrspacecast ptr %ref.tmp.i20 to ptr addrspace(4)
108-
call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %ref.tmp.i20)
109103
store ptr addrspace(4) %lambda.i, ptr %ref.tmp.i20, align 8
110104
%call.i22 = call spir_func noundef target("spirv.CooperativeMatrixKHR", i16, 3, 8, 16, 0) @_Z43__spirv_CooperativeMatrixApplyFunctionINTEL(ptr addrspace(4) noundef align 8 dereferenceable(8) %ref.tmp.ascast.i21, target("spirv.CooperativeMatrixKHR", i16, 3, 8, 16, 0) noundef %call.i18)
111-
call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %ref.tmp.i20)
112105
%6 = load ptr addrspace(1), ptr %0, align 8
113106
%7 = load i64, ptr %__SYCLKernel, align 8
114107
%8 = load i64, ptr %arrayidx.i29.i.i.i.i, align 8
@@ -122,16 +115,9 @@ entry:
122115
%div14.i = and i64 %sub5.i, -16
123116
%add.ptr.i44 = getelementptr inbounds %"class.sycl::_V1::ext::oneapi::bfloat16", ptr addrspace(1) %add.ptr.i43, i64 %div14.i
124117
call spir_func void @_Z33__spirv_CooperativeMatrixStoreKHRPU3AS4iPU3AS144__spirv_CooperativeMatrixKHR__uint_3_12_12_3ili(ptr addrspace(1) noundef %add.ptr.i44, target("spirv.CooperativeMatrixKHR", i16, 3, 8, 16, 0) noundef %call.i22, i32 noundef 0, i64 noundef 0)
125-
call void @llvm.lifetime.end.p0(i64 64, ptr nonnull %__SYCLKernel)
126118
ret void
127119
}
128120

129-
; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
130-
declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none))
131-
132-
; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
133-
declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none))
134-
135121
; Function Attrs: convergent nounwind
136122
declare dso_local spir_func noundef target("spirv.CooperativeMatrixKHR", i16, 3, 8, 16, 0) @_Z26__spirv_CompositeConstruct(ptr noundef byval(%"class.sycl::_V1::ext::oneapi::bfloat16") align 2) local_unnamed_addr
137123

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_joint_matrix/cooperative_matrix_checked.ll

Lines changed: 0 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -84,7 +84,6 @@ entry:
8484
%sub.i = sub nsw i64 %1, %4
8585
%cmp.i58.i = icmp ult i64 %5, 2147483648
8686
%sub5.i = sub nsw i64 %2, %5
87-
call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %sub_c.sroa.0.i)
8887
%call.i.i = tail call spir_func noundef target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) @_Z46__spirv_CooperativeMatrixConstructCheckedINTEL(i32 noundef 4, i32 noundef 4, i32 noundef 12, i32 noundef 12, i32 noundef %_arg_Initvalue) #4
8988
store target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) %call.i.i, ptr %sub_c.sroa.0.i, align 8
9089
%mul.i = mul nsw i64 %sub.i, 12
@@ -118,13 +117,11 @@ for.body.i: ; preds = %for.cond.i
118117
%add.ptr.i111.i = getelementptr i8, ptr addrspace(1) %add.ptr.i108140.i, i64 %mul23.i
119118
%call.ascast.i72.i = addrspacecast ptr addrspace(1) %add.ptr.i111.i to ptr addrspace(4)
120119
%call1.i73.i = tail call spir_func noundef target("spirv.CooperativeMatrixKHR", i8, 2, 48, 12, 1) @_Z41__spirv_CooperativeMatrixLoadCheckedINTEL_2(ptr addrspace(4) noundef %call.ascast.i72.i, i32 noundef 0, i32 noundef 0, i32 noundef 0, i32 noundef 48, i32 noundef 12, i64 noundef %mul22.i) #4
121-
call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %ref.tmp29.sroa.0.i)
122120
%sub_c.sroa.0.i.0.sub_c.sroa.0.i.0.sub_c.sroa.0.0.sub_c.sroa.0.0.sub_c.sroa.0.0.125.i = load target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2), ptr %sub_c.sroa.0.i, align 8
123121
%call.i77.i = tail call spir_func noundef target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) @_Z34__spirv_CooperativeMatrixMulAddKHR(target("spirv.CooperativeMatrixKHR", i8, 3, 12, 48, 0) noundef %call1.i.i, target("spirv.CooperativeMatrixKHR", i8, 2, 48, 12, 1) noundef %call1.i73.i, target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) noundef %sub_c.sroa.0.i.0.sub_c.sroa.0.i.0.sub_c.sroa.0.0.sub_c.sroa.0.0.sub_c.sroa.0.0.125.i, i32 noundef 12) #4
124122
store target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) %call.i77.i, ptr %ref.tmp29.sroa.0.i, align 8
125123
%ref.tmp29.sroa.0.i.0.ref.tmp29.sroa.0.i.0.ref.tmp29.sroa.0.0.ref.tmp29.sroa.0.0.ref.tmp29.sroa.0.0..i = load i64, ptr %ref.tmp29.sroa.0.i, align 8
126124
store i64 %ref.tmp29.sroa.0.i.0.ref.tmp29.sroa.0.i.0.ref.tmp29.sroa.0.0.ref.tmp29.sroa.0.0.ref.tmp29.sroa.0.0..i, ptr %sub_c.sroa.0.i, align 8
127-
call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %ref.tmp29.sroa.0.i)
128125
%add.i = add nuw nsw i32 %k.0.i, 1
129126
br label %for.cond.i
130127

@@ -136,7 +133,6 @@ _ZZZ15matrix_multiplyIiaLm24ELm96ELm24ELm96ELm24ELm24EEvR10big_matrixIT_XT5_EXT6
136133
%call.ascast.i.i = addrspacecast ptr addrspace(1) %add.ptr.i81.i to ptr addrspace(4)
137134
%sub_c.sroa.0.i.0.sub_c.sroa.0.i.0.sub_c.sroa.0.0.sub_c.sroa.0.0.sub_c.sroa.0.0..i = load target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2), ptr %sub_c.sroa.0.i, align 8
138135
tail call spir_func void @_Z42__spirv_CooperativeMatrixStoreCheckedINTEL(ptr addrspace(4) noundef %call.ascast.i.i, i32 noundef 0, i32 noundef 0, target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) noundef %sub_c.sroa.0.i.0.sub_c.sroa.0.i.0.sub_c.sroa.0.0.sub_c.sroa.0.0.sub_c.sroa.0.0..i, i32 noundef 0, i32 noundef 12, i32 noundef 12, i64 noundef %_arg_N, i32 noundef 1) #4
139-
call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %sub_c.sroa.0.i)
140136
ret void
141137
}
142138

@@ -157,12 +153,6 @@ declare dso_local spir_func noundef target("spirv.CooperativeMatrixKHR", i32, 3,
157153
; Function Attrs: convergent
158154
declare dso_local spir_func void @_Z42__spirv_CooperativeMatrixStoreCheckedINTEL(ptr addrspace(4) noundef, i32 noundef, i32 noundef, target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) noundef, i32 noundef, i32 noundef, i32 noundef, i64 noundef, i32 noundef) local_unnamed_addr #2
159155

160-
; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
161-
declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #3
162-
163-
; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
164-
declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #3
165-
166156
attributes #0 = { convergent norecurse "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="matrix-int8-test.cpp" "uniform-work-group-size"="true" }
167157
attributes #1 = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: readwrite) }
168158
attributes #2 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }

‎llvm-spirv/test/extensions/INTEL/SPV_INTEL_joint_matrix/cooperative_matrix_prefetch.ll

Lines changed: 0 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -87,7 +87,6 @@ entry:
8787
%sub.i = sub nsw i64 %1, %4
8888
%cmp.i58.i = icmp ult i64 %5, 2147483648
8989
%sub5.i = sub nsw i64 %2, %5
90-
call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %sub_c.sroa.0.i)
9190
%call.i.i = tail call spir_func noundef target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) @_Z26__spirv_CompositeConstruct(i32 noundef 0) #4
9291
store target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) %call.i.i, ptr %sub_c.sroa.0.i, align 8
9392
%mul.i = mul nsw i64 %sub.i, 12
@@ -123,13 +122,11 @@ for.body.i: ; preds = %for.cond.i
123122
%call.ascast.i72.i = addrspacecast ptr addrspace(1) %add.ptr.i111.i to ptr addrspace(4)
124123
tail call spir_func void @_Z38__spirv_CooperativeMatrixPrefetchINTEL(ptr addrspace(4) noundef %call.ascast.i72.i, i32 noundef 12, i32 noundef 48, i32 noundef 0, i32 noundef 0, i64 noundef %mul22.i)
125124
%call1.i73.i = tail call spir_func noundef target("spirv.CooperativeMatrixKHR", i8, 2, 48, 12, 1) @_Z32__spirv_CooperativeMatrixLoadKHR_2(ptr addrspace(4) noundef %call.ascast.i72.i, i32 noundef 0, i64 noundef %mul22.i) #4
126-
call void @llvm.lifetime.start.p0(i64 8, ptr nonnull %ref.tmp29.sroa.0.i)
127125
%sub_c.sroa.0.i.0.sub_c.sroa.0.i.0.sub_c.sroa.0.0.sub_c.sroa.0.0.sub_c.sroa.0.0.125.i = load target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2), ptr %sub_c.sroa.0.i, align 8
128126
%call.i77.i = tail call spir_func noundef target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) @_Z34__spirv_CooperativeMatrixMulAddKHR(target("spirv.CooperativeMatrixKHR", i8, 3, 12, 48, 0) noundef %call1.i.i, target("spirv.CooperativeMatrixKHR", i8, 2, 48, 12, 1) noundef %call1.i73.i, target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) noundef %sub_c.sroa.0.i.0.sub_c.sroa.0.i.0.sub_c.sroa.0.0.sub_c.sroa.0.0.sub_c.sroa.0.0.125.i, i32 noundef 12) #4
129127
store target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) %call.i77.i, ptr %ref.tmp29.sroa.0.i, align 8
130128
%ref.tmp29.sroa.0.i.0.ref.tmp29.sroa.0.i.0.ref.tmp29.sroa.0.0.ref.tmp29.sroa.0.0.ref.tmp29.sroa.0.0..i = load i64, ptr %ref.tmp29.sroa.0.i, align 8
131129
store i64 %ref.tmp29.sroa.0.i.0.ref.tmp29.sroa.0.i.0.ref.tmp29.sroa.0.0.ref.tmp29.sroa.0.0.ref.tmp29.sroa.0.0..i, ptr %sub_c.sroa.0.i, align 8
132-
call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %ref.tmp29.sroa.0.i)
133130
%add.i = add nuw nsw i32 %k.0.i, 1
134131
br label %for.cond.i
135132

@@ -141,7 +138,6 @@ _ZZZ15matrix_multiplyIiaLm24ELm96ELm24ELm96ELm24ELm24EEvR10big_matrixIT_XT5_EXT6
141138
%call.ascast.i.i = addrspacecast ptr addrspace(1) %add.ptr.i81.i to ptr addrspace(4)
142139
%sub_c.sroa.0.i.0.sub_c.sroa.0.i.0.sub_c.sroa.0.0.sub_c.sroa.0.0.sub_c.sroa.0.0..i = load target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2), ptr %sub_c.sroa.0.i, align 8
143140
tail call spir_func void @_Z33__spirv_CooperativeMatrixStoreKHR(ptr addrspace(4) noundef %call.ascast.i.i, target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) noundef %sub_c.sroa.0.i.0.sub_c.sroa.0.i.0.sub_c.sroa.0.0.sub_c.sroa.0.0.sub_c.sroa.0.0..i, i32 noundef 0, i64 noundef %_arg_N, i32 noundef 1) #4
144-
call void @llvm.lifetime.end.p0(i64 8, ptr nonnull %sub_c.sroa.0.i)
145141
ret void
146142
}
147143

@@ -165,12 +161,6 @@ declare dso_local spir_func noundef target("spirv.CooperativeMatrixKHR", i32, 3,
165161
; Function Attrs: convergent
166162
declare dso_local spir_func void @_Z33__spirv_CooperativeMatrixStoreKHR(ptr addrspace(4) noundef, target("spirv.CooperativeMatrixKHR", i32, 3, 12, 12, 2) noundef, i32 noundef, i64 noundef, i32 noundef) local_unnamed_addr #2
167163

168-
; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
169-
declare void @llvm.lifetime.start.p0(i64 immarg, ptr captures(none)) #3
170-
171-
; Function Attrs: nocallback nofree nosync nounwind willreturn memory(argmem: readwrite)
172-
declare void @llvm.lifetime.end.p0(i64 immarg, ptr captures(none)) #3
173-
174164
attributes #0 = { convergent norecurse "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "sycl-module-id"="matrix-int8-test.cpp" "uniform-work-group-size"="true" }
175165
attributes #1 = { nocallback nofree nosync nounwind willreturn memory(inaccessiblemem: readwrite) }
176166
attributes #2 = { convergent "frame-pointer"="all" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }

0 commit comments

Comments
 (0)