Lines Matching +full:foo +full:- +full:queue
1 //===-- NVPTXLowerArgs.cpp - Lower arguments ------------------------------===//
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
7 //===----------------------------------------------------------------------===//
12 // http://docs.nvidia.com/cuda/parallel-thread-execution/#state-spaces
14 // Kernel parameters are read-only and accessible only via ld.param
29 // 1. Convert non-byval pointer arguments of CUDA kernels to pointers in the
34 // define void @foo(float* %input) {
41 // define void @foo(float* %input) {
50 // define void @foo(float* %input) {
65 // __global__ void foo(S s) {
72 // define void @foo(ptr byval %input) {
80 // define void @foo({i32*, i32*}* byval %input) {
92 // define void @foo(ptr byval(%struct.s) align 4 %input) {
99 // define void @foo(ptr byval(%struct.s) align 4 %input) #1 {
117 // define void @foo(ptr byval(%struct.s) %input) {
125 // define void @foo(ptr byval(%struct.s) %input) {
128 // ; to prevent generic -> param -> generic from getting cancelled out
136 //===----------------------------------------------------------------------===//
153 #include <queue>
155 #define DEBUG_TYPE "nvptx-lower-args"
174 // NVPTXInferAddressSpaces to fold the global-to-generic cast into
192 INITIALIZE_PASS_BEGIN(NVPTXLowerArgs, "nvptx-lower-args",
195 INITIALIZE_PASS_END(NVPTXLowerArgs, "nvptx-lower-args", in INITIALIZE_PASS_DEPENDENCY()
199 // If the function had a byval struct ptr arg, say foo(%struct.x* byval %d), in INITIALIZE_PASS_DEPENDENCY()
221 Instruction *I = dyn_cast<Instruction>(OldUse->getUser()); in INITIALIZE_PASS_DEPENDENCY()
231 auto CloneInstInParamAS = [GridConstant](const IP &I) -> Value * { in INITIALIZE_PASS_DEPENDENCY()
233 LI->setOperand(0, I.NewParam); in INITIALIZE_PASS_DEPENDENCY()
237 SmallVector<Value *, 4> Indices(GEP->indices()); in INITIALIZE_PASS_DEPENDENCY()
239 GEP->getSourceElementType(), I.NewParam, Indices, GEP->getName(), in INITIALIZE_PASS_DEPENDENCY()
240 GEP->getIterator()); in INITIALIZE_PASS_DEPENDENCY()
241 NewGEP->setIsInBounds(GEP->isInBounds()); in INITIALIZE_PASS_DEPENDENCY()
245 auto *NewBCType = PointerType::get(BC->getContext(), ADDRESS_SPACE_PARAM); in INITIALIZE_PASS_DEPENDENCY()
246 return BitCastInst::Create(BC->getOpcode(), I.NewParam, NewBCType, in INITIALIZE_PASS_DEPENDENCY()
247 BC->getName(), BC->getIterator()); in INITIALIZE_PASS_DEPENDENCY()
250 assert(ASC->getDestAddressSpace() == ADDRESS_SPACE_PARAM); in INITIALIZE_PASS_DEPENDENCY()
258 [](Value *Addr, Instruction *OriginalUser) -> Value * { in INITIALIZE_PASS_DEPENDENCY()
260 PointerType::get(OriginalUser->getContext(), ADDRESS_SPACE_GENERIC); in INITIALIZE_PASS_DEPENDENCY()
262 OriginalUser->getModule(), Intrinsic::nvvm_ptr_param_to_gen, in INITIALIZE_PASS_DEPENDENCY()
263 {ReturnTy, PointerType::get(OriginalUser->getContext(), in INITIALIZE_PASS_DEPENDENCY()
268 CallInst::Create(CvtToGen, Addr, Addr->getName() + ".gen", in INITIALIZE_PASS_DEPENDENCY()
269 OriginalUser->getIterator()); in INITIALIZE_PASS_DEPENDENCY()
274 I.OldUse->set(GetParamAddrCastToGeneric(I.NewParam, CI)); in INITIALIZE_PASS_DEPENDENCY()
279 if (SI->getValueOperand() == I.OldUse->get()) in INITIALIZE_PASS_DEPENDENCY()
280 SI->setOperand(0, GetParamAddrCastToGeneric(I.NewParam, SI)); in INITIALIZE_PASS_DEPENDENCY()
284 if (PI->getPointerOperand() == I.OldUse->get()) in INITIALIZE_PASS_DEPENDENCY()
285 PI->setOperand(0, GetParamAddrCastToGeneric(I.NewParam, PI)); in INITIALIZE_PASS_DEPENDENCY()
300 // We've created a new instruction. Queue users of the old instruction to in INITIALIZE_PASS_DEPENDENCY()
303 for (Use &U : I.OldInstruction->uses()) in INITIALIZE_PASS_DEPENDENCY()
318 I->eraseFromParent(); in INITIALIZE_PASS_DEPENDENCY()
329 Function *Func = Arg->getParent(); in adjustByValArgAlignment()
330 Type *StructType = Arg->getParamByValType(); in adjustByValArgAlignment()
331 const DataLayout DL(Func->getParent()); in adjustByValArgAlignment()
334 TLI->getFunctionParamOptimizedAlign(Func, StructType, DL).value(); in adjustByValArgAlignment()
336 Arg->getAttribute(Attribute::Alignment).getValueAsInt(); in adjustByValArgAlignment()
345 Attribute::get(Func->getContext(), Attribute::Alignment, NewArgAlign); in adjustByValArgAlignment()
346 Arg->removeAttr(Attribute::Alignment); in adjustByValArgAlignment()
347 Arg->addAttr(NewAlignAttr); in adjustByValArgAlignment()
360 std::queue<LoadContext> Worklist; in adjustByValArgAlignment()
368 for (User *CurUser : Ctx.InitialVal->users()) { in adjustByValArgAlignment()
383 if (!I->accumulateConstantOffset(DL, OffsetAccumulated)) in adjustByValArgAlignment()
386 uint64_t OffsetLimit = -1; in adjustByValArgAlignment()
407 Align CurLoadAlign(CurLoad.Inst->getAlign()); in adjustByValArgAlignment()
408 CurLoad.Inst->setAlignment(std::max(NewLoadAlign, CurLoadAlign)); in adjustByValArgAlignment()
415 Function *Func = Arg->getParent(); in handleByValParam()
416 BasicBlock::iterator FirstInst = Func->getEntryBlock().begin(); in handleByValParam()
417 Type *StructType = Arg->getParamByValType(); in handleByValParam()
422 auto IsSupportedUse = [IsGridConstant](Value *V) -> bool { in handleByValParam()
425 // ASC to param space are OK, too -- we'll just strip them. in handleByValParam()
427 if (ASC->getDestAddressSpace() == ADDRESS_SPACE_PARAM) in handleByValParam()
449 llvm::append_range(ValuesToCheck, V->users()); in handleByValParam()
454 if (llvm::all_of(Arg->users(), AreSupportedUsers)) { in handleByValParam()
458 for (Use &U : Arg->uses()) in handleByValParam()
462 Arg, PointerType::get(StructType, ADDRESS_SPACE_PARAM), Arg->getName(), in handleByValParam()
469 cast<NVPTXTargetLowering>(TM.getSubtargetImpl()->getTargetLowering()); in handleByValParam()
476 const DataLayout &DL = Func->getDataLayout(); in handleByValParam()
483 IRBuilder<> IRB(&Func->getEntryBlock().front()); in handleByValParam()
487 Arg, IRB.getPtrTy(ADDRESS_SPACE_PARAM), Arg->getName() + ".param")); in handleByValParam()
491 // generic address space, and a `generic -> param` cast followed by a `param in handleByValParam()
492 // -> generic` cast will be folded away. The `param -> generic` intrinsic in handleByValParam()
496 CastToParam, nullptr, CastToParam->getName() + ".gen"); in handleByValParam()
498 Arg->replaceAllUsesWith(CvtToGenCall); in handleByValParam()
501 CastToParam->setOperand(0, Arg); in handleByValParam()
505 new AllocaInst(StructType, AS, Arg->getName(), FirstInst); in handleByValParam()
509 AllocA->setAlignment(Func->getParamAlign(Arg->getArgNo()) in handleByValParam()
511 Arg->replaceAllUsesWith(AllocA); in handleByValParam()
514 Arg, PointerType::get(Arg->getContext(), ADDRESS_SPACE_PARAM), in handleByValParam()
515 Arg->getName(), FirstInst); in handleByValParam()
520 new LoadInst(StructType, ArgInParam, Arg->getName(), in handleByValParam()
521 /*isVolatile=*/false, AllocA->getAlign(), FirstInst); in handleByValParam()
527 if (Ptr->getType()->getPointerAddressSpace() != ADDRESS_SPACE_GENERIC) in markPointerAsGlobal()
534 InsertPt = Arg->getParent()->getEntryBlock().begin(); in markPointerAsGlobal()
537 InsertPt = ++cast<Instruction>(Ptr)->getIterator(); in markPointerAsGlobal()
538 assert(InsertPt != InsertPt->getParent()->end() && in markPointerAsGlobal()
543 Ptr, PointerType::get(Ptr->getContext(), ADDRESS_SPACE_GLOBAL), in markPointerAsGlobal()
544 Ptr->getName(), InsertPt); in markPointerAsGlobal()
545 Value *PtrInGeneric = new AddrSpaceCastInst(PtrInGlobal, Ptr->getType(), in markPointerAsGlobal()
546 Ptr->getName(), InsertPt); in markPointerAsGlobal()
548 Ptr->replaceAllUsesWith(PtrInGeneric); in markPointerAsGlobal()
549 PtrInGlobal->setOperand(0, Ptr); in markPointerAsGlobal()
573 if (LI->getType()->isPointerTy() || LI->getType()->isIntegerTy()) { in runOnKernelFunction()
574 Value *UO = getUnderlyingObject(LI->getPointerOperand()); in runOnKernelFunction()
576 if (Arg->hasByValAttr()) { in runOnKernelFunction()
578 if (LI->getType()->isPointerTy()) in runOnKernelFunction()
592 if (Arg.getType()->isPointerTy()) { in runOnKernelFunction()
597 } else if (Arg.getType()->isIntegerTy() && in runOnKernelFunction()
610 if (Arg.getType()->isPointerTy() && Arg.hasByValAttr()) in runOnDeviceFunction()