1709 lines · cpp
1//===-- AMDGPUPromoteAlloca.cpp - Promote Allocas -------------------------===//2//3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.4// See https://llvm.org/LICENSE.txt for license information.5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception6//7//===----------------------------------------------------------------------===//8//9// Eliminates allocas by either converting them into vectors or by migrating10// them to local address space.11//12// Two passes are exposed by this file:13// - "promote-alloca-to-vector", which runs early in the pipeline and only14// promotes to vector. Promotion to vector is almost always profitable15// except when the alloca is too big and the promotion would result in16// very high register pressure.17// - "promote-alloca", which does both promotion to vector and LDS and runs18// much later in the pipeline. This runs after SROA because promoting to19// LDS is of course less profitable than getting rid of the alloca or20// vectorizing it, thus we only want to do it when the only alternative is21// lowering the alloca to stack.22//23// Note that both of them exist for the old and new PMs. The new PM passes are24// declared in AMDGPU.h and the legacy PM ones are declared here.s25//26//===----------------------------------------------------------------------===//27 28#include "AMDGPU.h"29#include "GCNSubtarget.h"30#include "Utils/AMDGPUBaseInfo.h"31#include "llvm/ADT/STLExtras.h"32#include "llvm/Analysis/CaptureTracking.h"33#include "llvm/Analysis/InstSimplifyFolder.h"34#include "llvm/Analysis/InstructionSimplify.h"35#include "llvm/Analysis/LoopInfo.h"36#include "llvm/Analysis/ValueTracking.h"37#include "llvm/CodeGen/TargetPassConfig.h"38#include "llvm/IR/IRBuilder.h"39#include "llvm/IR/IntrinsicInst.h"40#include "llvm/IR/IntrinsicsAMDGPU.h"41#include "llvm/IR/IntrinsicsR600.h"42#include "llvm/IR/PatternMatch.h"43#include "llvm/InitializePasses.h"44#include "llvm/Pass.h"45#include "llvm/Target/TargetMachine.h"46#include "llvm/Transforms/Utils/SSAUpdater.h"47 48#define DEBUG_TYPE "amdgpu-promote-alloca"49 50using namespace llvm;51 52namespace {53 54static cl::opt<bool>55 DisablePromoteAllocaToVector("disable-promote-alloca-to-vector",56 cl::desc("Disable promote alloca to vector"),57 cl::init(false));58 59static cl::opt<bool>60 DisablePromoteAllocaToLDS("disable-promote-alloca-to-lds",61 cl::desc("Disable promote alloca to LDS"),62 cl::init(false));63 64static cl::opt<unsigned> PromoteAllocaToVectorLimit(65 "amdgpu-promote-alloca-to-vector-limit",66 cl::desc("Maximum byte size to consider promote alloca to vector"),67 cl::init(0));68 69static cl::opt<unsigned> PromoteAllocaToVectorMaxRegs(70 "amdgpu-promote-alloca-to-vector-max-regs",71 cl::desc(72 "Maximum vector size (in 32b registers) to use when promoting alloca"),73 cl::init(32));74 75// Use up to 1/4 of available register budget for vectorization.76// FIXME: Increase the limit for whole function budgets? Perhaps x2?77static cl::opt<unsigned> PromoteAllocaToVectorVGPRRatio(78 "amdgpu-promote-alloca-to-vector-vgpr-ratio",79 cl::desc("Ratio of VGPRs to budget for promoting alloca to vectors"),80 cl::init(4));81 82static cl::opt<unsigned>83 LoopUserWeight("promote-alloca-vector-loop-user-weight",84 cl::desc("The bonus weight of users of allocas within loop "85 "when sorting profitable allocas"),86 cl::init(4));87 88// Shared implementation which can do both promotion to vector and to LDS.89class AMDGPUPromoteAllocaImpl {90private:91 const TargetMachine &TM;92 LoopInfo &LI;93 Module *Mod = nullptr;94 const DataLayout *DL = nullptr;95 96 // FIXME: This should be per-kernel.97 uint32_t LocalMemLimit = 0;98 uint32_t CurrentLocalMemUsage = 0;99 unsigned MaxVGPRs;100 unsigned VGPRBudgetRatio;101 unsigned MaxVectorRegs;102 103 bool IsAMDGCN = false;104 bool IsAMDHSA = false;105 106 std::pair<Value *, Value *> getLocalSizeYZ(IRBuilder<> &Builder);107 Value *getWorkitemID(IRBuilder<> &Builder, unsigned N);108 109 /// BaseAlloca is the alloca root the search started from.110 /// Val may be that alloca or a recursive user of it.111 bool collectUsesWithPtrTypes(Value *BaseAlloca, Value *Val,112 std::vector<Value *> &WorkList) const;113 114 /// Val is a derived pointer from Alloca. OpIdx0/OpIdx1 are the operand115 /// indices to an instruction with 2 pointer inputs (e.g. select, icmp).116 /// Returns true if both operands are derived from the same alloca. Val should117 /// be the same value as one of the input operands of UseInst.118 bool binaryOpIsDerivedFromSameAlloca(Value *Alloca, Value *Val,119 Instruction *UseInst, int OpIdx0,120 int OpIdx1) const;121 122 /// Check whether we have enough local memory for promotion.123 bool hasSufficientLocalMem(const Function &F);124 125 bool tryPromoteAllocaToVector(AllocaInst &I);126 bool tryPromoteAllocaToLDS(AllocaInst &I, bool SufficientLDS);127 128 void sortAllocasToPromote(SmallVectorImpl<AllocaInst *> &Allocas);129 130 void setFunctionLimits(const Function &F);131 132public:133 AMDGPUPromoteAllocaImpl(TargetMachine &TM, LoopInfo &LI) : TM(TM), LI(LI) {134 135 const Triple &TT = TM.getTargetTriple();136 IsAMDGCN = TT.isAMDGCN();137 IsAMDHSA = TT.getOS() == Triple::AMDHSA;138 }139 140 bool run(Function &F, bool PromoteToLDS);141};142 143// FIXME: This can create globals so should be a module pass.144class AMDGPUPromoteAlloca : public FunctionPass {145public:146 static char ID;147 148 AMDGPUPromoteAlloca() : FunctionPass(ID) {}149 150 bool runOnFunction(Function &F) override {151 if (skipFunction(F))152 return false;153 if (auto *TPC = getAnalysisIfAvailable<TargetPassConfig>())154 return AMDGPUPromoteAllocaImpl(155 TPC->getTM<TargetMachine>(),156 getAnalysis<LoopInfoWrapperPass>().getLoopInfo())157 .run(F, /*PromoteToLDS*/ true);158 return false;159 }160 161 StringRef getPassName() const override { return "AMDGPU Promote Alloca"; }162 163 void getAnalysisUsage(AnalysisUsage &AU) const override {164 AU.setPreservesCFG();165 AU.addRequired<LoopInfoWrapperPass>();166 FunctionPass::getAnalysisUsage(AU);167 }168};169 170static unsigned getMaxVGPRs(unsigned LDSBytes, const TargetMachine &TM,171 const Function &F) {172 if (!TM.getTargetTriple().isAMDGCN())173 return 128;174 175 const GCNSubtarget &ST = TM.getSubtarget<GCNSubtarget>(F);176 177 unsigned DynamicVGPRBlockSize = AMDGPU::getDynamicVGPRBlockSize(F);178 // Temporarily check both the attribute and the subtarget feature, until the179 // latter is removed.180 if (DynamicVGPRBlockSize == 0 && ST.isDynamicVGPREnabled())181 DynamicVGPRBlockSize = ST.getDynamicVGPRBlockSize();182 183 unsigned MaxVGPRs = ST.getMaxNumVGPRs(184 ST.getWavesPerEU(ST.getFlatWorkGroupSizes(F), LDSBytes, F).first,185 DynamicVGPRBlockSize);186 187 // A non-entry function has only 32 caller preserved registers.188 // Do not promote alloca which will force spilling unless we know the function189 // will be inlined.190 if (!F.hasFnAttribute(Attribute::AlwaysInline) &&191 !AMDGPU::isEntryFunctionCC(F.getCallingConv()))192 MaxVGPRs = std::min(MaxVGPRs, 32u);193 return MaxVGPRs;194}195 196} // end anonymous namespace197 198char AMDGPUPromoteAlloca::ID = 0;199 200INITIALIZE_PASS_BEGIN(AMDGPUPromoteAlloca, DEBUG_TYPE,201 "AMDGPU promote alloca to vector or LDS", false, false)202// Move LDS uses from functions to kernels before promote alloca for accurate203// estimation of LDS available204INITIALIZE_PASS_DEPENDENCY(AMDGPULowerModuleLDSLegacy)205INITIALIZE_PASS_DEPENDENCY(LoopInfoWrapperPass)206INITIALIZE_PASS_END(AMDGPUPromoteAlloca, DEBUG_TYPE,207 "AMDGPU promote alloca to vector or LDS", false, false)208 209char &llvm::AMDGPUPromoteAllocaID = AMDGPUPromoteAlloca::ID;210 211PreservedAnalyses AMDGPUPromoteAllocaPass::run(Function &F,212 FunctionAnalysisManager &AM) {213 auto &LI = AM.getResult<LoopAnalysis>(F);214 bool Changed = AMDGPUPromoteAllocaImpl(TM, LI).run(F, /*PromoteToLDS=*/true);215 if (Changed) {216 PreservedAnalyses PA;217 PA.preserveSet<CFGAnalyses>();218 return PA;219 }220 return PreservedAnalyses::all();221}222 223PreservedAnalyses224AMDGPUPromoteAllocaToVectorPass::run(Function &F, FunctionAnalysisManager &AM) {225 auto &LI = AM.getResult<LoopAnalysis>(F);226 bool Changed = AMDGPUPromoteAllocaImpl(TM, LI).run(F, /*PromoteToLDS=*/false);227 if (Changed) {228 PreservedAnalyses PA;229 PA.preserveSet<CFGAnalyses>();230 return PA;231 }232 return PreservedAnalyses::all();233}234 235FunctionPass *llvm::createAMDGPUPromoteAlloca() {236 return new AMDGPUPromoteAlloca();237}238 239static void collectAllocaUses(AllocaInst &Alloca,240 SmallVectorImpl<Use *> &Uses) {241 SmallVector<Instruction *, 4> WorkList({&Alloca});242 while (!WorkList.empty()) {243 auto *Cur = WorkList.pop_back_val();244 for (auto &U : Cur->uses()) {245 Uses.push_back(&U);246 247 if (isa<GetElementPtrInst>(U.getUser()))248 WorkList.push_back(cast<Instruction>(U.getUser()));249 }250 }251}252 253void AMDGPUPromoteAllocaImpl::sortAllocasToPromote(254 SmallVectorImpl<AllocaInst *> &Allocas) {255 DenseMap<AllocaInst *, unsigned> Scores;256 257 for (auto *Alloca : Allocas) {258 LLVM_DEBUG(dbgs() << "Scoring: " << *Alloca << "\n");259 unsigned &Score = Scores[Alloca];260 // Increment score by one for each user + a bonus for users within loops.261 SmallVector<Use *, 8> Uses;262 collectAllocaUses(*Alloca, Uses);263 for (auto *U : Uses) {264 Instruction *Inst = cast<Instruction>(U->getUser());265 if (isa<GetElementPtrInst>(Inst))266 continue;267 unsigned UserScore =268 1 + (LoopUserWeight * LI.getLoopDepth(Inst->getParent()));269 LLVM_DEBUG(dbgs() << " [+" << UserScore << "]:\t" << *Inst << "\n");270 Score += UserScore;271 }272 LLVM_DEBUG(dbgs() << " => Final Score:" << Score << "\n");273 }274 275 stable_sort(Allocas, [&](AllocaInst *A, AllocaInst *B) {276 return Scores.at(A) > Scores.at(B);277 });278 279 // clang-format off280 LLVM_DEBUG(281 dbgs() << "Sorted Worklist:\n";282 for (auto *A: Allocas)283 dbgs() << " " << *A << "\n";284 );285 // clang-format on286}287 288void AMDGPUPromoteAllocaImpl::setFunctionLimits(const Function &F) {289 // Load per function limits, overriding with global options where appropriate.290 // R600 register tuples/aliasing are fragile with large vector promotions so291 // apply architecture specific limit here.292 const int R600MaxVectorRegs = 16;293 MaxVectorRegs = F.getFnAttributeAsParsedInteger(294 "amdgpu-promote-alloca-to-vector-max-regs",295 IsAMDGCN ? PromoteAllocaToVectorMaxRegs : R600MaxVectorRegs);296 if (PromoteAllocaToVectorMaxRegs.getNumOccurrences())297 MaxVectorRegs = PromoteAllocaToVectorMaxRegs;298 VGPRBudgetRatio = F.getFnAttributeAsParsedInteger(299 "amdgpu-promote-alloca-to-vector-vgpr-ratio",300 PromoteAllocaToVectorVGPRRatio);301 if (PromoteAllocaToVectorVGPRRatio.getNumOccurrences())302 VGPRBudgetRatio = PromoteAllocaToVectorVGPRRatio;303}304 305bool AMDGPUPromoteAllocaImpl::run(Function &F, bool PromoteToLDS) {306 Mod = F.getParent();307 DL = &Mod->getDataLayout();308 309 const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, F);310 if (!ST.isPromoteAllocaEnabled())311 return false;312 313 bool SufficientLDS = PromoteToLDS && hasSufficientLocalMem(F);314 MaxVGPRs = getMaxVGPRs(CurrentLocalMemUsage, TM, F);315 setFunctionLimits(F);316 317 unsigned VectorizationBudget =318 (PromoteAllocaToVectorLimit ? PromoteAllocaToVectorLimit * 8319 : (MaxVGPRs * 32)) /320 VGPRBudgetRatio;321 322 SmallVector<AllocaInst *, 16> Allocas;323 for (Instruction &I : F.getEntryBlock()) {324 if (AllocaInst *AI = dyn_cast<AllocaInst>(&I)) {325 // Array allocations are probably not worth handling, since an allocation326 // of the array type is the canonical form.327 if (!AI->isStaticAlloca() || AI->isArrayAllocation())328 continue;329 Allocas.push_back(AI);330 }331 }332 333 sortAllocasToPromote(Allocas);334 335 bool Changed = false;336 for (AllocaInst *AI : Allocas) {337 const unsigned AllocaCost = DL->getTypeSizeInBits(AI->getAllocatedType());338 // First, check if we have enough budget to vectorize this alloca.339 if (AllocaCost <= VectorizationBudget) {340 // If we do, attempt vectorization, otherwise, fall through and try341 // promoting to LDS instead.342 if (tryPromoteAllocaToVector(*AI)) {343 Changed = true;344 assert((VectorizationBudget - AllocaCost) < VectorizationBudget &&345 "Underflow!");346 VectorizationBudget -= AllocaCost;347 LLVM_DEBUG(dbgs() << " Remaining vectorization budget:"348 << VectorizationBudget << "\n");349 continue;350 }351 } else {352 LLVM_DEBUG(dbgs() << "Alloca too big for vectorization (size:"353 << AllocaCost << ", budget:" << VectorizationBudget354 << "): " << *AI << "\n");355 }356 357 if (PromoteToLDS && tryPromoteAllocaToLDS(*AI, SufficientLDS))358 Changed = true;359 }360 361 // NOTE: tryPromoteAllocaToVector removes the alloca, so Allocas contains362 // dangling pointers. If we want to reuse it past this point, the loop above363 // would need to be updated to remove successfully promoted allocas.364 365 return Changed;366}367 368struct MemTransferInfo {369 ConstantInt *SrcIndex = nullptr;370 ConstantInt *DestIndex = nullptr;371};372 373// Checks if the instruction I is a memset user of the alloca AI that we can374// deal with. Currently, only non-volatile memsets that affect the whole alloca375// are handled.376static bool isSupportedMemset(MemSetInst *I, AllocaInst *AI,377 const DataLayout &DL) {378 using namespace PatternMatch;379 // For now we only care about non-volatile memsets that affect the whole type380 // (start at index 0 and fill the whole alloca).381 //382 // TODO: Now that we moved to PromoteAlloca we could handle any memsets383 // (except maybe volatile ones?) - we just need to use shufflevector if it384 // only affects a subset of the vector.385 const unsigned Size = DL.getTypeStoreSize(AI->getAllocatedType());386 return I->getOperand(0) == AI &&387 match(I->getOperand(2), m_SpecificInt(Size)) && !I->isVolatile();388}389 390static Value *calculateVectorIndex(391 Value *Ptr, const std::map<GetElementPtrInst *, WeakTrackingVH> &GEPIdx) {392 auto *GEP = dyn_cast<GetElementPtrInst>(Ptr->stripPointerCasts());393 if (!GEP)394 return ConstantInt::getNullValue(Type::getInt32Ty(Ptr->getContext()));395 396 auto I = GEPIdx.find(GEP);397 assert(I != GEPIdx.end() && "Must have entry for GEP!");398 399 Value *IndexValue = I->second;400 assert(IndexValue && "index value missing from GEP index map");401 return IndexValue;402}403 404static Value *GEPToVectorIndex(GetElementPtrInst *GEP, AllocaInst *Alloca,405 Type *VecElemTy, const DataLayout &DL,406 SmallVector<Instruction *> &NewInsts) {407 // TODO: Extracting a "multiple of X" from a GEP might be a useful generic408 // helper.409 LLVMContext &Ctx = GEP->getContext();410 unsigned BW = DL.getIndexTypeSizeInBits(GEP->getType());411 SmallMapVector<Value *, APInt, 4> VarOffsets;412 APInt ConstOffset(BW, 0);413 414 // Walk backwards through nested GEPs to collect both constant and variable415 // offsets, so that nested vector GEP chains can be lowered in one step.416 //417 // Given this IR fragment as input:418 //419 // %0 = alloca [10 x <2 x i32>], align 8, addrspace(5)420 // %1 = getelementptr [10 x <2 x i32>], ptr addrspace(5) %0, i32 0, i32 %j421 // %2 = getelementptr i8, ptr addrspace(5) %1, i32 4422 // %3 = load i32, ptr addrspace(5) %2, align 4423 //424 // Combine both GEP operations in a single pass, producing:425 // BasePtr = %0426 // ConstOffset = 4427 // VarOffsets = { %j -> element_size(<2 x i32>) }428 //429 // That lets us emit a single buffer_load directly into a VGPR, without ever430 // allocating scratch memory for the intermediate pointer.431 Value *CurPtr = GEP;432 while (auto *CurGEP = dyn_cast<GetElementPtrInst>(CurPtr)) {433 if (!CurGEP->collectOffset(DL, BW, VarOffsets, ConstOffset))434 return nullptr;435 436 // Move to the next outer pointer.437 CurPtr = CurGEP->getPointerOperand();438 }439 440 assert(CurPtr == Alloca && "GEP not based on alloca");441 442 int64_t VecElemSize = DL.getTypeAllocSize(VecElemTy);443 if (VarOffsets.size() > 1)444 return nullptr;445 446 APInt IndexQuot;447 int64_t Rem;448 APInt::sdivrem(ConstOffset, VecElemSize, IndexQuot, Rem);449 if (Rem != 0)450 return nullptr;451 if (VarOffsets.size() == 0)452 return ConstantInt::get(Ctx, IndexQuot);453 454 IRBuilder<> Builder(GEP);455 456 const auto &VarOffset = VarOffsets.front();457 APInt OffsetQuot;458 APInt::sdivrem(VarOffset.second, VecElemSize, OffsetQuot, Rem);459 if (Rem != 0 || OffsetQuot.isZero())460 return nullptr;461 462 Value *Offset = VarOffset.first;463 auto *OffsetType = dyn_cast<IntegerType>(Offset->getType());464 if (!OffsetType)465 return nullptr;466 467 if (!OffsetQuot.isOne()) {468 ConstantInt *ConstMul =469 ConstantInt::get(Ctx, OffsetQuot.sext(OffsetType->getBitWidth()));470 Offset = Builder.CreateMul(Offset, ConstMul);471 if (Instruction *NewInst = dyn_cast<Instruction>(Offset))472 NewInsts.push_back(NewInst);473 }474 if (ConstOffset.isZero())475 return Offset;476 477 ConstantInt *ConstIndex =478 ConstantInt::get(Ctx, IndexQuot.sext(OffsetType->getBitWidth()));479 Value *IndexAdd = Builder.CreateAdd(Offset, ConstIndex);480 if (Instruction *NewInst = dyn_cast<Instruction>(IndexAdd))481 NewInsts.push_back(NewInst);482 return IndexAdd;483}484 485/// Promotes a single user of the alloca to a vector form.486///487/// \param Inst Instruction to be promoted.488/// \param DL Module Data Layout.489/// \param VectorTy Vectorized Type.490/// \param VecStoreSize Size of \p VectorTy in bytes.491/// \param ElementSize Size of \p VectorTy element type in bytes.492/// \param TransferInfo MemTransferInst info map.493/// \param GEPVectorIdx GEP -> VectorIdx cache.494/// \param CurVal Current value of the vector (e.g. last stored value)495/// \param[out] DeferredLoads \p Inst is added to this vector if it can't496/// be promoted now. This happens when promoting requires \p497/// CurVal, but \p CurVal is nullptr.498/// \return the stored value if \p Inst would have written to the alloca, or499/// nullptr otherwise.500static Value *promoteAllocaUserToVector(501 Instruction *Inst, const DataLayout &DL, FixedVectorType *VectorTy,502 unsigned VecStoreSize, unsigned ElementSize,503 DenseMap<MemTransferInst *, MemTransferInfo> &TransferInfo,504 std::map<GetElementPtrInst *, WeakTrackingVH> &GEPVectorIdx, Value *CurVal,505 SmallVectorImpl<LoadInst *> &DeferredLoads) {506 // Note: we use InstSimplifyFolder because it can leverage the DataLayout507 // to do more folding, especially in the case of vector splats.508 IRBuilder<InstSimplifyFolder> Builder(Inst->getContext(),509 InstSimplifyFolder(DL));510 Builder.SetInsertPoint(Inst);511 512 const auto GetOrLoadCurrentVectorValue = [&]() -> Value * {513 if (CurVal)514 return CurVal;515 516 // If the current value is not known, insert a dummy load and lower it on517 // the second pass.518 LoadInst *Dummy =519 Builder.CreateLoad(VectorTy, PoisonValue::get(Builder.getPtrTy()),520 "promotealloca.dummyload");521 DeferredLoads.push_back(Dummy);522 return Dummy;523 };524 525 const auto CreateTempPtrIntCast = [&Builder, DL](Value *Val,526 Type *PtrTy) -> Value * {527 assert(DL.getTypeStoreSize(Val->getType()) == DL.getTypeStoreSize(PtrTy));528 const unsigned Size = DL.getTypeStoreSizeInBits(PtrTy);529 if (!PtrTy->isVectorTy())530 return Builder.CreateBitOrPointerCast(Val, Builder.getIntNTy(Size));531 const unsigned NumPtrElts = cast<FixedVectorType>(PtrTy)->getNumElements();532 // If we want to cast to cast, e.g. a <2 x ptr> into a <4 x i32>, we need to533 // first cast the ptr vector to <2 x i64>.534 assert((Size % NumPtrElts == 0) && "Vector size not divisble");535 Type *EltTy = Builder.getIntNTy(Size / NumPtrElts);536 return Builder.CreateBitOrPointerCast(537 Val, FixedVectorType::get(EltTy, NumPtrElts));538 };539 540 Type *VecEltTy = VectorTy->getElementType();541 542 switch (Inst->getOpcode()) {543 case Instruction::Load: {544 // Loads can only be lowered if the value is known.545 if (!CurVal) {546 DeferredLoads.push_back(cast<LoadInst>(Inst));547 return nullptr;548 }549 550 Value *Index = calculateVectorIndex(551 cast<LoadInst>(Inst)->getPointerOperand(), GEPVectorIdx);552 553 // We're loading the full vector.554 Type *AccessTy = Inst->getType();555 TypeSize AccessSize = DL.getTypeStoreSize(AccessTy);556 if (Constant *CI = dyn_cast<Constant>(Index)) {557 if (CI->isZeroValue() && AccessSize == VecStoreSize) {558 if (AccessTy->isPtrOrPtrVectorTy())559 CurVal = CreateTempPtrIntCast(CurVal, AccessTy);560 else if (CurVal->getType()->isPtrOrPtrVectorTy())561 CurVal = CreateTempPtrIntCast(CurVal, CurVal->getType());562 Value *NewVal = Builder.CreateBitOrPointerCast(CurVal, AccessTy);563 Inst->replaceAllUsesWith(NewVal);564 return nullptr;565 }566 }567 568 // Loading a subvector.569 if (isa<FixedVectorType>(AccessTy)) {570 assert(AccessSize.isKnownMultipleOf(DL.getTypeStoreSize(VecEltTy)));571 const unsigned NumLoadedElts = AccessSize / DL.getTypeStoreSize(VecEltTy);572 auto *SubVecTy = FixedVectorType::get(VecEltTy, NumLoadedElts);573 assert(DL.getTypeStoreSize(SubVecTy) == DL.getTypeStoreSize(AccessTy));574 575 Value *SubVec = PoisonValue::get(SubVecTy);576 for (unsigned K = 0; K < NumLoadedElts; ++K) {577 Value *CurIdx =578 Builder.CreateAdd(Index, ConstantInt::get(Index->getType(), K));579 SubVec = Builder.CreateInsertElement(580 SubVec, Builder.CreateExtractElement(CurVal, CurIdx), K);581 }582 583 if (AccessTy->isPtrOrPtrVectorTy())584 SubVec = CreateTempPtrIntCast(SubVec, AccessTy);585 else if (SubVecTy->isPtrOrPtrVectorTy())586 SubVec = CreateTempPtrIntCast(SubVec, SubVecTy);587 588 SubVec = Builder.CreateBitOrPointerCast(SubVec, AccessTy);589 Inst->replaceAllUsesWith(SubVec);590 return nullptr;591 }592 593 // We're loading one element.594 Value *ExtractElement = Builder.CreateExtractElement(CurVal, Index);595 if (AccessTy != VecEltTy)596 ExtractElement = Builder.CreateBitOrPointerCast(ExtractElement, AccessTy);597 598 Inst->replaceAllUsesWith(ExtractElement);599 return nullptr;600 }601 case Instruction::Store: {602 // For stores, it's a bit trickier and it depends on whether we're storing603 // the full vector or not. If we're storing the full vector, we don't need604 // to know the current value. If this is a store of a single element, we605 // need to know the value.606 StoreInst *SI = cast<StoreInst>(Inst);607 Value *Index = calculateVectorIndex(SI->getPointerOperand(), GEPVectorIdx);608 Value *Val = SI->getValueOperand();609 610 // We're storing the full vector, we can handle this without knowing CurVal.611 Type *AccessTy = Val->getType();612 TypeSize AccessSize = DL.getTypeStoreSize(AccessTy);613 if (Constant *CI = dyn_cast<Constant>(Index)) {614 if (CI->isZeroValue() && AccessSize == VecStoreSize) {615 if (AccessTy->isPtrOrPtrVectorTy())616 Val = CreateTempPtrIntCast(Val, AccessTy);617 else if (VectorTy->isPtrOrPtrVectorTy())618 Val = CreateTempPtrIntCast(Val, VectorTy);619 return Builder.CreateBitOrPointerCast(Val, VectorTy);620 }621 }622 623 // Storing a subvector.624 if (isa<FixedVectorType>(AccessTy)) {625 assert(AccessSize.isKnownMultipleOf(DL.getTypeStoreSize(VecEltTy)));626 const unsigned NumWrittenElts =627 AccessSize / DL.getTypeStoreSize(VecEltTy);628 const unsigned NumVecElts = VectorTy->getNumElements();629 auto *SubVecTy = FixedVectorType::get(VecEltTy, NumWrittenElts);630 assert(DL.getTypeStoreSize(SubVecTy) == DL.getTypeStoreSize(AccessTy));631 632 if (SubVecTy->isPtrOrPtrVectorTy())633 Val = CreateTempPtrIntCast(Val, SubVecTy);634 else if (AccessTy->isPtrOrPtrVectorTy())635 Val = CreateTempPtrIntCast(Val, AccessTy);636 637 Val = Builder.CreateBitOrPointerCast(Val, SubVecTy);638 639 Value *CurVec = GetOrLoadCurrentVectorValue();640 for (unsigned K = 0, NumElts = std::min(NumWrittenElts, NumVecElts);641 K < NumElts; ++K) {642 Value *CurIdx =643 Builder.CreateAdd(Index, ConstantInt::get(Index->getType(), K));644 CurVec = Builder.CreateInsertElement(645 CurVec, Builder.CreateExtractElement(Val, K), CurIdx);646 }647 return CurVec;648 }649 650 if (Val->getType() != VecEltTy)651 Val = Builder.CreateBitOrPointerCast(Val, VecEltTy);652 return Builder.CreateInsertElement(GetOrLoadCurrentVectorValue(), Val,653 Index);654 }655 case Instruction::Call: {656 if (auto *MTI = dyn_cast<MemTransferInst>(Inst)) {657 // For memcpy, we need to know curval.658 ConstantInt *Length = cast<ConstantInt>(MTI->getLength());659 unsigned NumCopied = Length->getZExtValue() / ElementSize;660 MemTransferInfo *TI = &TransferInfo[MTI];661 unsigned SrcBegin = TI->SrcIndex->getZExtValue();662 unsigned DestBegin = TI->DestIndex->getZExtValue();663 664 SmallVector<int> Mask;665 for (unsigned Idx = 0; Idx < VectorTy->getNumElements(); ++Idx) {666 if (Idx >= DestBegin && Idx < DestBegin + NumCopied) {667 Mask.push_back(SrcBegin < VectorTy->getNumElements()668 ? SrcBegin++669 : PoisonMaskElem);670 } else {671 Mask.push_back(Idx);672 }673 }674 675 return Builder.CreateShuffleVector(GetOrLoadCurrentVectorValue(), Mask);676 }677 678 if (auto *MSI = dyn_cast<MemSetInst>(Inst)) {679 // For memset, we don't need to know the previous value because we680 // currently only allow memsets that cover the whole alloca.681 Value *Elt = MSI->getOperand(1);682 const unsigned BytesPerElt = DL.getTypeStoreSize(VecEltTy);683 if (BytesPerElt > 1) {684 Value *EltBytes = Builder.CreateVectorSplat(BytesPerElt, Elt);685 686 // If the element type of the vector is a pointer, we need to first cast687 // to an integer, then use a PtrCast.688 if (VecEltTy->isPointerTy()) {689 Type *PtrInt = Builder.getIntNTy(BytesPerElt * 8);690 Elt = Builder.CreateBitCast(EltBytes, PtrInt);691 Elt = Builder.CreateIntToPtr(Elt, VecEltTy);692 } else693 Elt = Builder.CreateBitCast(EltBytes, VecEltTy);694 }695 696 return Builder.CreateVectorSplat(VectorTy->getElementCount(), Elt);697 }698 699 if (auto *Intr = dyn_cast<IntrinsicInst>(Inst)) {700 if (Intr->getIntrinsicID() == Intrinsic::objectsize) {701 Intr->replaceAllUsesWith(702 Builder.getIntN(Intr->getType()->getIntegerBitWidth(),703 DL.getTypeAllocSize(VectorTy)));704 return nullptr;705 }706 }707 708 llvm_unreachable("Unsupported call when promoting alloca to vector");709 }710 711 default:712 llvm_unreachable("Inconsistency in instructions promotable to vector");713 }714 715 llvm_unreachable("Did not return after promoting instruction!");716}717 718static bool isSupportedAccessType(FixedVectorType *VecTy, Type *AccessTy,719 const DataLayout &DL) {720 // Access as a vector type can work if the size of the access vector is a721 // multiple of the size of the alloca's vector element type.722 //723 // Examples:724 // - VecTy = <8 x float>, AccessTy = <4 x float> -> OK725 // - VecTy = <4 x double>, AccessTy = <2 x float> -> OK726 // - VecTy = <4 x double>, AccessTy = <3 x float> -> NOT OK727 // - 3*32 is not a multiple of 64728 //729 // We could handle more complicated cases, but it'd make things a lot more730 // complicated.731 if (isa<FixedVectorType>(AccessTy)) {732 TypeSize AccTS = DL.getTypeStoreSize(AccessTy);733 // If the type size and the store size don't match, we would need to do more734 // than just bitcast to translate between an extracted/insertable subvectors735 // and the accessed value.736 if (AccTS * 8 != DL.getTypeSizeInBits(AccessTy))737 return false;738 TypeSize VecTS = DL.getTypeStoreSize(VecTy->getElementType());739 return AccTS.isKnownMultipleOf(VecTS);740 }741 742 return CastInst::isBitOrNoopPointerCastable(VecTy->getElementType(), AccessTy,743 DL);744}745 746/// Iterates over an instruction worklist that may contain multiple instructions747/// from the same basic block, but in a different order.748template <typename InstContainer>749static void forEachWorkListItem(const InstContainer &WorkList,750 std::function<void(Instruction *)> Fn) {751 // Bucket up uses of the alloca by the block they occur in.752 // This is important because we have to handle multiple defs/uses in a block753 // ourselves: SSAUpdater is purely for cross-block references.754 DenseMap<BasicBlock *, SmallDenseSet<Instruction *>> UsesByBlock;755 for (Instruction *User : WorkList)756 UsesByBlock[User->getParent()].insert(User);757 758 for (Instruction *User : WorkList) {759 BasicBlock *BB = User->getParent();760 auto &BlockUses = UsesByBlock[BB];761 762 // Already processed, skip.763 if (BlockUses.empty())764 continue;765 766 // Only user in the block, directly process it.767 if (BlockUses.size() == 1) {768 Fn(User);769 continue;770 }771 772 // Multiple users in the block, do a linear scan to see users in order.773 for (Instruction &Inst : *BB) {774 if (!BlockUses.contains(&Inst))775 continue;776 777 Fn(&Inst);778 }779 780 // Clear the block so we know it's been processed.781 BlockUses.clear();782 }783}784 785/// Find an insert point after an alloca, after all other allocas clustered at786/// the start of the block.787static BasicBlock::iterator skipToNonAllocaInsertPt(BasicBlock &BB,788 BasicBlock::iterator I) {789 for (BasicBlock::iterator E = BB.end(); I != E && isa<AllocaInst>(*I); ++I)790 ;791 return I;792}793 794// FIXME: Should try to pick the most likely to be profitable allocas first.795bool AMDGPUPromoteAllocaImpl::tryPromoteAllocaToVector(AllocaInst &Alloca) {796 LLVM_DEBUG(dbgs() << "Trying to promote to vector: " << Alloca << '\n');797 798 if (DisablePromoteAllocaToVector) {799 LLVM_DEBUG(dbgs() << " Promote alloca to vector is disabled\n");800 return false;801 }802 803 Type *AllocaTy = Alloca.getAllocatedType();804 auto *VectorTy = dyn_cast<FixedVectorType>(AllocaTy);805 if (auto *ArrayTy = dyn_cast<ArrayType>(AllocaTy)) {806 uint64_t NumElems = 1;807 Type *ElemTy;808 do {809 NumElems *= ArrayTy->getNumElements();810 ElemTy = ArrayTy->getElementType();811 } while ((ArrayTy = dyn_cast<ArrayType>(ElemTy)));812 813 // Check for array of vectors814 auto *InnerVectorTy = dyn_cast<FixedVectorType>(ElemTy);815 if (InnerVectorTy) {816 NumElems *= InnerVectorTy->getNumElements();817 ElemTy = InnerVectorTy->getElementType();818 }819 820 if (VectorType::isValidElementType(ElemTy) && NumElems > 0) {821 unsigned ElementSize = DL->getTypeSizeInBits(ElemTy) / 8;822 if (ElementSize > 0) {823 unsigned AllocaSize = DL->getTypeStoreSize(AllocaTy);824 // Expand vector if required to match padding of inner type,825 // i.e. odd size subvectors.826 // Storage size of new vector must match that of alloca for correct827 // behaviour of byte offsets and GEP computation.828 if (NumElems * ElementSize != AllocaSize)829 NumElems = AllocaSize / ElementSize;830 if (NumElems > 0 && (AllocaSize % ElementSize) == 0)831 VectorTy = FixedVectorType::get(ElemTy, NumElems);832 }833 }834 }835 836 if (!VectorTy) {837 LLVM_DEBUG(dbgs() << " Cannot convert type to vector\n");838 return false;839 }840 841 const unsigned MaxElements =842 (MaxVectorRegs * 32) / DL->getTypeSizeInBits(VectorTy->getElementType());843 844 if (VectorTy->getNumElements() > MaxElements ||845 VectorTy->getNumElements() < 2) {846 LLVM_DEBUG(dbgs() << " " << *VectorTy847 << " has an unsupported number of elements\n");848 return false;849 }850 851 std::map<GetElementPtrInst *, WeakTrackingVH> GEPVectorIdx;852 SmallVector<Instruction *> WorkList;853 SmallVector<Instruction *> UsersToRemove;854 SmallVector<Instruction *> DeferredInsts;855 SmallVector<Instruction *> NewGEPInsts;856 DenseMap<MemTransferInst *, MemTransferInfo> TransferInfo;857 858 const auto RejectUser = [&](Instruction *Inst, Twine Msg) {859 LLVM_DEBUG(dbgs() << " Cannot promote alloca to vector: " << Msg << "\n"860 << " " << *Inst << "\n");861 for (auto *Inst : reverse(NewGEPInsts))862 Inst->eraseFromParent();863 return false;864 };865 866 SmallVector<Use *, 8> Uses;867 collectAllocaUses(Alloca, Uses);868 869 LLVM_DEBUG(dbgs() << " Attempting promotion to: " << *VectorTy << "\n");870 871 Type *VecEltTy = VectorTy->getElementType();872 unsigned ElementSizeInBits = DL->getTypeSizeInBits(VecEltTy);873 if (ElementSizeInBits != DL->getTypeAllocSizeInBits(VecEltTy)) {874 LLVM_DEBUG(dbgs() << " Cannot convert to vector if the allocation size "875 "does not match the type's size\n");876 return false;877 }878 unsigned ElementSize = ElementSizeInBits / 8;879 assert(ElementSize > 0);880 for (auto *U : Uses) {881 Instruction *Inst = cast<Instruction>(U->getUser());882 883 if (Value *Ptr = getLoadStorePointerOperand(Inst)) {884 // This is a store of the pointer, not to the pointer.885 if (isa<StoreInst>(Inst) &&886 U->getOperandNo() != StoreInst::getPointerOperandIndex())887 return RejectUser(Inst, "pointer is being stored");888 889 Type *AccessTy = getLoadStoreType(Inst);890 if (AccessTy->isAggregateType())891 return RejectUser(Inst, "unsupported load/store as aggregate");892 assert(!AccessTy->isAggregateType() || AccessTy->isArrayTy());893 894 // Check that this is a simple access of a vector element.895 bool IsSimple = isa<LoadInst>(Inst) ? cast<LoadInst>(Inst)->isSimple()896 : cast<StoreInst>(Inst)->isSimple();897 if (!IsSimple)898 return RejectUser(Inst, "not a simple load or store");899 900 Ptr = Ptr->stripPointerCasts();901 902 // Alloca already accessed as vector.903 if (Ptr == &Alloca && DL->getTypeStoreSize(Alloca.getAllocatedType()) ==904 DL->getTypeStoreSize(AccessTy)) {905 WorkList.push_back(Inst);906 continue;907 }908 909 if (!isSupportedAccessType(VectorTy, AccessTy, *DL))910 return RejectUser(Inst, "not a supported access type");911 912 WorkList.push_back(Inst);913 continue;914 }915 916 if (auto *GEP = dyn_cast<GetElementPtrInst>(Inst)) {917 // If we can't compute a vector index from this GEP, then we can't918 // promote this alloca to vector.919 Value *Index = GEPToVectorIndex(GEP, &Alloca, VecEltTy, *DL, NewGEPInsts);920 if (!Index)921 return RejectUser(Inst, "cannot compute vector index for GEP");922 923 GEPVectorIdx[GEP] = Index;924 UsersToRemove.push_back(Inst);925 continue;926 }927 928 if (MemSetInst *MSI = dyn_cast<MemSetInst>(Inst);929 MSI && isSupportedMemset(MSI, &Alloca, *DL)) {930 WorkList.push_back(Inst);931 continue;932 }933 934 if (MemTransferInst *TransferInst = dyn_cast<MemTransferInst>(Inst)) {935 if (TransferInst->isVolatile())936 return RejectUser(Inst, "mem transfer inst is volatile");937 938 ConstantInt *Len = dyn_cast<ConstantInt>(TransferInst->getLength());939 if (!Len || (Len->getZExtValue() % ElementSize))940 return RejectUser(Inst, "mem transfer inst length is non-constant or "941 "not a multiple of the vector element size");942 943 if (TransferInfo.try_emplace(TransferInst).second) {944 DeferredInsts.push_back(Inst);945 WorkList.push_back(Inst);946 }947 948 auto getPointerIndexOfAlloca = [&](Value *Ptr) -> ConstantInt * {949 GetElementPtrInst *GEP = dyn_cast<GetElementPtrInst>(Ptr);950 if (Ptr != &Alloca && !GEPVectorIdx.count(GEP))951 return nullptr;952 953 return dyn_cast<ConstantInt>(calculateVectorIndex(Ptr, GEPVectorIdx));954 };955 956 unsigned OpNum = U->getOperandNo();957 MemTransferInfo *TI = &TransferInfo[TransferInst];958 if (OpNum == 0) {959 Value *Dest = TransferInst->getDest();960 ConstantInt *Index = getPointerIndexOfAlloca(Dest);961 if (!Index)962 return RejectUser(Inst, "could not calculate constant dest index");963 TI->DestIndex = Index;964 } else {965 assert(OpNum == 1);966 Value *Src = TransferInst->getSource();967 ConstantInt *Index = getPointerIndexOfAlloca(Src);968 if (!Index)969 return RejectUser(Inst, "could not calculate constant src index");970 TI->SrcIndex = Index;971 }972 continue;973 }974 975 if (auto *Intr = dyn_cast<IntrinsicInst>(Inst)) {976 if (Intr->getIntrinsicID() == Intrinsic::objectsize) {977 WorkList.push_back(Inst);978 continue;979 }980 }981 982 // Ignore assume-like intrinsics and comparisons used in assumes.983 if (isAssumeLikeIntrinsic(Inst)) {984 if (!Inst->use_empty())985 return RejectUser(Inst, "assume-like intrinsic cannot have any users");986 UsersToRemove.push_back(Inst);987 continue;988 }989 990 if (isa<ICmpInst>(Inst) && all_of(Inst->users(), [](User *U) {991 return isAssumeLikeIntrinsic(cast<Instruction>(U));992 })) {993 UsersToRemove.push_back(Inst);994 continue;995 }996 997 return RejectUser(Inst, "unhandled alloca user");998 }999 1000 while (!DeferredInsts.empty()) {1001 Instruction *Inst = DeferredInsts.pop_back_val();1002 MemTransferInst *TransferInst = cast<MemTransferInst>(Inst);1003 // TODO: Support the case if the pointers are from different alloca or1004 // from different address spaces.1005 MemTransferInfo &Info = TransferInfo[TransferInst];1006 if (!Info.SrcIndex || !Info.DestIndex)1007 return RejectUser(1008 Inst, "mem transfer inst is missing constant src and/or dst index");1009 }1010 1011 LLVM_DEBUG(dbgs() << " Converting alloca to vector " << *AllocaTy << " -> "1012 << *VectorTy << '\n');1013 const unsigned VecStoreSize = DL->getTypeStoreSize(VectorTy);1014 1015 // Alloca is uninitialized memory. Imitate that by making the first value1016 // undef.1017 SSAUpdater Updater;1018 Updater.Initialize(VectorTy, "promotealloca");1019 1020 BasicBlock *EntryBB = Alloca.getParent();1021 BasicBlock::iterator InitInsertPos =1022 skipToNonAllocaInsertPt(*EntryBB, Alloca.getIterator());1023 // Alloca memory is undefined to begin, not poison.1024 Value *AllocaInitValue =1025 new FreezeInst(PoisonValue::get(VectorTy), "", InitInsertPos);1026 AllocaInitValue->takeName(&Alloca);1027 1028 Updater.AddAvailableValue(EntryBB, AllocaInitValue);1029 1030 // First handle the initial worklist.1031 SmallVector<LoadInst *, 4> DeferredLoads;1032 forEachWorkListItem(WorkList, [&](Instruction *I) {1033 BasicBlock *BB = I->getParent();1034 // On the first pass, we only take values that are trivially known, i.e.1035 // where AddAvailableValue was already called in this block.1036 Value *Result = promoteAllocaUserToVector(1037 I, *DL, VectorTy, VecStoreSize, ElementSize, TransferInfo, GEPVectorIdx,1038 Updater.FindValueForBlock(BB), DeferredLoads);1039 if (Result)1040 Updater.AddAvailableValue(BB, Result);1041 });1042 1043 // Then handle deferred loads.1044 forEachWorkListItem(DeferredLoads, [&](Instruction *I) {1045 SmallVector<LoadInst *, 0> NewDLs;1046 BasicBlock *BB = I->getParent();1047 // On the second pass, we use GetValueInMiddleOfBlock to guarantee we always1048 // get a value, inserting PHIs as needed.1049 Value *Result = promoteAllocaUserToVector(1050 I, *DL, VectorTy, VecStoreSize, ElementSize, TransferInfo, GEPVectorIdx,1051 Updater.GetValueInMiddleOfBlock(I->getParent()), NewDLs);1052 if (Result)1053 Updater.AddAvailableValue(BB, Result);1054 assert(NewDLs.empty() && "No more deferred loads should be queued!");1055 });1056 1057 // Delete all instructions. On the first pass, new dummy loads may have been1058 // added so we need to collect them too.1059 DenseSet<Instruction *> InstsToDelete(WorkList.begin(), WorkList.end());1060 InstsToDelete.insert_range(DeferredLoads);1061 for (Instruction *I : InstsToDelete) {1062 assert(I->use_empty());1063 I->eraseFromParent();1064 }1065 1066 // Delete all the users that are known to be removeable.1067 for (Instruction *I : reverse(UsersToRemove)) {1068 I->dropDroppableUses();1069 assert(I->use_empty());1070 I->eraseFromParent();1071 }1072 1073 // Alloca should now be dead too.1074 assert(Alloca.use_empty());1075 Alloca.eraseFromParent();1076 return true;1077}1078 1079std::pair<Value *, Value *>1080AMDGPUPromoteAllocaImpl::getLocalSizeYZ(IRBuilder<> &Builder) {1081 Function &F = *Builder.GetInsertBlock()->getParent();1082 const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, F);1083 1084 if (!IsAMDHSA) {1085 CallInst *LocalSizeY =1086 Builder.CreateIntrinsic(Intrinsic::r600_read_local_size_y, {});1087 CallInst *LocalSizeZ =1088 Builder.CreateIntrinsic(Intrinsic::r600_read_local_size_z, {});1089 1090 ST.makeLIDRangeMetadata(LocalSizeY);1091 ST.makeLIDRangeMetadata(LocalSizeZ);1092 1093 return std::pair(LocalSizeY, LocalSizeZ);1094 }1095 1096 // We must read the size out of the dispatch pointer.1097 assert(IsAMDGCN);1098 1099 // We are indexing into this struct, and want to extract the workgroup_size_*1100 // fields.1101 //1102 // typedef struct hsa_kernel_dispatch_packet_s {1103 // uint16_t header;1104 // uint16_t setup;1105 // uint16_t workgroup_size_x ;1106 // uint16_t workgroup_size_y;1107 // uint16_t workgroup_size_z;1108 // uint16_t reserved0;1109 // uint32_t grid_size_x ;1110 // uint32_t grid_size_y ;1111 // uint32_t grid_size_z;1112 //1113 // uint32_t private_segment_size;1114 // uint32_t group_segment_size;1115 // uint64_t kernel_object;1116 //1117 // #ifdef HSA_LARGE_MODEL1118 // void *kernarg_address;1119 // #elif defined HSA_LITTLE_ENDIAN1120 // void *kernarg_address;1121 // uint32_t reserved1;1122 // #else1123 // uint32_t reserved1;1124 // void *kernarg_address;1125 // #endif1126 // uint64_t reserved2;1127 // hsa_signal_t completion_signal; // uint64_t wrapper1128 // } hsa_kernel_dispatch_packet_t1129 //1130 CallInst *DispatchPtr =1131 Builder.CreateIntrinsic(Intrinsic::amdgcn_dispatch_ptr, {});1132 DispatchPtr->addRetAttr(Attribute::NoAlias);1133 DispatchPtr->addRetAttr(Attribute::NonNull);1134 F.removeFnAttr("amdgpu-no-dispatch-ptr");1135 1136 // Size of the dispatch packet struct.1137 DispatchPtr->addDereferenceableRetAttr(64);1138 1139 Type *I32Ty = Type::getInt32Ty(Mod->getContext());1140 1141 // We could do a single 64-bit load here, but it's likely that the basic1142 // 32-bit and extract sequence is already present, and it is probably easier1143 // to CSE this. The loads should be mergeable later anyway.1144 Value *GEPXY = Builder.CreateConstInBoundsGEP1_64(I32Ty, DispatchPtr, 1);1145 LoadInst *LoadXY = Builder.CreateAlignedLoad(I32Ty, GEPXY, Align(4));1146 1147 Value *GEPZU = Builder.CreateConstInBoundsGEP1_64(I32Ty, DispatchPtr, 2);1148 LoadInst *LoadZU = Builder.CreateAlignedLoad(I32Ty, GEPZU, Align(4));1149 1150 MDNode *MD = MDNode::get(Mod->getContext(), {});1151 LoadXY->setMetadata(LLVMContext::MD_invariant_load, MD);1152 LoadZU->setMetadata(LLVMContext::MD_invariant_load, MD);1153 ST.makeLIDRangeMetadata(LoadZU);1154 1155 // Extract y component. Upper half of LoadZU should be zero already.1156 Value *Y = Builder.CreateLShr(LoadXY, 16);1157 1158 return std::pair(Y, LoadZU);1159}1160 1161Value *AMDGPUPromoteAllocaImpl::getWorkitemID(IRBuilder<> &Builder,1162 unsigned N) {1163 Function *F = Builder.GetInsertBlock()->getParent();1164 const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, *F);1165 Intrinsic::ID IntrID = Intrinsic::not_intrinsic;1166 StringRef AttrName;1167 1168 switch (N) {1169 case 0:1170 IntrID = IsAMDGCN ? (Intrinsic::ID)Intrinsic::amdgcn_workitem_id_x1171 : (Intrinsic::ID)Intrinsic::r600_read_tidig_x;1172 AttrName = "amdgpu-no-workitem-id-x";1173 break;1174 case 1:1175 IntrID = IsAMDGCN ? (Intrinsic::ID)Intrinsic::amdgcn_workitem_id_y1176 : (Intrinsic::ID)Intrinsic::r600_read_tidig_y;1177 AttrName = "amdgpu-no-workitem-id-y";1178 break;1179 1180 case 2:1181 IntrID = IsAMDGCN ? (Intrinsic::ID)Intrinsic::amdgcn_workitem_id_z1182 : (Intrinsic::ID)Intrinsic::r600_read_tidig_z;1183 AttrName = "amdgpu-no-workitem-id-z";1184 break;1185 default:1186 llvm_unreachable("invalid dimension");1187 }1188 1189 Function *WorkitemIdFn = Intrinsic::getOrInsertDeclaration(Mod, IntrID);1190 CallInst *CI = Builder.CreateCall(WorkitemIdFn);1191 ST.makeLIDRangeMetadata(CI);1192 F->removeFnAttr(AttrName);1193 1194 return CI;1195}1196 1197static bool isCallPromotable(CallInst *CI) {1198 IntrinsicInst *II = dyn_cast<IntrinsicInst>(CI);1199 if (!II)1200 return false;1201 1202 switch (II->getIntrinsicID()) {1203 case Intrinsic::memcpy:1204 case Intrinsic::memmove:1205 case Intrinsic::memset:1206 case Intrinsic::lifetime_start:1207 case Intrinsic::lifetime_end:1208 case Intrinsic::invariant_start:1209 case Intrinsic::invariant_end:1210 case Intrinsic::launder_invariant_group:1211 case Intrinsic::strip_invariant_group:1212 case Intrinsic::objectsize:1213 return true;1214 default:1215 return false;1216 }1217}1218 1219bool AMDGPUPromoteAllocaImpl::binaryOpIsDerivedFromSameAlloca(1220 Value *BaseAlloca, Value *Val, Instruction *Inst, int OpIdx0,1221 int OpIdx1) const {1222 // Figure out which operand is the one we might not be promoting.1223 Value *OtherOp = Inst->getOperand(OpIdx0);1224 if (Val == OtherOp)1225 OtherOp = Inst->getOperand(OpIdx1);1226 1227 if (isa<ConstantPointerNull, ConstantAggregateZero>(OtherOp))1228 return true;1229 1230 // TODO: getUnderlyingObject will not work on a vector getelementptr1231 Value *OtherObj = getUnderlyingObject(OtherOp);1232 if (!isa<AllocaInst>(OtherObj))1233 return false;1234 1235 // TODO: We should be able to replace undefs with the right pointer type.1236 1237 // TODO: If we know the other base object is another promotable1238 // alloca, not necessarily this alloca, we can do this. The1239 // important part is both must have the same address space at1240 // the end.1241 if (OtherObj != BaseAlloca) {1242 LLVM_DEBUG(1243 dbgs() << "Found a binary instruction with another alloca object\n");1244 return false;1245 }1246 1247 return true;1248}1249 1250bool AMDGPUPromoteAllocaImpl::collectUsesWithPtrTypes(1251 Value *BaseAlloca, Value *Val, std::vector<Value *> &WorkList) const {1252 1253 for (User *User : Val->users()) {1254 if (is_contained(WorkList, User))1255 continue;1256 1257 if (CallInst *CI = dyn_cast<CallInst>(User)) {1258 if (!isCallPromotable(CI))1259 return false;1260 1261 WorkList.push_back(User);1262 continue;1263 }1264 1265 Instruction *UseInst = cast<Instruction>(User);1266 if (UseInst->getOpcode() == Instruction::PtrToInt)1267 return false;1268 1269 if (LoadInst *LI = dyn_cast<LoadInst>(UseInst)) {1270 if (LI->isVolatile())1271 return false;1272 continue;1273 }1274 1275 if (StoreInst *SI = dyn_cast<StoreInst>(UseInst)) {1276 if (SI->isVolatile())1277 return false;1278 1279 // Reject if the stored value is not the pointer operand.1280 if (SI->getPointerOperand() != Val)1281 return false;1282 continue;1283 }1284 1285 if (AtomicRMWInst *RMW = dyn_cast<AtomicRMWInst>(UseInst)) {1286 if (RMW->isVolatile())1287 return false;1288 continue;1289 }1290 1291 if (AtomicCmpXchgInst *CAS = dyn_cast<AtomicCmpXchgInst>(UseInst)) {1292 if (CAS->isVolatile())1293 return false;1294 continue;1295 }1296 1297 // Only promote a select if we know that the other select operand1298 // is from another pointer that will also be promoted.1299 if (ICmpInst *ICmp = dyn_cast<ICmpInst>(UseInst)) {1300 if (!binaryOpIsDerivedFromSameAlloca(BaseAlloca, Val, ICmp, 0, 1))1301 return false;1302 1303 // May need to rewrite constant operands.1304 WorkList.push_back(ICmp);1305 continue;1306 }1307 1308 if (GetElementPtrInst *GEP = dyn_cast<GetElementPtrInst>(UseInst)) {1309 // Be conservative if an address could be computed outside the bounds of1310 // the alloca.1311 if (!GEP->isInBounds())1312 return false;1313 } else if (SelectInst *SI = dyn_cast<SelectInst>(UseInst)) {1314 // Only promote a select if we know that the other select operand is from1315 // another pointer that will also be promoted.1316 if (!binaryOpIsDerivedFromSameAlloca(BaseAlloca, Val, SI, 1, 2))1317 return false;1318 } else if (PHINode *Phi = dyn_cast<PHINode>(UseInst)) {1319 // Repeat for phis.1320 1321 // TODO: Handle more complex cases. We should be able to replace loops1322 // over arrays.1323 switch (Phi->getNumIncomingValues()) {1324 case 1:1325 break;1326 case 2:1327 if (!binaryOpIsDerivedFromSameAlloca(BaseAlloca, Val, Phi, 0, 1))1328 return false;1329 break;1330 default:1331 return false;1332 }1333 } else if (!isa<ExtractElementInst>(User)) {1334 // Do not promote vector/aggregate type instructions. It is hard to track1335 // their users.1336 1337 // Do not promote addrspacecast.1338 //1339 // TODO: If we know the address is only observed through flat pointers, we1340 // could still promote.1341 return false;1342 }1343 1344 WorkList.push_back(User);1345 if (!collectUsesWithPtrTypes(BaseAlloca, User, WorkList))1346 return false;1347 }1348 1349 return true;1350}1351 1352bool AMDGPUPromoteAllocaImpl::hasSufficientLocalMem(const Function &F) {1353 1354 FunctionType *FTy = F.getFunctionType();1355 const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, F);1356 1357 // If the function has any arguments in the local address space, then it's1358 // possible these arguments require the entire local memory space, so1359 // we cannot use local memory in the pass.1360 for (Type *ParamTy : FTy->params()) {1361 PointerType *PtrTy = dyn_cast<PointerType>(ParamTy);1362 if (PtrTy && PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) {1363 LocalMemLimit = 0;1364 LLVM_DEBUG(dbgs() << "Function has local memory argument. Promoting to "1365 "local memory disabled.\n");1366 return false;1367 }1368 }1369 1370 LocalMemLimit = ST.getAddressableLocalMemorySize();1371 if (LocalMemLimit == 0)1372 return false;1373 1374 SmallVector<const Constant *, 16> Stack;1375 SmallPtrSet<const Constant *, 8> VisitedConstants;1376 SmallPtrSet<const GlobalVariable *, 8> UsedLDS;1377 1378 auto visitUsers = [&](const GlobalVariable *GV, const Constant *Val) -> bool {1379 for (const User *U : Val->users()) {1380 if (const Instruction *Use = dyn_cast<Instruction>(U)) {1381 if (Use->getFunction() == &F)1382 return true;1383 } else {1384 const Constant *C = cast<Constant>(U);1385 if (VisitedConstants.insert(C).second)1386 Stack.push_back(C);1387 }1388 }1389 1390 return false;1391 };1392 1393 for (GlobalVariable &GV : Mod->globals()) {1394 if (GV.getAddressSpace() != AMDGPUAS::LOCAL_ADDRESS)1395 continue;1396 1397 if (visitUsers(&GV, &GV)) {1398 UsedLDS.insert(&GV);1399 Stack.clear();1400 continue;1401 }1402 1403 // For any ConstantExpr uses, we need to recursively search the users until1404 // we see a function.1405 while (!Stack.empty()) {1406 const Constant *C = Stack.pop_back_val();1407 if (visitUsers(&GV, C)) {1408 UsedLDS.insert(&GV);1409 Stack.clear();1410 break;1411 }1412 }1413 }1414 1415 const DataLayout &DL = Mod->getDataLayout();1416 SmallVector<std::pair<uint64_t, Align>, 16> AllocatedSizes;1417 AllocatedSizes.reserve(UsedLDS.size());1418 1419 for (const GlobalVariable *GV : UsedLDS) {1420 Align Alignment =1421 DL.getValueOrABITypeAlignment(GV->getAlign(), GV->getValueType());1422 uint64_t AllocSize = DL.getTypeAllocSize(GV->getValueType());1423 1424 // HIP uses an extern unsized array in local address space for dynamically1425 // allocated shared memory. In that case, we have to disable the promotion.1426 if (GV->hasExternalLinkage() && AllocSize == 0) {1427 LocalMemLimit = 0;1428 LLVM_DEBUG(dbgs() << "Function has a reference to externally allocated "1429 "local memory. Promoting to local memory "1430 "disabled.\n");1431 return false;1432 }1433 1434 AllocatedSizes.emplace_back(AllocSize, Alignment);1435 }1436 1437 // Sort to try to estimate the worst case alignment padding1438 //1439 // FIXME: We should really do something to fix the addresses to a more optimal1440 // value instead1441 llvm::sort(AllocatedSizes, llvm::less_second());1442 1443 // Check how much local memory is being used by global objects1444 CurrentLocalMemUsage = 0;1445 1446 // FIXME: Try to account for padding here. The real padding and address is1447 // currently determined from the inverse order of uses in the function when1448 // legalizing, which could also potentially change. We try to estimate the1449 // worst case here, but we probably should fix the addresses earlier.1450 for (auto Alloc : AllocatedSizes) {1451 CurrentLocalMemUsage = alignTo(CurrentLocalMemUsage, Alloc.second);1452 CurrentLocalMemUsage += Alloc.first;1453 }1454 1455 unsigned MaxOccupancy =1456 ST.getWavesPerEU(ST.getFlatWorkGroupSizes(F), CurrentLocalMemUsage, F)1457 .second;1458 1459 // Round up to the next tier of usage.1460 unsigned MaxSizeWithWaveCount =1461 ST.getMaxLocalMemSizeWithWaveCount(MaxOccupancy, F);1462 1463 // Program may already use more LDS than is usable at maximum occupancy.1464 if (CurrentLocalMemUsage > MaxSizeWithWaveCount)1465 return false;1466 1467 LocalMemLimit = MaxSizeWithWaveCount;1468 1469 LLVM_DEBUG(dbgs() << F.getName() << " uses " << CurrentLocalMemUsage1470 << " bytes of LDS\n"1471 << " Rounding size to " << MaxSizeWithWaveCount1472 << " with a maximum occupancy of " << MaxOccupancy << '\n'1473 << " and " << (LocalMemLimit - CurrentLocalMemUsage)1474 << " available for promotion\n");1475 1476 return true;1477}1478 1479// FIXME: Should try to pick the most likely to be profitable allocas first.1480bool AMDGPUPromoteAllocaImpl::tryPromoteAllocaToLDS(AllocaInst &I,1481 bool SufficientLDS) {1482 LLVM_DEBUG(dbgs() << "Trying to promote to LDS: " << I << '\n');1483 1484 if (DisablePromoteAllocaToLDS) {1485 LLVM_DEBUG(dbgs() << " Promote alloca to LDS is disabled\n");1486 return false;1487 }1488 1489 const DataLayout &DL = Mod->getDataLayout();1490 IRBuilder<> Builder(&I);1491 1492 const Function &ContainingFunction = *I.getFunction();1493 CallingConv::ID CC = ContainingFunction.getCallingConv();1494 1495 // Don't promote the alloca to LDS for shader calling conventions as the work1496 // item ID intrinsics are not supported for these calling conventions.1497 // Furthermore not all LDS is available for some of the stages.1498 switch (CC) {1499 case CallingConv::AMDGPU_KERNEL:1500 case CallingConv::SPIR_KERNEL:1501 break;1502 default:1503 LLVM_DEBUG(1504 dbgs()1505 << " promote alloca to LDS not supported with calling convention.\n");1506 return false;1507 }1508 1509 // Not likely to have sufficient local memory for promotion.1510 if (!SufficientLDS)1511 return false;1512 1513 const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, ContainingFunction);1514 unsigned WorkGroupSize = ST.getFlatWorkGroupSizes(ContainingFunction).second;1515 1516 Align Alignment =1517 DL.getValueOrABITypeAlignment(I.getAlign(), I.getAllocatedType());1518 1519 // FIXME: This computed padding is likely wrong since it depends on inverse1520 // usage order.1521 //1522 // FIXME: It is also possible that if we're allowed to use all of the memory1523 // could end up using more than the maximum due to alignment padding.1524 1525 uint32_t NewSize = alignTo(CurrentLocalMemUsage, Alignment);1526 uint32_t AllocSize =1527 WorkGroupSize * DL.getTypeAllocSize(I.getAllocatedType());1528 NewSize += AllocSize;1529 1530 if (NewSize > LocalMemLimit) {1531 LLVM_DEBUG(dbgs() << " " << AllocSize1532 << " bytes of local memory not available to promote\n");1533 return false;1534 }1535 1536 CurrentLocalMemUsage = NewSize;1537 1538 std::vector<Value *> WorkList;1539 1540 if (!collectUsesWithPtrTypes(&I, &I, WorkList)) {1541 LLVM_DEBUG(dbgs() << " Do not know how to convert all uses\n");1542 return false;1543 }1544 1545 LLVM_DEBUG(dbgs() << "Promoting alloca to local memory\n");1546 1547 Function *F = I.getFunction();1548 1549 Type *GVTy = ArrayType::get(I.getAllocatedType(), WorkGroupSize);1550 GlobalVariable *GV = new GlobalVariable(1551 *Mod, GVTy, false, GlobalValue::InternalLinkage, PoisonValue::get(GVTy),1552 Twine(F->getName()) + Twine('.') + I.getName(), nullptr,1553 GlobalVariable::NotThreadLocal, AMDGPUAS::LOCAL_ADDRESS);1554 GV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);1555 GV->setAlignment(I.getAlign());1556 1557 Value *TCntY, *TCntZ;1558 1559 std::tie(TCntY, TCntZ) = getLocalSizeYZ(Builder);1560 Value *TIdX = getWorkitemID(Builder, 0);1561 Value *TIdY = getWorkitemID(Builder, 1);1562 Value *TIdZ = getWorkitemID(Builder, 2);1563 1564 Value *Tmp0 = Builder.CreateMul(TCntY, TCntZ, "", true, true);1565 Tmp0 = Builder.CreateMul(Tmp0, TIdX);1566 Value *Tmp1 = Builder.CreateMul(TIdY, TCntZ, "", true, true);1567 Value *TID = Builder.CreateAdd(Tmp0, Tmp1);1568 TID = Builder.CreateAdd(TID, TIdZ);1569 1570 LLVMContext &Context = Mod->getContext();1571 Value *Indices[] = {Constant::getNullValue(Type::getInt32Ty(Context)), TID};1572 1573 Value *Offset = Builder.CreateInBoundsGEP(GVTy, GV, Indices);1574 I.mutateType(Offset->getType());1575 I.replaceAllUsesWith(Offset);1576 I.eraseFromParent();1577 1578 SmallVector<IntrinsicInst *> DeferredIntrs;1579 1580 PointerType *NewPtrTy = PointerType::get(Context, AMDGPUAS::LOCAL_ADDRESS);1581 1582 for (Value *V : WorkList) {1583 CallInst *Call = dyn_cast<CallInst>(V);1584 if (!Call) {1585 if (ICmpInst *CI = dyn_cast<ICmpInst>(V)) {1586 Value *LHS = CI->getOperand(0);1587 Value *RHS = CI->getOperand(1);1588 1589 Type *NewTy = LHS->getType()->getWithNewType(NewPtrTy);1590 if (isa<ConstantPointerNull, ConstantAggregateZero>(LHS))1591 CI->setOperand(0, Constant::getNullValue(NewTy));1592 1593 if (isa<ConstantPointerNull, ConstantAggregateZero>(RHS))1594 CI->setOperand(1, Constant::getNullValue(NewTy));1595 1596 continue;1597 }1598 1599 // The operand's value should be corrected on its own and we don't want to1600 // touch the users.1601 if (isa<AddrSpaceCastInst>(V))1602 continue;1603 1604 assert(V->getType()->isPtrOrPtrVectorTy());1605 1606 Type *NewTy = V->getType()->getWithNewType(NewPtrTy);1607 V->mutateType(NewTy);1608 1609 // Adjust the types of any constant operands.1610 if (SelectInst *SI = dyn_cast<SelectInst>(V)) {1611 if (isa<ConstantPointerNull, ConstantAggregateZero>(SI->getOperand(1)))1612 SI->setOperand(1, Constant::getNullValue(NewTy));1613 1614 if (isa<ConstantPointerNull, ConstantAggregateZero>(SI->getOperand(2)))1615 SI->setOperand(2, Constant::getNullValue(NewTy));1616 } else if (PHINode *Phi = dyn_cast<PHINode>(V)) {1617 for (unsigned I = 0, E = Phi->getNumIncomingValues(); I != E; ++I) {1618 if (isa<ConstantPointerNull, ConstantAggregateZero>(1619 Phi->getIncomingValue(I)))1620 Phi->setIncomingValue(I, Constant::getNullValue(NewTy));1621 }1622 }1623 1624 continue;1625 }1626 1627 IntrinsicInst *Intr = cast<IntrinsicInst>(Call);1628 Builder.SetInsertPoint(Intr);1629 switch (Intr->getIntrinsicID()) {1630 case Intrinsic::lifetime_start:1631 case Intrinsic::lifetime_end:1632 // These intrinsics are for address space 0 only1633 Intr->eraseFromParent();1634 continue;1635 case Intrinsic::memcpy:1636 case Intrinsic::memmove:1637 // These have 2 pointer operands. In case if second pointer also needs1638 // to be replaced we defer processing of these intrinsics until all1639 // other values are processed.1640 DeferredIntrs.push_back(Intr);1641 continue;1642 case Intrinsic::memset: {1643 MemSetInst *MemSet = cast<MemSetInst>(Intr);1644 Builder.CreateMemSet(MemSet->getRawDest(), MemSet->getValue(),1645 MemSet->getLength(), MemSet->getDestAlign(),1646 MemSet->isVolatile());1647 Intr->eraseFromParent();1648 continue;1649 }1650 case Intrinsic::invariant_start:1651 case Intrinsic::invariant_end:1652 case Intrinsic::launder_invariant_group:1653 case Intrinsic::strip_invariant_group: {1654 SmallVector<Value *> Args;1655 if (Intr->getIntrinsicID() == Intrinsic::invariant_start) {1656 Args.emplace_back(Intr->getArgOperand(0));1657 } else if (Intr->getIntrinsicID() == Intrinsic::invariant_end) {1658 Args.emplace_back(Intr->getArgOperand(0));1659 Args.emplace_back(Intr->getArgOperand(1));1660 }1661 Args.emplace_back(Offset);1662 Function *F = Intrinsic::getOrInsertDeclaration(1663 Intr->getModule(), Intr->getIntrinsicID(), Offset->getType());1664 CallInst *NewIntr =1665 CallInst::Create(F, Args, Intr->getName(), Intr->getIterator());1666 Intr->mutateType(NewIntr->getType());1667 Intr->replaceAllUsesWith(NewIntr);1668 Intr->eraseFromParent();1669 continue;1670 }1671 case Intrinsic::objectsize: {1672 Value *Src = Intr->getOperand(0);1673 1674 CallInst *NewCall = Builder.CreateIntrinsic(1675 Intrinsic::objectsize,1676 {Intr->getType(), PointerType::get(Context, AMDGPUAS::LOCAL_ADDRESS)},1677 {Src, Intr->getOperand(1), Intr->getOperand(2), Intr->getOperand(3)});1678 Intr->replaceAllUsesWith(NewCall);1679 Intr->eraseFromParent();1680 continue;1681 }1682 default:1683 Intr->print(errs());1684 llvm_unreachable("Don't know how to promote alloca intrinsic use.");1685 }1686 }1687 1688 for (IntrinsicInst *Intr : DeferredIntrs) {1689 Builder.SetInsertPoint(Intr);1690 Intrinsic::ID ID = Intr->getIntrinsicID();1691 assert(ID == Intrinsic::memcpy || ID == Intrinsic::memmove);1692 1693 MemTransferInst *MI = cast<MemTransferInst>(Intr);1694 auto *B = Builder.CreateMemTransferInst(1695 ID, MI->getRawDest(), MI->getDestAlign(), MI->getRawSource(),1696 MI->getSourceAlign(), MI->getLength(), MI->isVolatile());1697 1698 for (unsigned I = 0; I != 2; ++I) {1699 if (uint64_t Bytes = Intr->getParamDereferenceableBytes(I)) {1700 B->addDereferenceableParamAttr(I, Bytes);1701 }1702 }1703 1704 Intr->eraseFromParent();1705 }1706 1707 return true;1708}1709