diff --git a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp index 424087f361a6a..38ed7b48a6cb1 100644 --- a/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVBuiltins.cpp @@ -1451,11 +1451,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) @@ -2055,16 +2066,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/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp index 5ef0be1cab722..bbd25dc85f52b 100644 --- a/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp +++ b/llvm/lib/Target/SPIRV/SPIRVEmitIntrinsics.cpp @@ -61,9 +61,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 +95,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 +217,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 +536,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 +579,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 +773,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 +878,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 +957,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 +979,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 +1172,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 +1191,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 +1407,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: } 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)) 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]]