From 9a75467d7e4cd185c7604007751192daea7ec0e5 Mon Sep 17 00:00:00 2001 From: "Levytskyy, Vyacheslav" Date: Wed, 5 Jun 2024 06:22:01 -0700 Subject: [PATCH 1/5] improve type inference --- llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp | 20 ++- llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp | 150 +++++++++++------- llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h | 16 +- llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp | 6 +- .../test/CodeGen/SPIRV/event-wait-ptr-type.ll | 16 +- ...Intrinsics-no-duplicate-spv_assign_type.ll | 4 +- 6 files changed, 142 insertions(+), 70 deletions(-) diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp index 424087f361a6a..9b9b8f7cbc089 100644 --- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp @@ -24,6 +24,13 @@ #define DEBUG_TYPE "spirv-builtins" namespace llvm { + +// Defined in SPIRVPreLegalizer.cpp. +extern Register insertAssignInstr(Register Reg, Type *Ty, SPIRVType *SpirvTy, + SPIRVGlobalRegistry *GR, + MachineIRBuilder &MIB, + MachineRegisterInfo &MRI); + namespace SPIRV { #define GET_BuiltinGroup_DECL #include "SPIRVGenTables.inc" @@ -1451,11 +1458,22 @@ static bool generateImageSizeQueryInst(const SPIRV::IncomingCall *Call, Component == 3 ? NumActualRetComponents - 1 : Component; assert(ExtractedComposite < NumActualRetComponents && "Invalid composite index!"); + Register TypeReg = GR->getSPIRVTypeID(Call->ReturnType); + SPIRVType *NewType = nullptr; + if (QueryResultType->getOpcode() == SPIRV::OpTypeVector) { + Register NewTypeReg = QueryResultType->getOperand(1).getReg(); + if (TypeReg != NewTypeReg && + (NewType = GR->getSPIRVTypeForVReg(NewTypeReg)) != nullptr) + TypeReg = NewTypeReg; + } MIRBuilder.buildInstr(SPIRV::OpCompositeExtract) .addDef(Call->ReturnRegister) - .addUse(GR->getSPIRVTypeID(Call->ReturnType)) + .addUse(TypeReg) .addUse(QueryResult) .addImm(ExtractedComposite); + if (NewType != nullptr) + insertAssignInstr(Call->ReturnRegister, nullptr, NewType, GR, MIRBuilder, + MIRBuilder.getMF().getRegInfo()); } else { // More than 1 component is expected, fill a new vector. auto MIB = MIRBuilder.buildInstr(SPIRV::OpVectorShuffle) diff --git a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp index 5ef0be1cab722..696706258ec40 100644 --- a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp @@ -50,6 +50,7 @@ void initializeSPIRVEmitIntrinsicsPass(PassRegistry &); } // namespace llvm namespace { + class SPIRVEmitIntrinsics : public ModulePass, public InstVisitor { @@ -61,9 +62,6 @@ class SPIRVEmitIntrinsics DenseMap AggrConstTypes; DenseSet AggrStores; - // a registry of created Intrinsic::spv_assign_ptr_type instructions - DenseMap AssignPtrTypeInstr; - // deduce element type of untyped pointers Type *deduceElementType(Value *I); Type *deduceElementTypeHelper(Value *I); @@ -98,14 +96,16 @@ class SPIRVEmitIntrinsics return B.CreateIntrinsic(IntrID, {Types}, Args); } + void buildAssignType(IRBuilder<> &B, Type *ElemTy, Value *Arg); void buildAssignPtr(IRBuilder<> &B, Type *ElemTy, Value *Arg); + void updateAssignType(CallInst *AssignCI, Value *Arg, Value *OfType); void replaceMemInstrUses(Instruction *Old, Instruction *New, IRBuilder<> &B); void processInstrAfterVisit(Instruction *I, IRBuilder<> &B); void insertAssignPtrTypeIntrs(Instruction *I, IRBuilder<> &B); void insertAssignTypeIntrs(Instruction *I, IRBuilder<> &B); - void insertAssignTypeInstrForTargetExtTypes(TargetExtType *AssignedType, - Value *V, IRBuilder<> &B); + void insertAssignPtrTypeTargetExt(TargetExtType *AssignedType, Value *V, + IRBuilder<> &B); void replacePointerOperandWithPtrCast(Instruction *I, Value *Pointer, Type *ExpectedElementType, unsigned OperandToReplace, @@ -218,15 +218,39 @@ static inline void reportFatalOnTokenType(const Instruction *I) { false); } +void SPIRVEmitIntrinsics::buildAssignType(IRBuilder<> &B, Type *Ty, + Value *Arg) { + Value *OfType = PoisonValue::get(Ty); + CallInst *AssignCI = buildIntrWithMD(Intrinsic::spv_assign_type, + {Arg->getType()}, OfType, Arg, {}, B); + GR->addAssignPtrTypeInstr(Arg, AssignCI); +} + void SPIRVEmitIntrinsics::buildAssignPtr(IRBuilder<> &B, Type *ElemTy, Value *Arg) { - CallInst *AssignPtrTyCI = - buildIntrWithMD(Intrinsic::spv_assign_ptr_type, {Arg->getType()}, - Constant::getNullValue(ElemTy), Arg, - {B.getInt32(getPointerAddressSpace(Arg->getType()))}, B); + Value *OfType = PoisonValue::get(ElemTy); + CallInst *AssignPtrTyCI = buildIntrWithMD( + Intrinsic::spv_assign_ptr_type, {Arg->getType()}, OfType, Arg, + {B.getInt32(getPointerAddressSpace(Arg->getType()))}, B); GR->addDeducedElementType(AssignPtrTyCI, ElemTy); GR->addDeducedElementType(Arg, ElemTy); - AssignPtrTypeInstr[Arg] = AssignPtrTyCI; + GR->addAssignPtrTypeInstr(Arg, AssignPtrTyCI); +} + +void SPIRVEmitIntrinsics::updateAssignType(CallInst *AssignCI, Value *Arg, + Value *OfType) { + LLVMContext &Ctx = Arg->getContext(); + AssignCI->setArgOperand( + 1, MetadataAsValue::get( + Ctx, MDNode::get(Ctx, ValueAsMetadata::getConstant(OfType)))); + if (cast(AssignCI)->getIntrinsicID() != + Intrinsic::spv_assign_ptr_type) + return; + + // update association with the pointee type + Type *ElemTy = OfType->getType(); + GR->addDeducedElementType(AssignCI, ElemTy); + GR->addDeducedElementType(Arg, ElemTy); } // Set element pointer type to the given value of ValueTy and tries to @@ -513,19 +537,16 @@ void SPIRVEmitIntrinsics::deduceOperandElementType(Instruction *I) { if (!Ty) { GR->addDeducedElementType(Op, KnownElemTy); // check if there is existing Intrinsic::spv_assign_ptr_type instruction - auto It = AssignPtrTypeInstr.find(Op); - if (It == AssignPtrTypeInstr.end()) { + CallInst *AssignCI = GR->findAssignPtrTypeInstr(Op); + if (AssignCI == nullptr) { Instruction *User = dyn_cast(Op->use_begin()->get()); setInsertPointSkippingPhis(B, User ? User->getNextNode() : I); CallInst *CI = buildIntrWithMD(Intrinsic::spv_assign_ptr_type, {OpTy}, OpTyVal, Op, {B.getInt32(getPointerAddressSpace(OpTy))}, B); - AssignPtrTypeInstr[Op] = CI; + GR->addAssignPtrTypeInstr(Op, CI); } else { - It->second->setArgOperand( - 1, - MetadataAsValue::get( - Ctx, MDNode::get(Ctx, ValueAsMetadata::getConstant(OpTyVal)))); + updateAssignType(AssignCI, Op, OpTyVal); } } else { if (auto *OpI = dyn_cast(Op)) { @@ -559,7 +580,9 @@ void SPIRVEmitIntrinsics::replaceMemInstrUses(Instruction *Old, if (isAssignTypeInstr(U)) { B.SetInsertPoint(U); SmallVector Args = {New, U->getOperand(1)}; - B.CreateIntrinsic(Intrinsic::spv_assign_type, {New->getType()}, Args); + CallInst *AssignCI = + B.CreateIntrinsic(Intrinsic::spv_assign_type, {New->getType()}, Args); + GR->addAssignPtrTypeInstr(New, AssignCI); U->eraseFromParent(); } else if (isMemInstrToReplace(U) || isa(U) || isa(U)) { @@ -751,33 +774,39 @@ Instruction *SPIRVEmitIntrinsics::visitBitCastInst(BitCastInst &I) { return NewI; } -void SPIRVEmitIntrinsics::insertAssignTypeInstrForTargetExtTypes( +void SPIRVEmitIntrinsics::insertAssignPtrTypeTargetExt( TargetExtType *AssignedType, Value *V, IRBuilder<> &B) { - // Do not emit spv_assign_type if the V is of the AssignedType already. - if (V->getType() == AssignedType) - return; + Type *VTy = V->getType(); - // Do not emit spv_assign_type if there is one already targetting V. If the - // found spv_assign_type assigns a type different than AssignedType, report an - // error. Builtin types cannot be redeclared or casted. - for (auto User : V->users()) { - auto *II = dyn_cast(User); - if (!II || II->getIntrinsicID() != Intrinsic::spv_assign_type) - continue; + // A couple of sanity checks. + assert(isPointerTy(VTy) && "Expect a pointer type!"); + if (auto PType = dyn_cast(VTy)) + if (PType->getElementType() != AssignedType) + report_fatal_error("Unexpected pointer element type!"); - MetadataAsValue *VMD = cast(II->getOperand(1)); - Type *BuiltinType = - dyn_cast(VMD->getMetadata())->getType(); - if (BuiltinType != AssignedType) - report_fatal_error("Type mismatch " + BuiltinType->getTargetExtName() + - "/" + AssignedType->getTargetExtName() + - " for value " + V->getName(), - false); + CallInst *AssignCI = GR->findAssignPtrTypeInstr(V); + if (!AssignCI) { + buildAssignType(B, AssignedType, V); return; } - Constant *Const = UndefValue::get(AssignedType); - buildIntrWithMD(Intrinsic::spv_assign_type, {V->getType()}, Const, V, {}, B); + Type *CurrentType = + dyn_cast( + cast(AssignCI->getOperand(1))->getMetadata()) + ->getType(); + if (CurrentType == AssignedType) + return; + + // Builtin types cannot be redeclared or casted. + if (CurrentType->isTargetExtTy()) + report_fatal_error("Type mismatch " + CurrentType->getTargetExtName() + + "/" + AssignedType->getTargetExtName() + + " for value " + V->getName(), + false); + + // Our previous guess about the type seems to be wrong, let's update + // inferred type according to a new, more precise type information. + updateAssignType(AssignCI, V, PoisonValue::get(AssignedType)); } void SPIRVEmitIntrinsics::replacePointerOperandWithPtrCast( @@ -850,7 +879,7 @@ void SPIRVEmitIntrinsics::replacePointerOperandWithPtrCast( ExpectedElementTypeConst, Pointer, {B.getInt32(AddressSpace)}, B); GR->addDeducedElementType(CI, ExpectedElementType); GR->addDeducedElementType(Pointer, ExpectedElementType); - AssignPtrTypeInstr[Pointer] = CI; + GR->addAssignPtrTypeInstr(Pointer, CI); return; } @@ -929,8 +958,7 @@ void SPIRVEmitIntrinsics::insertPtrCastOrAssignTypeInstr(Instruction *I, for (unsigned OpIdx = 0; OpIdx < CI->arg_size(); OpIdx++) { Value *ArgOperand = CI->getArgOperand(OpIdx); - if (!isa(ArgOperand->getType()) && - !isa(ArgOperand->getType())) + if (!isPointerTy(ArgOperand->getType())) continue; // Constants (nulls/undefs) are handled in insertAssignPtrTypeIntrs() @@ -952,8 +980,8 @@ void SPIRVEmitIntrinsics::insertPtrCastOrAssignTypeInstr(Instruction *I, continue; if (ExpectedType->isTargetExtTy()) - insertAssignTypeInstrForTargetExtTypes(cast(ExpectedType), - ArgOperand, B); + insertAssignPtrTypeTargetExt(cast(ExpectedType), + ArgOperand, B); else replacePointerOperandWithPtrCast(CI, ArgOperand, ExpectedType, OpIdx, B); } @@ -1145,7 +1173,7 @@ void SPIRVEmitIntrinsics::insertAssignPtrTypeIntrs(Instruction *I, CallInst *CI = buildIntrWithMD(Intrinsic::spv_assign_ptr_type, {I->getType()}, EltTyConst, I, {B.getInt32(AddressSpace)}, B); GR->addDeducedElementType(CI, ElemTy); - AssignPtrTypeInstr[I] = CI; + GR->addAssignPtrTypeInstr(I, CI); } void SPIRVEmitIntrinsics::insertAssignTypeIntrs(Instruction *I, @@ -1164,20 +1192,32 @@ void SPIRVEmitIntrinsics::insertAssignTypeIntrs(Instruction *I, TypeToAssign = It->second; } } - Constant *Const = UndefValue::get(TypeToAssign); - buildIntrWithMD(Intrinsic::spv_assign_type, {Ty}, Const, I, {}, B); + buildAssignType(B, TypeToAssign, I); } for (const auto &Op : I->operands()) { if (isa(Op) || isa(Op) || // Check GetElementPtrConstantExpr case. (isa(Op) && isa(Op))) { setInsertPointSkippingPhis(B, I); - if (isa(Op) && Op->getType()->isAggregateType()) - buildIntrWithMD(Intrinsic::spv_assign_type, {B.getInt32Ty()}, Op, - UndefValue::get(B.getInt32Ty()), {}, B); - else if (!isa(Op)) - buildIntrWithMD(Intrinsic::spv_assign_type, {Op->getType()}, Op, Op, {}, - B); + Type *OpTy = Op->getType(); + if (isa(Op) && OpTy->isAggregateType()) { + CallInst *AssignCI = + buildIntrWithMD(Intrinsic::spv_assign_type, {B.getInt32Ty()}, Op, + UndefValue::get(B.getInt32Ty()), {}, B); + GR->addAssignPtrTypeInstr(Op, AssignCI); + } else if (!isa(Op)) { + Type *OpTy = Op->getType(); + if (auto PType = dyn_cast(OpTy)) { + buildAssignPtr(B, PType->getElementType(), Op); + } else if (isPointerTy(OpTy)) { + Type *ElemTy = GR->findDeducedElementType(Op); + buildAssignPtr(B, ElemTy ? ElemTy : deduceElementType(Op), Op); + } else { + CallInst *AssignCI = buildIntrWithMD(Intrinsic::spv_assign_type, + {OpTy}, Op, Op, {}, B); + GR->addAssignPtrTypeInstr(Op, AssignCI); + } + } } } } @@ -1368,14 +1408,12 @@ bool SPIRVEmitIntrinsics::runOnFunction(Function &Func) { continue; insertAssignPtrTypeIntrs(I, B); + deduceOperandElementType(I); insertAssignTypeIntrs(I, B); insertPtrCastOrAssignTypeInstr(I, B); insertSpirvDecorations(I, B); } - for (auto &I : instructions(Func)) - deduceOperandElementType(&I); - for (auto *I : Worklist) { TrackConstants = true; if (!I->getType()->isVoidTy() || isa(I)) diff --git a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h index 55979ba403a0e..0103fb8214341 100644 --- a/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h +++ b/llvm/lib/Target/SPIRV/SPIRVGlobalRegistry.h @@ -72,8 +72,11 @@ class SPIRVGlobalRegistry { // untyped pointers. DenseMap DeducedElTys; // Maps composite values to deduced types where untyped pointers are replaced - // with typed ones + // with typed ones. DenseMap DeducedNestedTys; + // Maps values to "assign type" calls, thus being a registry of created + // Intrinsic::spv_assign_ptr_type instructions. + DenseMap AssignPtrTypeInstr; // Add a new OpTypeXXX instruction without checking for duplicates. SPIRVType *createSPIRVType(const Type *Type, MachineIRBuilder &MIRBuilder, @@ -148,6 +151,17 @@ class SPIRVGlobalRegistry { return It == FunResPointerTypes.end() ? nullptr : It->second; } + // A registry of "assign type" records: + // - Add a record. + void addAssignPtrTypeInstr(Value *Val, CallInst *AssignPtrTyCI) { + AssignPtrTypeInstr[Val] = AssignPtrTyCI; + } + // - Find a record. + CallInst *findAssignPtrTypeInstr(const Value *Val) { + auto It = AssignPtrTypeInstr.find(Val); + return It == AssignPtrTypeInstr.end() ? nullptr : It->second; + } + // Deduced element types of untyped pointers and composites: // - Add a record to the map of deduced element types. void addDeducedElementType(Value *Val, Type *Ty) { DeducedElTys[Val] = Ty; } diff --git a/llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp b/llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp index 3d536085b78aa..a0a253c23b1e8 100644 --- a/llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVPreLegalizer.cpp @@ -417,7 +417,8 @@ generateAssignInstrs(MachineFunction &MF, SPIRVGlobalRegistry *GR, MachineInstr *Def = MRI.getVRegDef(Reg); assert(Def && "Expecting an instruction that defines the register"); // G_GLOBAL_VALUE already has type info. - if (Def->getOpcode() != TargetOpcode::G_GLOBAL_VALUE) + if (Def->getOpcode() != TargetOpcode::G_GLOBAL_VALUE && + Def->getOpcode() != SPIRV::ASSIGN_TYPE) insertAssignInstr(Reg, nullptr, AssignedPtrType, GR, MIB, MF.getRegInfo()); ToErase.push_back(&MI); @@ -427,7 +428,8 @@ generateAssignInstrs(MachineFunction &MF, SPIRVGlobalRegistry *GR, MachineInstr *Def = MRI.getVRegDef(Reg); assert(Def && "Expecting an instruction that defines the register"); // G_GLOBAL_VALUE already has type info. - if (Def->getOpcode() != TargetOpcode::G_GLOBAL_VALUE) + if (Def->getOpcode() != TargetOpcode::G_GLOBAL_VALUE && + Def->getOpcode() != SPIRV::ASSIGN_TYPE) insertAssignInstr(Reg, Ty, nullptr, GR, MIB, MF.getRegInfo()); ToErase.push_back(&MI); } else if (MIOp == TargetOpcode::G_CONSTANT || diff --git a/llvm/test/CodeGen/SPIRV/event-wait-ptr-type.ll b/llvm/test/CodeGen/SPIRV/event-wait-ptr-type.ll index d6fb70bb59a7e..ec9afc789944d 100644 --- a/llvm/test/CodeGen/SPIRV/event-wait-ptr-type.ll +++ b/llvm/test/CodeGen/SPIRV/event-wait-ptr-type.ll @@ -4,16 +4,16 @@ ; RUN: llc -O0 -mtriple=spirv32-unknown-unknown %s -o - | FileCheck %s ; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv32-unknown-unknown %s -o - -filetype=obj | spirv-val %} -; CHECK: %[[#EventTy:]] = OpTypeEvent -; CHECK: %[[#StructEventTy:]] = OpTypeStruct %[[#EventTy]] -; CHECK: %[[#GenPtrStructEventTy:]] = OpTypePointer Generic %[[#StructEventTy]] -; CHECK: %[[#FunPtrStructEventTy:]] = OpTypePointer Function %[[#StructEventTy]] -; CHECK: %[[#GenPtrEventTy:]] = OpTypePointer Generic %[[#EventTy:]] +; CHECK-DAG: %[[#EventTy:]] = OpTypeEvent +; CHECK-DAG: %[[#StructEventTy:]] = OpTypeStruct %[[#EventTy]] +; CHECK-DAG: %[[#FunPtrStructEventTy:]] = OpTypePointer Function %[[#StructEventTy]] +; CHECK-DAG: %[[#GenPtrEventTy:]] = OpTypePointer Generic %[[#EventTy]] +; CHECK-DAG: %[[#FunPtrEventTy:]] = OpTypePointer Function %[[#EventTy]] ; CHECK: OpFunction ; CHECK: %[[#Var:]] = OpVariable %[[#FunPtrStructEventTy]] Function -; CHECK-NEXT: %[[#AddrspacecastVar:]] = OpPtrCastToGeneric %[[#GenPtrStructEventTy]] %[[#Var]] -; CHECK-NEXT: %[[#BitcastVar:]] = OpBitcast %[[#GenPtrEventTy]] %[[#AddrspacecastVar]] -; CHECK-NEXT: OpGroupWaitEvents %[[#]] %[[#]] %[[#BitcastVar]] +; CHECK-NEXT: %[[#FunEvent:]] = OpBitcast %[[#FunPtrEventTy]] %[[#Var]] +; CHECK-NEXT: %[[#GenEvent:]] = OpPtrCastToGeneric %[[#GenPtrEventTy]] %[[#FunEvent]] +; CHECK-NEXT: OpGroupWaitEvents %[[#]] %[[#]] %[[#GenEvent]] %"class.sycl::_V1::device_event" = type { target("spirv.Event") } diff --git a/llvm/test/CodeGen/SPIRV/passes/SPIRVEmitIntrinsics-no-duplicate-spv_assign_type.ll b/llvm/test/CodeGen/SPIRV/passes/SPIRVEmitIntrinsics-no-duplicate-spv_assign_type.ll index 7056b9cb1230d..9db4f26a27d4f 100644 --- a/llvm/test/CodeGen/SPIRV/passes/SPIRVEmitIntrinsics-no-duplicate-spv_assign_type.ll +++ b/llvm/test/CodeGen/SPIRV/passes/SPIRVEmitIntrinsics-no-duplicate-spv_assign_type.ll @@ -3,9 +3,9 @@ ; CHECK: *** IR Dump After SPIRV emit intrinsics (emit-intrinsics) *** define spir_kernel void @test(ptr addrspace(1) %srcimg) { -; CHECK: call void @llvm.spv.assign.type.p1(ptr addrspace(1) %srcimg, metadata target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) undef) +; CHECK: call void @llvm.spv.assign.type.p1(ptr addrspace(1) %srcimg, metadata target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) poison) %call1 = call spir_func <2 x i32> @_Z13get_image_dim14ocl_image2d_ro(ptr addrspace(1) %srcimg) -; CHECK-NOT: call void @llvm.spv.assign.type.p1(ptr addrspace(1) %srcimg, metadata target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) undef) +; CHECK-NOT: call void @llvm.spv.assign.type.p1(ptr addrspace(1) %srcimg, metadata target("spirv.Image", void, 1, 0, 0, 0, 0, 0, 0) poison) %call2 = call spir_func <2 x i32> @_Z13get_image_dim14ocl_image2d_ro(ptr addrspace(1) %srcimg) ret void ; CHECK: } From 4f40eb8c3eec26cb6905471858ecb33facb36522 Mon Sep 17 00:00:00 2001 From: "Levytskyy, Vyacheslav" Date: Wed, 5 Jun 2024 06:30:17 -0700 Subject: [PATCH 2/5] add validation --- llvm/test/CodeGen/SPIRV/transcoding/check_ro_qualifier.ll | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/llvm/test/CodeGen/SPIRV/transcoding/check_ro_qualifier.ll b/llvm/test/CodeGen/SPIRV/transcoding/check_ro_qualifier.ll index 824ca1b2d6924..6f61aba23a46f 100644 --- a/llvm/test/CodeGen/SPIRV/transcoding/check_ro_qualifier.ll +++ b/llvm/test/CodeGen/SPIRV/transcoding/check_ro_qualifier.ll @@ -1,5 +1,5 @@ ; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %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 %} ; CHECK-SPIRV: %[[#IMAGE_TYPE:]] = OpTypeImage ; CHECK-SPIRV: %[[#IMAGE_ARG:]] = OpFunctionParameter %[[#IMAGE_TYPE]] From 31b490628ce2ab7fc5d6653f7ec717a93ebf615d Mon Sep 17 00:00:00 2001 From: "Levytskyy, Vyacheslav" Date: Wed, 5 Jun 2024 09:11:53 -0700 Subject: [PATCH 3/5] add a test case --- llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp | 40 +++++++++++-------- llvm/lib/Target/SPIRV/SPIRVBuiltins.td | 1 + .../transcoding/OpGroupAsyncCopy-strided.ll | 36 +++++++++++++++++ 3 files changed, 61 insertions(+), 16 deletions(-) create mode 100644 llvm/test/CodeGen/SPIRV/transcoding/OpGroupAsyncCopy-strided.ll diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp index 9b9b8f7cbc089..93209e1728008 100644 --- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp @@ -25,12 +25,6 @@ namespace llvm { -// Defined in SPIRVPreLegalizer.cpp. -extern Register insertAssignInstr(Register Reg, Type *Ty, SPIRVType *SpirvTy, - SPIRVGlobalRegistry *GR, - MachineIRBuilder &MIB, - MachineRegisterInfo &MRI); - namespace SPIRV { #define GET_BuiltinGroup_DECL #include "SPIRVGenTables.inc" @@ -2073,16 +2067,30 @@ static bool generateAsyncCopy(const SPIRV::IncomingCall *Call, auto Scope = buildConstantIntReg(SPIRV::Scope::Workgroup, MIRBuilder, GR); switch (Opcode) { - case SPIRV::OpGroupAsyncCopy: - return MIRBuilder.buildInstr(Opcode) - .addDef(Call->ReturnRegister) - .addUse(GR->getSPIRVTypeID(Call->ReturnType)) - .addUse(Scope) - .addUse(Call->Arguments[0]) - .addUse(Call->Arguments[1]) - .addUse(Call->Arguments[2]) - .addUse(buildConstantIntReg(1, MIRBuilder, GR)) - .addUse(Call->Arguments[3]); + case SPIRV::OpGroupAsyncCopy: { + SPIRVType *NewType = + Call->ReturnType->getOpcode() == SPIRV::OpTypeEvent + ? nullptr + : GR->getOrCreateSPIRVTypeByName("spirv.Event", MIRBuilder); + Register TypeReg = GR->getSPIRVTypeID(NewType ? NewType : Call->ReturnType); + unsigned NumArgs = Call->Arguments.size(); + Register EventReg = Call->Arguments[NumArgs - 1]; + bool Res = MIRBuilder.buildInstr(Opcode) + .addDef(Call->ReturnRegister) + .addUse(TypeReg) + .addUse(Scope) + .addUse(Call->Arguments[0]) + .addUse(Call->Arguments[1]) + .addUse(Call->Arguments[2]) + .addUse(Call->Arguments.size() > 4 + ? Call->Arguments[3] + : buildConstantIntReg(1, MIRBuilder, GR)) + .addUse(EventReg); + if (NewType != nullptr) + insertAssignInstr(Call->ReturnRegister, nullptr, NewType, GR, MIRBuilder, + MIRBuilder.getMF().getRegInfo()); + return Res; + } case SPIRV::OpGroupWaitEvents: return MIRBuilder.buildInstr(Opcode) .addUse(Scope) diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.td b/llvm/lib/Target/SPIRV/SPIRVBuiltins.td index 692234c405ab6..da547cbab4e98 100644 --- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.td +++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.td @@ -586,6 +586,7 @@ defm : DemangledNativeBuiltin<"__spirv_SpecConstantComposite", OpenCL_std, SpecC // Async Copy and Prefetch builtin records: defm : DemangledNativeBuiltin<"async_work_group_copy", OpenCL_std, AsyncCopy, 4, 4, OpGroupAsyncCopy>; +defm : DemangledNativeBuiltin<"async_work_group_strided_copy", OpenCL_std, AsyncCopy, 5, 5, OpGroupAsyncCopy>; defm : DemangledNativeBuiltin<"__spirv_GroupAsyncCopy", OpenCL_std, AsyncCopy, 6, 6, OpGroupAsyncCopy>; defm : DemangledNativeBuiltin<"wait_group_events", OpenCL_std, AsyncCopy, 2, 2, OpGroupWaitEvents>; defm : DemangledNativeBuiltin<"__spirv_GroupWaitEvents", OpenCL_std, AsyncCopy, 3, 3, OpGroupWaitEvents>; diff --git a/llvm/test/CodeGen/SPIRV/transcoding/OpGroupAsyncCopy-strided.ll b/llvm/test/CodeGen/SPIRV/transcoding/OpGroupAsyncCopy-strided.ll new file mode 100644 index 0000000000000..96d6016083f06 --- /dev/null +++ b/llvm/test/CodeGen/SPIRV/transcoding/OpGroupAsyncCopy-strided.ll @@ -0,0 +1,36 @@ +; RUN: llc -O0 -mtriple=spirv64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-SPIRV +; RUN: %if spirv-tools %{ llc -O0 -mtriple=spirv64-unknown-unknown %s -o - -filetype=obj | spirv-val %} + +; 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 %} + +; CHECK-SPIRV-DAG: %[[#LongTy:]] = OpTypeInt 64 0 +; CHECK-SPIRV-DAG: %[[#IntTy:]] = OpTypeInt 32 0 +; CHECK-SPIRV-DAG: %[[#EventTy:]] = OpTypeEvent +; CHECK-SPIRV-DAG: %[[#Scope:]] = OpConstant %[[#IntTy]] 2 +; CHECK-SPIRV-DAG: %[[#Num:]] = OpConstant %[[#LongTy]] 123 +; CHECK-SPIRV-DAG: %[[#Null:]] = OpConstantNull +; CHECK-SPIRV-DAG: %[[#Stride:]] = OpConstant %[[#LongTy]] 1 +; CHECK-SPIRV-DAG: %[[#GenPtrEventTy:]] = OpTypePointer Generic %[[#EventTy]] +; CHECK-SPIRV-DAG: %[[#FunPtrEventTy:]] = OpTypePointer Function %[[#EventTy]] +; CHECK-SPIRV: OpFunction +; CHECK-SPIRV: %[[#Var:]] = OpVariable %[[#]] Function +; CHECK-SPIRV: %[[#ResEvent:]] = OpGroupAsyncCopy %[[#EventTy]] %[[#Scope]] %[[#Null]] %[[#Null]] %[[#Num]] %[[#Stride]] %[[#Null]] +; CHECK-SPIRV: %[[#VarPtrEvent:]] = OpBitcast %[[#FunPtrEventTy]] %[[#Var]] +; CHECK-SPIRV: OpStore %[[#VarPtrEvent]] %[[#ResEvent]] +; CHECK-SPIRV: %[[#VarPtrEvent2:]] = OpBitcast %[[#FunPtrEventTy]] %[[#Var]] +; CHECK-SPIRV: %[[#PtrEventGen:]] = OpPtrCastToGeneric %[[#]] %[[#VarPtrEvent2]] +; CHECK-SPIRV: OpGroupWaitEvents %[[#Scope]] %[[#Num]] %[[#PtrEventGen]] +; CHECK-SPIRV: OpFunctionEnd + +define spir_kernel void @foo() { + %event = alloca ptr, align 8 + %call = call spir_func ptr @_Z29async_work_group_strided_copyPU3AS3hPU3AS1Khmm9ocl_event(ptr null, ptr null, i64 123, i64 1, ptr null) + store ptr %call, ptr %event, align 8 + %event.ascast = addrspacecast ptr %event to ptr addrspace(4) + call spir_func void @_Z17wait_group_eventsiPU3AS49ocl_event(i64 123, ptr addrspace(4) %event.ascast) + ret void +} + +declare spir_func ptr @_Z29async_work_group_strided_copyPU3AS3hPU3AS1Khmm9ocl_event(ptr, ptr, i64, i64, ptr) +declare spir_func void @_Z17wait_group_eventsiPU3AS49ocl_event(i64, ptr addrspace(4)) From 462048ce1c740e7e83750b30dc10a7f629c10def Mon Sep 17 00:00:00 2001 From: "Levytskyy, Vyacheslav" Date: Wed, 5 Jun 2024 09:16:03 -0700 Subject: [PATCH 4/5] minor edit --- llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp index 93209e1728008..38ed7b48a6cb1 100644 --- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp @@ -24,7 +24,6 @@ #define DEBUG_TYPE "spirv-builtins" namespace llvm { - namespace SPIRV { #define GET_BuiltinGroup_DECL #include "SPIRVGenTables.inc" From d1b491a7e80e8c7b7a4b2b79cf2b8b74306226e8 Mon Sep 17 00:00:00 2001 From: "Levytskyy, Vyacheslav" Date: Wed, 5 Jun 2024 09:17:10 -0700 Subject: [PATCH 5/5] minor edit --- llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp index 696706258ec40..bbd25dc85f52b 100644 --- a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp @@ -50,7 +50,6 @@ void initializeSPIRVEmitIntrinsicsPass(PassRegistry &); } // namespace llvm namespace { - class SPIRVEmitIntrinsics : public ModulePass, public InstVisitor {