LCOV - code coverage report
Current view: top level - lib/Target/AMDGPU - AMDGPUPromoteAlloca.cpp (source / functions) Hit Total Coverage
Test: llvm-toolchain.info Lines: 289 297 97.3 %
Date: 2018-06-17 00:07:59 Functions: 19 21 90.5 %
Legend: Lines: hit not hit

          Line data    Source code
       1             : //===-- AMDGPUPromoteAlloca.cpp - Promote Allocas -------------------------===//
       2             : //
       3             : //                     The LLVM Compiler Infrastructure
       4             : //
       5             : // This file is distributed under the University of Illinois Open Source
       6             : // License. See LICENSE.TXT for details.
       7             : //
       8             : //===----------------------------------------------------------------------===//
       9             : //
      10             : // This pass eliminates allocas by either converting them into vectors or
      11             : // by migrating them to local address space.
      12             : //
      13             : //===----------------------------------------------------------------------===//
      14             : 
      15             : #include "AMDGPU.h"
      16             : #include "AMDGPUSubtarget.h"
      17             : #include "Utils/AMDGPUBaseInfo.h"
      18             : #include "llvm/ADT/APInt.h"
      19             : #include "llvm/ADT/None.h"
      20             : #include "llvm/ADT/STLExtras.h"
      21             : #include "llvm/ADT/StringRef.h"
      22             : #include "llvm/ADT/Triple.h"
      23             : #include "llvm/ADT/Twine.h"
      24             : #include "llvm/Analysis/CaptureTracking.h"
      25             : #include "llvm/Analysis/ValueTracking.h"
      26             : #include "llvm/CodeGen/TargetPassConfig.h"
      27             : #include "llvm/IR/Attributes.h"
      28             : #include "llvm/IR/BasicBlock.h"
      29             : #include "llvm/IR/Constant.h"
      30             : #include "llvm/IR/Constants.h"
      31             : #include "llvm/IR/DataLayout.h"
      32             : #include "llvm/IR/DerivedTypes.h"
      33             : #include "llvm/IR/Function.h"
      34             : #include "llvm/IR/GlobalValue.h"
      35             : #include "llvm/IR/GlobalVariable.h"
      36             : #include "llvm/IR/IRBuilder.h"
      37             : #include "llvm/IR/Instruction.h"
      38             : #include "llvm/IR/Instructions.h"
      39             : #include "llvm/IR/IntrinsicInst.h"
      40             : #include "llvm/IR/Intrinsics.h"
      41             : #include "llvm/IR/LLVMContext.h"
      42             : #include "llvm/IR/Metadata.h"
      43             : #include "llvm/IR/Module.h"
      44             : #include "llvm/IR/Type.h"
      45             : #include "llvm/IR/User.h"
      46             : #include "llvm/IR/Value.h"
      47             : #include "llvm/Pass.h"
      48             : #include "llvm/Support/Casting.h"
      49             : #include "llvm/Support/Debug.h"
      50             : #include "llvm/Support/ErrorHandling.h"
      51             : #include "llvm/Support/MathExtras.h"
      52             : #include "llvm/Support/raw_ostream.h"
      53             : #include "llvm/Target/TargetMachine.h"
      54             : #include <algorithm>
      55             : #include <cassert>
      56             : #include <cstdint>
      57             : #include <map>
      58             : #include <tuple>
      59             : #include <utility>
      60             : #include <vector>
      61             : 
      62             : #define DEBUG_TYPE "amdgpu-promote-alloca"
      63             : 
      64             : using namespace llvm;
      65             : 
      66             : namespace {
      67             : 
      68      101169 : static cl::opt<bool> DisablePromoteAllocaToVector(
      69             :   "disable-promote-alloca-to-vector",
      70      101169 :   cl::desc("Disable promote alloca to vector"),
      71      303507 :   cl::init(false));
      72             : 
      73             : // FIXME: This can create globals so should be a module pass.
      74        4080 : class AMDGPUPromoteAlloca : public FunctionPass {
      75             : private:
      76             :   const TargetMachine *TM;
      77             :   Module *Mod = nullptr;
      78             :   const DataLayout *DL = nullptr;
      79             :   AMDGPUAS AS;
      80             : 
      81             :   // FIXME: This should be per-kernel.
      82             :   uint32_t LocalMemLimit = 0;
      83             :   uint32_t CurrentLocalMemUsage = 0;
      84             : 
      85             :   bool IsAMDGCN = false;
      86             :   bool IsAMDHSA = false;
      87             : 
      88             :   std::pair<Value *, Value *> getLocalSizeYZ(IRBuilder<> &Builder);
      89             :   Value *getWorkitemID(IRBuilder<> &Builder, unsigned N);
      90             : 
      91             :   /// BaseAlloca is the alloca root the search started from.
      92             :   /// Val may be that alloca or a recursive user of it.
      93             :   bool collectUsesWithPtrTypes(Value *BaseAlloca,
      94             :                                Value *Val,
      95             :                                std::vector<Value*> &WorkList) const;
      96             : 
      97             :   /// Val is a derived pointer from Alloca. OpIdx0/OpIdx1 are the operand
      98             :   /// indices to an instruction with 2 pointer inputs (e.g. select, icmp).
      99             :   /// Returns true if both operands are derived from the same alloca. Val should
     100             :   /// be the same value as one of the input operands of UseInst.
     101             :   bool binaryOpIsDerivedFromSameAlloca(Value *Alloca, Value *Val,
     102             :                                        Instruction *UseInst,
     103             :                                        int OpIdx0, int OpIdx1) const;
     104             : 
     105             :   /// Check whether we have enough local memory for promotion.
     106             :   bool hasSufficientLocalMem(const Function &F);
     107             : 
     108             : public:
     109             :   static char ID;
     110             : 
     111        4096 :   AMDGPUPromoteAlloca() : FunctionPass(ID) {}
     112             : 
     113             :   bool doInitialization(Module &M) override;
     114             :   bool runOnFunction(Function &F) override;
     115             : 
     116           0 :   StringRef getPassName() const override { return "AMDGPU Promote Alloca"; }
     117             : 
     118             :   bool handleAlloca(AllocaInst &I, bool SufficientLDS);
     119             : 
     120        2037 :   void getAnalysisUsage(AnalysisUsage &AU) const override {
     121        2037 :     AU.setPreservesCFG();
     122        2037 :     FunctionPass::getAnalysisUsage(AU);
     123        2037 :   }
     124             : };
     125             : 
     126             : } // end anonymous namespace
     127             : 
     128             : char AMDGPUPromoteAlloca::ID = 0;
     129             : 
     130      357084 : INITIALIZE_PASS(AMDGPUPromoteAlloca, DEBUG_TYPE,
     131             :                 "AMDGPU promote alloca to vector or LDS", false, false)
     132             : 
     133             : char &llvm::AMDGPUPromoteAllocaID = AMDGPUPromoteAlloca::ID;
     134             : 
     135        2037 : bool AMDGPUPromoteAlloca::doInitialization(Module &M) {
     136        2037 :   Mod = &M;
     137        2037 :   DL = &Mod->getDataLayout();
     138             : 
     139        2037 :   return false;
     140             : }
     141             : 
     142       19857 : bool AMDGPUPromoteAlloca::runOnFunction(Function &F) {
     143       19857 :   if (skipFunction(F))
     144             :     return false;
     145             : 
     146       19855 :   if (auto *TPC = getAnalysisIfAvailable<TargetPassConfig>())
     147       19855 :     TM = &TPC->getTM<TargetMachine>();
     148             :   else
     149             :     return false;
     150             : 
     151             :   const Triple &TT = TM->getTargetTriple();
     152       19855 :   IsAMDGCN = TT.getArch() == Triple::amdgcn;
     153       19855 :   IsAMDHSA = TT.getOS() == Triple::AMDHSA;
     154             : 
     155             :   const AMDGPUSubtarget &ST = TM->getSubtarget<AMDGPUSubtarget>(F);
     156       19855 :   if (!ST.isPromoteAllocaEnabled())
     157             :     return false;
     158             : 
     159       19271 :   AS = AMDGPU::getAMDGPUAS(*F.getParent());
     160             : 
     161       19271 :   bool SufficientLDS = hasSufficientLocalMem(F);
     162             :   bool Changed = false;
     163             :   BasicBlock &EntryBB = *F.begin();
     164      128689 :   for (auto I = EntryBB.begin(), E = EntryBB.end(); I != E; ) {
     165             :     AllocaInst *AI = dyn_cast<AllocaInst>(I);
     166             : 
     167             :     ++I;
     168      109418 :     if (AI)
     169         525 :       Changed |= handleAlloca(*AI, SufficientLDS);
     170             :   }
     171             : 
     172             :   return Changed;
     173             : }
     174             : 
     175             : std::pair<Value *, Value *>
     176         165 : AMDGPUPromoteAlloca::getLocalSizeYZ(IRBuilder<> &Builder) {
     177         165 :   const AMDGPUSubtarget &ST = TM->getSubtarget<AMDGPUSubtarget>(
     178         165 :                                 *Builder.GetInsertBlock()->getParent());
     179             : 
     180         165 :   if (!IsAMDHSA) {
     181             :     Function *LocalSizeYFn
     182          99 :       = Intrinsic::getDeclaration(Mod, Intrinsic::r600_read_local_size_y);
     183             :     Function *LocalSizeZFn
     184          99 :       = Intrinsic::getDeclaration(Mod, Intrinsic::r600_read_local_size_z);
     185             : 
     186          99 :     CallInst *LocalSizeY = Builder.CreateCall(LocalSizeYFn, {});
     187          99 :     CallInst *LocalSizeZ = Builder.CreateCall(LocalSizeZFn, {});
     188             : 
     189          99 :     ST.makeLIDRangeMetadata(LocalSizeY);
     190          99 :     ST.makeLIDRangeMetadata(LocalSizeZ);
     191             : 
     192          99 :     return std::make_pair(LocalSizeY, LocalSizeZ);
     193             :   }
     194             : 
     195             :   // We must read the size out of the dispatch pointer.
     196             :   assert(IsAMDGCN);
     197             : 
     198             :   // We are indexing into this struct, and want to extract the workgroup_size_*
     199             :   // fields.
     200             :   //
     201             :   //   typedef struct hsa_kernel_dispatch_packet_s {
     202             :   //     uint16_t header;
     203             :   //     uint16_t setup;
     204             :   //     uint16_t workgroup_size_x ;
     205             :   //     uint16_t workgroup_size_y;
     206             :   //     uint16_t workgroup_size_z;
     207             :   //     uint16_t reserved0;
     208             :   //     uint32_t grid_size_x ;
     209             :   //     uint32_t grid_size_y ;
     210             :   //     uint32_t grid_size_z;
     211             :   //
     212             :   //     uint32_t private_segment_size;
     213             :   //     uint32_t group_segment_size;
     214             :   //     uint64_t kernel_object;
     215             :   //
     216             :   // #ifdef HSA_LARGE_MODEL
     217             :   //     void *kernarg_address;
     218             :   // #elif defined HSA_LITTLE_ENDIAN
     219             :   //     void *kernarg_address;
     220             :   //     uint32_t reserved1;
     221             :   // #else
     222             :   //     uint32_t reserved1;
     223             :   //     void *kernarg_address;
     224             :   // #endif
     225             :   //     uint64_t reserved2;
     226             :   //     hsa_signal_t completion_signal; // uint64_t wrapper
     227             :   //   } hsa_kernel_dispatch_packet_t
     228             :   //
     229             :   Function *DispatchPtrFn
     230          66 :     = Intrinsic::getDeclaration(Mod, Intrinsic::amdgcn_dispatch_ptr);
     231             : 
     232          66 :   CallInst *DispatchPtr = Builder.CreateCall(DispatchPtrFn, {});
     233          66 :   DispatchPtr->addAttribute(AttributeList::ReturnIndex, Attribute::NoAlias);
     234          66 :   DispatchPtr->addAttribute(AttributeList::ReturnIndex, Attribute::NonNull);
     235             : 
     236             :   // Size of the dispatch packet struct.
     237          66 :   DispatchPtr->addDereferenceableAttr(AttributeList::ReturnIndex, 64);
     238             : 
     239          66 :   Type *I32Ty = Type::getInt32Ty(Mod->getContext());
     240          66 :   Value *CastDispatchPtr = Builder.CreateBitCast(
     241          66 :     DispatchPtr, PointerType::get(I32Ty, AS.CONSTANT_ADDRESS));
     242             : 
     243             :   // We could do a single 64-bit load here, but it's likely that the basic
     244             :   // 32-bit and extract sequence is already present, and it is probably easier
     245             :   // to CSE this. The loads should be mergable later anyway.
     246          66 :   Value *GEPXY = Builder.CreateConstInBoundsGEP1_64(CastDispatchPtr, 1);
     247          66 :   LoadInst *LoadXY = Builder.CreateAlignedLoad(GEPXY, 4);
     248             : 
     249          66 :   Value *GEPZU = Builder.CreateConstInBoundsGEP1_64(CastDispatchPtr, 2);
     250          66 :   LoadInst *LoadZU = Builder.CreateAlignedLoad(GEPZU, 4);
     251             : 
     252          66 :   MDNode *MD = MDNode::get(Mod->getContext(), None);
     253          66 :   LoadXY->setMetadata(LLVMContext::MD_invariant_load, MD);
     254          66 :   LoadZU->setMetadata(LLVMContext::MD_invariant_load, MD);
     255          66 :   ST.makeLIDRangeMetadata(LoadZU);
     256             : 
     257             :   // Extract y component. Upper half of LoadZU should be zero already.
     258          66 :   Value *Y = Builder.CreateLShr(LoadXY, 16);
     259             : 
     260          66 :   return std::make_pair(Y, LoadZU);
     261             : }
     262             : 
     263         495 : Value *AMDGPUPromoteAlloca::getWorkitemID(IRBuilder<> &Builder, unsigned N) {
     264         495 :   const AMDGPUSubtarget &ST = TM->getSubtarget<AMDGPUSubtarget>(
     265         495 :                                 *Builder.GetInsertBlock()->getParent());
     266             :   Intrinsic::ID IntrID = Intrinsic::ID::not_intrinsic;
     267             : 
     268         495 :   switch (N) {
     269         165 :   case 0:
     270         165 :     IntrID = IsAMDGCN ? Intrinsic::amdgcn_workitem_id_x
     271             :       : Intrinsic::r600_read_tidig_x;
     272             :     break;
     273         165 :   case 1:
     274         165 :     IntrID = IsAMDGCN ? Intrinsic::amdgcn_workitem_id_y
     275             :       : Intrinsic::r600_read_tidig_y;
     276             :     break;
     277             : 
     278         165 :   case 2:
     279         165 :     IntrID = IsAMDGCN ? Intrinsic::amdgcn_workitem_id_z
     280             :       : Intrinsic::r600_read_tidig_z;
     281             :     break;
     282           0 :   default:
     283           0 :     llvm_unreachable("invalid dimension");
     284             :   }
     285             : 
     286         495 :   Function *WorkitemIdFn = Intrinsic::getDeclaration(Mod, IntrID);
     287         495 :   CallInst *CI = Builder.CreateCall(WorkitemIdFn);
     288         495 :   ST.makeLIDRangeMetadata(CI);
     289             : 
     290         495 :   return CI;
     291             : }
     292             : 
     293             : static VectorType *arrayTypeToVecType(ArrayType *ArrayTy) {
     294          60 :   return VectorType::get(ArrayTy->getElementType(),
     295         120 :                          ArrayTy->getNumElements());
     296             : }
     297             : 
     298             : static Value *
     299             : calculateVectorIndex(Value *Ptr,
     300             :                      const std::map<GetElementPtrInst *, Value *> &GEPIdx) {
     301             :   GetElementPtrInst *GEP = cast<GetElementPtrInst>(Ptr);
     302             : 
     303             :   auto I = GEPIdx.find(GEP);
     304         211 :   return I == GEPIdx.end() ? nullptr : I->second;
     305             : }
     306             : 
     307         271 : static Value* GEPToVectorIndex(GetElementPtrInst *GEP) {
     308             :   // FIXME we only support simple cases
     309         271 :   if (GEP->getNumOperands() != 3)
     310             :     return nullptr;
     311             : 
     312             :   ConstantInt *I0 = dyn_cast<ConstantInt>(GEP->getOperand(1));
     313         271 :   if (!I0 || !I0->isZero())
     314             :     return nullptr;
     315             : 
     316             :   return GEP->getOperand(2);
     317             : }
     318             : 
     319             : // Not an instruction handled below to turn into a vector.
     320             : //
     321             : // TODO: Check isTriviallyVectorizable for calls and handle other
     322             : // instructions.
     323         288 : static bool canVectorizeInst(Instruction *Inst, User *User) {
     324         288 :   switch (Inst->getOpcode()) {
     325             :   case Instruction::Load: {
     326             :     // Currently only handle the case where the Pointer Operand is a GEP.
     327             :     // Also we could not vectorize volatile or atomic loads.
     328             :     LoadInst *LI = cast<LoadInst>(Inst);
     329             :     return isa<GetElementPtrInst>(LI->getPointerOperand()) && LI->isSimple();
     330             :   }
     331             :   case Instruction::BitCast:
     332             :     return true;
     333             :   case Instruction::Store: {
     334             :     // Must be the stored pointer operand, not a stored value, plus
     335             :     // since it should be canonical form, the User should be a GEP.
     336             :     // Also we could not vectorize volatile or atomic stores.
     337             :     StoreInst *SI = cast<StoreInst>(Inst);
     338         179 :     return (SI->getPointerOperand() == User) && isa<GetElementPtrInst>(User) && SI->isSimple();
     339             :   }
     340          22 :   default:
     341          22 :     return false;
     342             :   }
     343             : }
     344             : 
     345         518 : static bool tryPromoteAllocaToVector(AllocaInst *Alloca, AMDGPUAS AS) {
     346             : 
     347         518 :   if (DisablePromoteAllocaToVector) {
     348             :     LLVM_DEBUG(dbgs() << "  Promotion alloca to vector is disabled\n");
     349             :     return false;
     350             :   }
     351             : 
     352         298 :   ArrayType *AllocaTy = dyn_cast<ArrayType>(Alloca->getAllocatedType());
     353             : 
     354             :   LLVM_DEBUG(dbgs() << "Alloca candidate for vectorization\n");
     355             : 
     356             :   // FIXME: There is no reason why we can't support larger arrays, we
     357             :   // are just being conservative for now.
     358             :   // FIXME: We also reject alloca's of the form [ 2 x [ 2 x i32 ]] or equivalent. Potentially these
     359             :   // could also be promoted but we don't currently handle this case
     360         199 :   if (!AllocaTy ||
     361         348 :       AllocaTy->getNumElements() > 16 ||
     362         143 :       AllocaTy->getNumElements() < 2 ||
     363         143 :       !VectorType::isValidElementType(AllocaTy->getElementType())) {
     364             :     LLVM_DEBUG(dbgs() << "  Cannot convert type to vector\n");
     365             :     return false;
     366             :   }
     367             : 
     368             :   std::map<GetElementPtrInst*, Value*> GEPVectorIdx;
     369             :   std::vector<Value*> WorkList;
     370         346 :   for (User *AllocaUser : Alloca->users()) {
     371         286 :     GetElementPtrInst *GEP = dyn_cast<GetElementPtrInst>(AllocaUser);
     372         292 :     if (!GEP) {
     373          15 :       if (!canVectorizeInst(cast<Instruction>(AllocaUser), Alloca))
     374          56 :         return false;
     375             : 
     376          12 :       WorkList.push_back(AllocaUser);
     377           6 :       continue;
     378             :     }
     379             : 
     380         271 :     Value *Index = GEPToVectorIndex(GEP);
     381             : 
     382             :     // If we can't compute a vector index from this GEP, then we can't
     383             :     // promote this alloca to vector.
     384         271 :     if (!Index) {
     385             :       LLVM_DEBUG(dbgs() << "  Cannot compute vector index for GEP " << *GEP
     386             :                         << '\n');
     387             :       return false;
     388             :     }
     389             : 
     390         270 :     GEPVectorIdx[GEP] = Index;
     391         497 :     for (User *GEPUser : AllocaUser->users()) {
     392         273 :       if (!canVectorizeInst(cast<Instruction>(GEPUser), AllocaUser))
     393             :         return false;
     394             : 
     395         454 :       WorkList.push_back(GEPUser);
     396             :     }
     397             :   }
     398             : 
     399             :   VectorType *VectorTy = arrayTypeToVecType(AllocaTy);
     400             : 
     401             :   LLVM_DEBUG(dbgs() << "  Converting alloca to vector " << *AllocaTy << " -> "
     402             :                     << *VectorTy << '\n');
     403             : 
     404         286 :   for (Value *V : WorkList) {
     405             :     Instruction *Inst = cast<Instruction>(V);
     406         226 :     IRBuilder<> Builder(Inst);
     407         226 :     switch (Inst->getOpcode()) {
     408          63 :     case Instruction::Load: {
     409          63 :       Type *VecPtrTy = VectorTy->getPointerTo(AS.PRIVATE_ADDRESS);
     410             :       Value *Ptr = cast<LoadInst>(Inst)->getPointerOperand();
     411             :       Value *Index = calculateVectorIndex(Ptr, GEPVectorIdx);
     412             : 
     413          63 :       Value *BitCast = Builder.CreateBitCast(Alloca, VecPtrTy);
     414          63 :       Value *VecValue = Builder.CreateLoad(BitCast);
     415          63 :       Value *ExtractElement = Builder.CreateExtractElement(VecValue, Index);
     416          63 :       Inst->replaceAllUsesWith(ExtractElement);
     417          63 :       Inst->eraseFromParent();
     418             :       break;
     419             :     }
     420         148 :     case Instruction::Store: {
     421         148 :       Type *VecPtrTy = VectorTy->getPointerTo(AS.PRIVATE_ADDRESS);
     422             : 
     423             :       StoreInst *SI = cast<StoreInst>(Inst);
     424             :       Value *Ptr = SI->getPointerOperand();
     425             :       Value *Index = calculateVectorIndex(Ptr, GEPVectorIdx);
     426         148 :       Value *BitCast = Builder.CreateBitCast(Alloca, VecPtrTy);
     427         148 :       Value *VecValue = Builder.CreateLoad(BitCast);
     428         148 :       Value *NewVecValue = Builder.CreateInsertElement(VecValue,
     429             :                                                        SI->getValueOperand(),
     430         148 :                                                        Index);
     431         148 :       Builder.CreateStore(NewVecValue, BitCast);
     432         148 :       Inst->eraseFromParent();
     433             :       break;
     434             :     }
     435             :     case Instruction::BitCast:
     436             :     case Instruction::AddrSpaceCast:
     437             :       break;
     438             : 
     439           0 :     default:
     440           0 :       llvm_unreachable("Inconsistency in instructions promotable to vector");
     441             :     }
     442             :   }
     443             :   return true;
     444             : }
     445             : 
     446          31 : static bool isCallPromotable(CallInst *CI) {
     447             :   IntrinsicInst *II = dyn_cast<IntrinsicInst>(CI);
     448             :   if (!II)
     449             :     return false;
     450             : 
     451             :   switch (II->getIntrinsicID()) {
     452             :   case Intrinsic::memcpy:
     453             :   case Intrinsic::memmove:
     454             :   case Intrinsic::memset:
     455             :   case Intrinsic::lifetime_start:
     456             :   case Intrinsic::lifetime_end:
     457             :   case Intrinsic::invariant_start:
     458             :   case Intrinsic::invariant_end:
     459             :   case Intrinsic::launder_invariant_group:
     460             :   case Intrinsic::objectsize:
     461             :     return true;
     462           1 :   default:
     463           1 :     return false;
     464             :   }
     465             : }
     466             : 
     467          23 : bool AMDGPUPromoteAlloca::binaryOpIsDerivedFromSameAlloca(Value *BaseAlloca,
     468             :                                                           Value *Val,
     469             :                                                           Instruction *Inst,
     470             :                                                           int OpIdx0,
     471             :                                                           int OpIdx1) const {
     472             :   // Figure out which operand is the one we might not be promoting.
     473          23 :   Value *OtherOp = Inst->getOperand(OpIdx0);
     474          23 :   if (Val == OtherOp)
     475           5 :     OtherOp = Inst->getOperand(OpIdx1);
     476             : 
     477          23 :   if (isa<ConstantPointerNull>(OtherOp))
     478             :     return true;
     479             : 
     480          17 :   Value *OtherObj = GetUnderlyingObject(OtherOp, *DL);
     481             :   if (!isa<AllocaInst>(OtherObj))
     482             :     return false;
     483             : 
     484             :   // TODO: We should be able to replace undefs with the right pointer type.
     485             : 
     486             :   // TODO: If we know the other base object is another promotable
     487             :   // alloca, not necessarily this alloca, we can do this. The
     488             :   // important part is both must have the same address space at
     489             :   // the end.
     490          10 :   if (OtherObj != BaseAlloca) {
     491             :     LLVM_DEBUG(
     492             :         dbgs() << "Found a binary instruction with another alloca object\n");
     493             :     return false;
     494             :   }
     495             : 
     496             :   return true;
     497             : }
     498             : 
     499         679 : bool AMDGPUPromoteAlloca::collectUsesWithPtrTypes(
     500             :   Value *BaseAlloca,
     501             :   Value *Val,
     502             :   std::vector<Value*> &WorkList) const {
     503             : 
     504        2581 :   for (User *User : Val->users()) {
     505        1024 :     if (is_contained(WorkList, User))
     506           9 :       continue;
     507             : 
     508        1026 :     if (CallInst *CI = dyn_cast<CallInst>(User)) {
     509          31 :       if (!isCallPromotable(CI))
     510         146 :         return false;
     511             : 
     512          22 :       WorkList.push_back(User);
     513          11 :       continue;
     514             :     }
     515             : 
     516             :     Instruction *UseInst = cast<Instruction>(User);
     517         984 :     if (UseInst->getOpcode() == Instruction::PtrToInt)
     518             :       return false;
     519             : 
     520         230 :     if (LoadInst *LI = dyn_cast<LoadInst>(UseInst)) {
     521         231 :       if (LI->isVolatile())
     522             :         return false;
     523             : 
     524         230 :       continue;
     525             :     }
     526             : 
     527             :     if (StoreInst *SI = dyn_cast<StoreInst>(UseInst)) {
     528         260 :       if (SI->isVolatile())
     529             :         return false;
     530             : 
     531             :       // Reject if the stored value is not the pointer operand.
     532         253 :       if (SI->getPointerOperand() != Val)
     533             :         return false;
     534             :     } else if (AtomicRMWInst *RMW = dyn_cast<AtomicRMWInst>(UseInst)) {
     535           2 :       if (RMW->isVolatile())
     536             :         return false;
     537             :     } else if (AtomicCmpXchgInst *CAS = dyn_cast<AtomicCmpXchgInst>(UseInst)) {
     538           2 :       if (CAS->isVolatile())
     539             :         return false;
     540             :     }
     541             : 
     542             :     // Only promote a select if we know that the other select operand
     543             :     // is from another pointer that will also be promoted.
     544             :     if (ICmpInst *ICmp = dyn_cast<ICmpInst>(UseInst)) {
     545           6 :       if (!binaryOpIsDerivedFromSameAlloca(BaseAlloca, Val, ICmp, 0, 1))
     546             :         return false;
     547             : 
     548             :       // May need to rewrite constant operands.
     549           6 :       WorkList.push_back(ICmp);
     550             :     }
     551             : 
     552         730 :     if (UseInst->getOpcode() == Instruction::AddrSpaceCast) {
     553             :       // Give up if the pointer may be captured.
     554           4 :       if (PointerMayBeCaptured(UseInst, true, true))
     555             :         return false;
     556             :       // Don't collect the users of this.
     557           2 :       WorkList.push_back(User);
     558           1 :       continue;
     559             :     }
     560             : 
     561        1450 :     if (!User->getType()->isPointerTy())
     562         259 :       continue;
     563             : 
     564             :     if (GetElementPtrInst *GEP = dyn_cast<GetElementPtrInst>(UseInst)) {
     565             :       // Be conservative if an address could be computed outside the bounds of
     566             :       // the alloca.
     567         440 :       if (!GEP->isInBounds())
     568             :         return false;
     569             :     }
     570             : 
     571             :     // Only promote a select if we know that the other select operand is from
     572             :     // another pointer that will also be promoted.
     573             :     if (SelectInst *SI = dyn_cast<SelectInst>(UseInst)) {
     574          11 :       if (!binaryOpIsDerivedFromSameAlloca(BaseAlloca, Val, SI, 1, 2))
     575             :         return false;
     576             :     }
     577             : 
     578             :     // Repeat for phis.
     579             :     if (PHINode *Phi = dyn_cast<PHINode>(UseInst)) {
     580             :       // TODO: Handle more complex cases. We should be able to replace loops
     581             :       // over arrays.
     582           7 :       switch (Phi->getNumIncomingValues()) {
     583             :       case 1:
     584             :         break;
     585           6 :       case 2:
     586           6 :         if (!binaryOpIsDerivedFromSameAlloca(BaseAlloca, Val, Phi, 0, 1))
     587             :           return false;
     588             :         break;
     589             :       default:
     590             :         return false;
     591             :       }
     592             :     }
     593             : 
     594         748 :     WorkList.push_back(User);
     595         374 :     if (!collectUsesWithPtrTypes(BaseAlloca, User, WorkList))
     596             :       return false;
     597             :   }
     598             : 
     599         533 :   return true;
     600             : }
     601             : 
     602       19271 : bool AMDGPUPromoteAlloca::hasSufficientLocalMem(const Function &F) {
     603             : 
     604             :   FunctionType *FTy = F.getFunctionType();
     605       19271 :   const AMDGPUSubtarget &ST = TM->getSubtarget<AMDGPUSubtarget>(F);
     606             : 
     607             :   // If the function has any arguments in the local address space, then it's
     608             :   // possible these arguments require the entire local memory space, so
     609             :   // we cannot use local memory in the pass.
     610       98857 :   for (Type *ParamTy : FTy->params()) {
     611             :     PointerType *PtrTy = dyn_cast<PointerType>(ParamTy);
     612       26098 :     if (PtrTy && PtrTy->getAddressSpace() == AS.LOCAL_ADDRESS) {
     613        1980 :       LocalMemLimit = 0;
     614             :       LLVM_DEBUG(dbgs() << "Function has local memory argument. Promoting to "
     615             :                            "local memory disabled.\n");
     616        1980 :       return false;
     617             :     }
     618             :   }
     619             : 
     620       17291 :   LocalMemLimit = ST.getLocalMemorySize();
     621       17291 :   if (LocalMemLimit == 0)
     622             :     return false;
     623             : 
     624       17265 :   const DataLayout &DL = Mod->getDataLayout();
     625             : 
     626             :   // Check how much local memory is being used by global objects
     627       17265 :   CurrentLocalMemUsage = 0;
     628       37626 :   for (GlobalVariable &GV : Mod->globals()) {
     629        3096 :     if (GV.getType()->getAddressSpace() != AS.LOCAL_ADDRESS)
     630         557 :       continue;
     631             : 
     632        8569 :     for (const User *U : GV.users()) {
     633             :       const Instruction *Use = dyn_cast<Instruction>(U);
     634         665 :       if (!Use)
     635         665 :         continue;
     636             : 
     637        5625 :       if (Use->getParent()->getParent() == &F) {
     638             :         unsigned Align = GV.getAlignment();
     639         260 :         if (Align == 0)
     640          20 :           Align = DL.getABITypeAlignment(GV.getValueType());
     641             : 
     642             :         // FIXME: Try to account for padding here. The padding is currently
     643             :         // determined from the inverse order of uses in the function. I'm not
     644             :         // sure if the use list order is in any way connected to this, so the
     645             :         // total reported size is likely incorrect.
     646         260 :         uint64_t AllocSize = DL.getTypeAllocSize(GV.getValueType());
     647         520 :         CurrentLocalMemUsage = alignTo(CurrentLocalMemUsage, Align);
     648         260 :         CurrentLocalMemUsage += AllocSize;
     649         260 :         break;
     650             :       }
     651             :     }
     652             :   }
     653             : 
     654       17265 :   unsigned MaxOccupancy = ST.getOccupancyWithLocalMemSize(CurrentLocalMemUsage,
     655       17265 :                                                           F);
     656             : 
     657             :   // Restrict local memory usage so that we don't drastically reduce occupancy,
     658             :   // unless it is already significantly reduced.
     659             : 
     660             :   // TODO: Have some sort of hint or other heuristics to guess occupancy based
     661             :   // on other factors..
     662       17265 :   unsigned OccupancyHint = ST.getWavesPerEU(F).second;
     663       17265 :   if (OccupancyHint == 0)
     664           0 :     OccupancyHint = 7;
     665             : 
     666             :   // Clamp to max value.
     667       34530 :   OccupancyHint = std::min(OccupancyHint, ST.getMaxWavesPerEU());
     668             : 
     669             :   // Check the hint but ignore it if it's obviously wrong from the existing LDS
     670             :   // usage.
     671       17265 :   MaxOccupancy = std::min(OccupancyHint, MaxOccupancy);
     672             : 
     673             : 
     674             :   // Round up to the next tier of usage.
     675             :   unsigned MaxSizeWithWaveCount
     676       17265 :     = ST.getMaxLocalMemSizeWithWaveCount(MaxOccupancy, F);
     677             : 
     678             :   // Program is possibly broken by using more local mem than available.
     679       17265 :   if (CurrentLocalMemUsage > MaxSizeWithWaveCount)
     680             :     return false;
     681             : 
     682       17261 :   LocalMemLimit = MaxSizeWithWaveCount;
     683             : 
     684             :   LLVM_DEBUG(dbgs() << F.getName() << " uses " << CurrentLocalMemUsage
     685             :                     << " bytes of LDS\n"
     686             :                     << "  Rounding size to " << MaxSizeWithWaveCount
     687             :                     << " with a maximum occupancy of " << MaxOccupancy << '\n'
     688             :                     << " and " << (LocalMemLimit - CurrentLocalMemUsage)
     689             :                     << " available for promotion\n");
     690             : 
     691       17261 :   return true;
     692             : }
     693             : 
     694             : // FIXME: Should try to pick the most likely to be profitable allocas first.
     695         525 : bool AMDGPUPromoteAlloca::handleAlloca(AllocaInst &I, bool SufficientLDS) {
     696             :   // Array allocations are probably not worth handling, since an allocation of
     697             :   // the array type is the canonical form.
     698         525 :   if (!I.isStaticAlloca() || I.isArrayAllocation())
     699             :     return false;
     700             : 
     701         518 :   IRBuilder<> Builder(&I);
     702             : 
     703             :   // First try to replace the alloca with a vector
     704         518 :   Type *AllocaTy = I.getAllocatedType();
     705             : 
     706             :   LLVM_DEBUG(dbgs() << "Trying to promote " << I << '\n');
     707             : 
     708         518 :   if (tryPromoteAllocaToVector(&I, AS))
     709             :     return true; // Promoted to vector.
     710             : 
     711         458 :   const Function &ContainingFunction = *I.getParent()->getParent();
     712             :   CallingConv::ID CC = ContainingFunction.getCallingConv();
     713             : 
     714             :   // Don't promote the alloca to LDS for shader calling conventions as the work
     715             :   // item ID intrinsics are not supported for these calling conventions.
     716             :   // Furthermore not all LDS is available for some of the stages.
     717         458 :   switch (CC) {
     718             :   case CallingConv::AMDGPU_KERNEL:
     719             :   case CallingConv::SPIR_KERNEL:
     720             :     break;
     721             :   default:
     722             :     LLVM_DEBUG(
     723             :         dbgs()
     724             :         << " promote alloca to LDS not supported with calling convention.\n");
     725             :     return false;
     726             :   }
     727             : 
     728             :   // Not likely to have sufficient local memory for promotion.
     729         377 :   if (!SufficientLDS)
     730             :     return false;
     731             : 
     732             :   const AMDGPUSubtarget &ST =
     733         376 :     TM->getSubtarget<AMDGPUSubtarget>(ContainingFunction);
     734         376 :   unsigned WorkGroupSize = ST.getFlatWorkGroupSizes(ContainingFunction).second;
     735             : 
     736         376 :   const DataLayout &DL = Mod->getDataLayout();
     737             : 
     738             :   unsigned Align = I.getAlignment();
     739         376 :   if (Align == 0)
     740         207 :     Align = DL.getABITypeAlignment(I.getAllocatedType());
     741             : 
     742             :   // FIXME: This computed padding is likely wrong since it depends on inverse
     743             :   // usage order.
     744             :   //
     745             :   // FIXME: It is also possible that if we're allowed to use all of the memory
     746             :   // could could end up using more than the maximum due to alignment padding.
     747             : 
     748         752 :   uint32_t NewSize = alignTo(CurrentLocalMemUsage, Align);
     749         376 :   uint32_t AllocSize = WorkGroupSize * DL.getTypeAllocSize(AllocaTy);
     750         376 :   NewSize += AllocSize;
     751             : 
     752         376 :   if (NewSize > LocalMemLimit) {
     753             :     LLVM_DEBUG(dbgs() << "  " << AllocSize
     754             :                       << " bytes of local memory not available to promote\n");
     755             :     return false;
     756             :   }
     757             : 
     758         305 :   CurrentLocalMemUsage = NewSize;
     759             : 
     760             :   std::vector<Value*> WorkList;
     761             : 
     762         305 :   if (!collectUsesWithPtrTypes(&I, &I, WorkList)) {
     763             :     LLVM_DEBUG(dbgs() << " Do not know how to convert all uses\n");
     764             :     return false;
     765             :   }
     766             : 
     767             :   LLVM_DEBUG(dbgs() << "Promoting alloca to local memory\n");
     768             : 
     769         165 :   Function *F = I.getParent()->getParent();
     770             : 
     771         165 :   Type *GVTy = ArrayType::get(I.getAllocatedType(), WorkGroupSize);
     772             :   GlobalVariable *GV = new GlobalVariable(
     773         165 :       *Mod, GVTy, false, GlobalValue::InternalLinkage,
     774         165 :       UndefValue::get(GVTy),
     775         495 :       Twine(F->getName()) + Twine('.') + I.getName(),
     776             :       nullptr,
     777             :       GlobalVariable::NotThreadLocal,
     778         165 :       AS.LOCAL_ADDRESS);
     779             :   GV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
     780         330 :   GV->setAlignment(I.getAlignment());
     781             : 
     782             :   Value *TCntY, *TCntZ;
     783             : 
     784         330 :   std::tie(TCntY, TCntZ) = getLocalSizeYZ(Builder);
     785         165 :   Value *TIdX = getWorkitemID(Builder, 0);
     786         165 :   Value *TIdY = getWorkitemID(Builder, 1);
     787         165 :   Value *TIdZ = getWorkitemID(Builder, 2);
     788             : 
     789         165 :   Value *Tmp0 = Builder.CreateMul(TCntY, TCntZ, "", true, true);
     790         165 :   Tmp0 = Builder.CreateMul(Tmp0, TIdX);
     791         165 :   Value *Tmp1 = Builder.CreateMul(TIdY, TCntZ, "", true, true);
     792         165 :   Value *TID = Builder.CreateAdd(Tmp0, Tmp1);
     793         165 :   TID = Builder.CreateAdd(TID, TIdZ);
     794             : 
     795             :   Value *Indices[] = {
     796         165 :     Constant::getNullValue(Type::getInt32Ty(Mod->getContext())),
     797             :     TID
     798         165 :   };
     799             : 
     800         165 :   Value *Offset = Builder.CreateInBoundsGEP(GVTy, GV, Indices);
     801         165 :   I.mutateType(Offset->getType());
     802         165 :   I.replaceAllUsesWith(Offset);
     803         165 :   I.eraseFromParent();
     804             : 
     805         522 :   for (Value *V : WorkList) {
     806             :     CallInst *Call = dyn_cast<CallInst>(V);
     807             :     if (!Call) {
     808             :       if (ICmpInst *CI = dyn_cast<ICmpInst>(V)) {
     809             :         Value *Src0 = CI->getOperand(0);
     810           3 :         Type *EltTy = Src0->getType()->getPointerElementType();
     811           3 :         PointerType *NewTy = PointerType::get(EltTy, AS.LOCAL_ADDRESS);
     812             : 
     813           3 :         if (isa<ConstantPointerNull>(CI->getOperand(0)))
     814           1 :           CI->setOperand(0, ConstantPointerNull::get(NewTy));
     815             : 
     816           3 :         if (isa<ConstantPointerNull>(CI->getOperand(1)))
     817           1 :           CI->setOperand(1, ConstantPointerNull::get(NewTy));
     818             : 
     819           3 :         continue;
     820             :       }
     821             : 
     822             :       // The operand's value should be corrected on its own and we don't want to
     823             :       // touch the users.
     824           1 :       if (isa<AddrSpaceCastInst>(V))
     825           1 :         continue;
     826             : 
     827         342 :       Type *EltTy = V->getType()->getPointerElementType();
     828         342 :       PointerType *NewTy = PointerType::get(EltTy, AS.LOCAL_ADDRESS);
     829             : 
     830             :       // FIXME: It doesn't really make sense to try to do this for all
     831             :       // instructions.
     832             :       V->mutateType(NewTy);
     833             : 
     834             :       // Adjust the types of any constant operands.
     835             :       if (SelectInst *SI = dyn_cast<SelectInst>(V)) {
     836           9 :         if (isa<ConstantPointerNull>(SI->getOperand(1)))
     837           1 :           SI->setOperand(1, ConstantPointerNull::get(NewTy));
     838             : 
     839           9 :         if (isa<ConstantPointerNull>(SI->getOperand(2)))
     840           1 :           SI->setOperand(2, ConstantPointerNull::get(NewTy));
     841             :       } else if (PHINode *Phi = dyn_cast<PHINode>(V)) {
     842          18 :         for (unsigned I = 0, E = Phi->getNumIncomingValues(); I != E; ++I) {
     843           7 :           if (isa<ConstantPointerNull>(Phi->getIncomingValue(I)))
     844           2 :             Phi->setIncomingValue(I, ConstantPointerNull::get(NewTy));
     845             :         }
     846             :       }
     847             : 
     848         342 :       continue;
     849             :     }
     850             : 
     851             :     IntrinsicInst *Intr = cast<IntrinsicInst>(Call);
     852          11 :     Builder.SetInsertPoint(Intr);
     853          11 :     switch (Intr->getIntrinsicID()) {
     854           2 :     case Intrinsic::lifetime_start:
     855             :     case Intrinsic::lifetime_end:
     856             :       // These intrinsics are for address space 0 only
     857           2 :       Intr->eraseFromParent();
     858           2 :       continue;
     859             :     case Intrinsic::memcpy: {
     860             :       MemCpyInst *MemCpy = cast<MemCpyInst>(Intr);
     861           4 :       Builder.CreateMemCpy(MemCpy->getRawDest(), MemCpy->getDestAlignment(),
     862             :                            MemCpy->getRawSource(), MemCpy->getSourceAlignment(),
     863           2 :                            MemCpy->getLength(), MemCpy->isVolatile());
     864           2 :       Intr->eraseFromParent();
     865           2 :       continue;
     866             :     }
     867             :     case Intrinsic::memmove: {
     868             :       MemMoveInst *MemMove = cast<MemMoveInst>(Intr);
     869           4 :       Builder.CreateMemMove(MemMove->getRawDest(), MemMove->getDestAlignment(),
     870             :                             MemMove->getRawSource(), MemMove->getSourceAlignment(),
     871           2 :                             MemMove->getLength(), MemMove->isVolatile());
     872           2 :       Intr->eraseFromParent();
     873           2 :       continue;
     874             :     }
     875             :     case Intrinsic::memset: {
     876             :       MemSetInst *MemSet = cast<MemSetInst>(Intr);
     877           2 :       Builder.CreateMemSet(MemSet->getRawDest(), MemSet->getValue(),
     878             :                            MemSet->getLength(), MemSet->getDestAlignment(),
     879           1 :                            MemSet->isVolatile());
     880           1 :       Intr->eraseFromParent();
     881           1 :       continue;
     882             :     }
     883           3 :     case Intrinsic::invariant_start:
     884             :     case Intrinsic::invariant_end:
     885             :     case Intrinsic::launder_invariant_group:
     886           3 :       Intr->eraseFromParent();
     887             :       // FIXME: I think the invariant marker should still theoretically apply,
     888             :       // but the intrinsics need to be changed to accept pointers with any
     889             :       // address space.
     890           3 :       continue;
     891           1 :     case Intrinsic::objectsize: {
     892           1 :       Value *Src = Intr->getOperand(0);
     893           1 :       Type *SrcTy = Src->getType()->getPointerElementType();
     894           4 :       Function *ObjectSize = Intrinsic::getDeclaration(Mod,
     895             :         Intrinsic::objectsize,
     896           2 :         { Intr->getType(), PointerType::get(SrcTy, AS.LOCAL_ADDRESS) }
     897           1 :       );
     898             : 
     899           4 :       CallInst *NewCall = Builder.CreateCall(
     900             :           ObjectSize, {Src, Intr->getOperand(1), Intr->getOperand(2)});
     901           1 :       Intr->replaceAllUsesWith(NewCall);
     902           1 :       Intr->eraseFromParent();
     903           1 :       continue;
     904             :     }
     905           0 :     default:
     906           0 :       Intr->print(errs());
     907           5 :       llvm_unreachable("Don't know how to promote alloca intrinsic use.");
     908             :     }
     909             :   }
     910             :   return true;
     911             : }
     912             : 
     913        2027 : FunctionPass *llvm::createAMDGPUPromoteAlloca() {
     914        4054 :   return new AMDGPUPromoteAlloca();
     915      303507 : }

Generated by: LCOV version 1.13