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-exception
6//
7//===----------------------------------------------------------------------===//
8//
9// Eliminates allocas by either converting them into vectors or by migrating
10// 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 only
14// promotes to vector. Promotion to vector is almost always profitable
15// except when the alloca is too big and the promotion would result in
16// very high register pressure.
17// - "promote-alloca", which does both promotion to vector and LDS and runs
18// much later in the pipeline. This runs after SROA because promoting to
19// LDS is of course less profitable than getting rid of the alloca or
20// vectorizing it, thus we only want to do it when the only alternative is
21// lowering the alloca to stack.
22//
23// Note that both of them exist for the old and new PMs. The new PM passes are
24// declared in AMDGPU.h and the legacy PM ones are declared here.s
25//
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/Support/MathExtras.h"
46#include "llvm/Target/TargetMachine.h"
47#include "llvm/Transforms/Utils/SSAUpdater.h"
48
49#define DEBUG_TYPE "amdgpu-promote-alloca"
50
51using namespace llvm;
52
53namespace {
54
55static cl::opt<bool>
56 DisablePromoteAllocaToVector("disable-promote-alloca-to-vector",
57 cl::desc("Disable promote alloca to vector"),
58 cl::init(Val: false));
59
60static cl::opt<bool>
61 DisablePromoteAllocaToLDS("disable-promote-alloca-to-lds",
62 cl::desc("Disable promote alloca to LDS"),
63 cl::init(Val: false));
64
65static cl::opt<unsigned> PromoteAllocaToVectorLimit(
66 "amdgpu-promote-alloca-to-vector-limit",
67 cl::desc("Maximum byte size to consider promote alloca to vector"),
68 cl::init(Val: 0));
69
70static cl::opt<unsigned> PromoteAllocaToVectorMaxRegs(
71 "amdgpu-promote-alloca-to-vector-max-regs",
72 cl::desc(
73 "Maximum vector size (in 32b registers) to use when promoting alloca"),
74 cl::init(Val: 32));
75
76// Use up to 1/4 of available register budget for vectorization.
77// FIXME: Increase the limit for whole function budgets? Perhaps x2?
78static cl::opt<unsigned> PromoteAllocaToVectorVGPRRatio(
79 "amdgpu-promote-alloca-to-vector-vgpr-ratio",
80 cl::desc("Ratio of VGPRs to budget for promoting alloca to vectors"),
81 cl::init(Val: 4));
82
83static cl::opt<unsigned>
84 LoopUserWeight("promote-alloca-vector-loop-user-weight",
85 cl::desc("The bonus weight of users of allocas within loop "
86 "when sorting profitable allocas"),
87 cl::init(Val: 4));
88
89// We support vector indices of the form ((A * stride) >> shift) + B
90// VarIndex is A, VarMul is stride, VarShift is shift and ConstIndex is B. All
91// parts are optional.
92struct GEPToVectorIndex {
93 WeakTrackingVH VarIndex = nullptr; // defaults to 0
94 ConstantInt *VarMul = nullptr; // defaults to 1
95 ConstantInt *VarShift = nullptr; // defaults to 0
96 ConstantInt *ConstIndex = nullptr; // defaults to 0
97 Value *Full = nullptr;
98};
99
100struct MemTransferInfo {
101 ConstantInt *SrcIndex = nullptr;
102 ConstantInt *DestIndex = nullptr;
103};
104
105// Analysis for planning the different strategies of alloca promotion.
106struct AllocaAnalysis {
107 AllocaInst *Alloca = nullptr;
108 DenseSet<Value *> Pointers;
109 SmallVector<Use *> Uses;
110 unsigned Score = 0;
111 bool HaveSelectOrPHI = false;
112 struct {
113 FixedVectorType *Ty = nullptr;
114 SmallVector<Instruction *> Worklist;
115 SmallVector<Instruction *> UsersToRemove;
116 MapVector<GetElementPtrInst *, GEPToVectorIndex> GEPVectorIdx;
117 MapVector<MemTransferInst *, MemTransferInfo> TransferInfo;
118 } Vector;
119 struct {
120 bool Enable = false;
121 SmallVector<User *> Worklist;
122 } LDS;
123
124 explicit AllocaAnalysis(AllocaInst *Alloca) : Alloca(Alloca) {}
125};
126
127// Shared implementation which can do both promotion to vector and to LDS.
128class AMDGPUPromoteAllocaImpl {
129private:
130 const TargetMachine &TM;
131 LoopInfo &LI;
132 Module &Mod;
133 const DataLayout &DL;
134
135 // FIXME: This should be per-kernel.
136 uint32_t LocalMemLimit = 0;
137 uint32_t CurrentLocalMemUsage = 0;
138 unsigned MaxVGPRs;
139 unsigned VGPRBudgetRatio;
140 unsigned MaxVectorRegs;
141
142 bool IsAMDGCN = false;
143 bool IsAMDHSA = false;
144
145 std::pair<Value *, Value *> getLocalSizeYZ(IRBuilder<> &Builder);
146 Value *getWorkitemID(IRBuilder<> &Builder, unsigned N);
147
148 bool collectAllocaUses(AllocaAnalysis &AA) const;
149
150 /// Val is a derived pointer from Alloca. OpIdx0/OpIdx1 are the operand
151 /// indices to an instruction with 2 pointer inputs (e.g. select, icmp).
152 /// Returns true if both operands are derived from the same alloca. Val should
153 /// be the same value as one of the input operands of UseInst.
154 bool binaryOpIsDerivedFromSameAlloca(Value *Alloca, Value *Val,
155 Instruction *UseInst, int OpIdx0,
156 int OpIdx1) const;
157
158 /// Check whether we have enough local memory for promotion.
159 bool hasSufficientLocalMem(const Function &F);
160
161 FixedVectorType *getVectorTypeForAlloca(Type *AllocaTy) const;
162 void analyzePromoteToVector(AllocaAnalysis &AA) const;
163 void promoteAllocaToVector(AllocaAnalysis &AA);
164 void analyzePromoteToLDS(AllocaAnalysis &AA) const;
165 bool tryPromoteAllocaToLDS(AllocaAnalysis &AA, bool SufficientLDS,
166 SetVector<IntrinsicInst *> &DeferredIntrs);
167 void
168 finishDeferredAllocaToLDSPromotion(SetVector<IntrinsicInst *> &DeferredIntrs);
169
170 void scoreAlloca(AllocaAnalysis &AA) const;
171
172 void setFunctionLimits(const Function &F);
173
174public:
175 AMDGPUPromoteAllocaImpl(TargetMachine &TM, Module &M, LoopInfo &LI)
176 : TM(TM), LI(LI), Mod(M), DL(M.getDataLayout()) {
177 const Triple &TT = M.getTargetTriple();
178 IsAMDGCN = TT.isAMDGCN();
179 IsAMDHSA = TT.getOS() == Triple::AMDHSA;
180 }
181
182 bool run(Function &F, bool PromoteToLDS);
183};
184
185// FIXME: This can create globals so should be a module pass.
186class AMDGPUPromoteAlloca : public FunctionPass {
187public:
188 static char ID;
189
190 AMDGPUPromoteAlloca() : FunctionPass(ID) {}
191
192 bool runOnFunction(Function &F) override {
193 if (skipFunction(F))
194 return false;
195 if (auto *TPC = getAnalysisIfAvailable<TargetPassConfig>())
196 return AMDGPUPromoteAllocaImpl(
197 TPC->getTM<TargetMachine>(), *F.getParent(),
198 getAnalysis<LoopInfoWrapperPass>().getLoopInfo())
199 .run(F, /*PromoteToLDS*/ true);
200 return false;
201 }
202
203 StringRef getPassName() const override { return "AMDGPU Promote Alloca"; }
204
205 void getAnalysisUsage(AnalysisUsage &AU) const override {
206 AU.setPreservesCFG();
207 AU.addRequired<LoopInfoWrapperPass>();
208 FunctionPass::getAnalysisUsage(AU);
209 }
210};
211
212static unsigned getMaxVGPRs(unsigned LDSBytes, const TargetMachine &TM,
213 const Function &F) {
214 const GCNSubtarget &ST = TM.getSubtarget<GCNSubtarget>(F);
215
216 unsigned DynamicVGPRBlockSize = AMDGPU::getDynamicVGPRBlockSize(F);
217 unsigned MaxVGPRs = ST.getMaxNumVGPRs(
218 WavesPerEU: ST.getWavesPerEU(FlatWorkGroupSizes: ST.getFlatWorkGroupSizes(F), LDSBytes, F).first,
219 DynamicVGPRBlockSize);
220
221 // A DVGPR wave launches with a single VGPR block allocated.
222 if (DynamicVGPRBlockSize != 0 &&
223 AMDGPU::isEntryFunctionCC(CC: F.getCallingConv()))
224 MaxVGPRs = std::min(a: MaxVGPRs, b: DynamicVGPRBlockSize);
225
226 // A non-entry function has only 32 caller preserved registers.
227 // Do not promote alloca which will force spilling unless we know the function
228 // will be inlined.
229 if (!F.hasFnAttribute(Kind: Attribute::AlwaysInline) &&
230 !AMDGPU::isEntryFunctionCC(CC: F.getCallingConv()))
231 MaxVGPRs = std::min(a: MaxVGPRs, b: 32u);
232 return MaxVGPRs;
233}
234
235} // end anonymous namespace
236
237char AMDGPUPromoteAlloca::ID = 0;
238
239INITIALIZE_PASS_BEGIN(AMDGPUPromoteAlloca, DEBUG_TYPE,
240 "AMDGPU promote alloca to vector or LDS", false, false)
241// Move LDS uses from functions to kernels before promote alloca for accurate
242// estimation of LDS available
243INITIALIZE_PASS_DEPENDENCY(AMDGPULowerModuleLDSLegacy)
244INITIALIZE_PASS_DEPENDENCY(LoopInfoWrapperPass)
245INITIALIZE_PASS_END(AMDGPUPromoteAlloca, DEBUG_TYPE,
246 "AMDGPU promote alloca to vector or LDS", false, false)
247
248char &llvm::AMDGPUPromoteAllocaID = AMDGPUPromoteAlloca::ID;
249
250PreservedAnalyses AMDGPUPromoteAllocaPass::run(Function &F,
251 FunctionAnalysisManager &AM) {
252 auto &LI = AM.getResult<LoopAnalysis>(IR&: F);
253 bool Changed = AMDGPUPromoteAllocaImpl(TM, *F.getParent(), LI)
254 .run(F, /*PromoteToLDS=*/true);
255 if (Changed) {
256 PreservedAnalyses PA;
257 PA.preserveSet<CFGAnalyses>();
258 return PA;
259 }
260 return PreservedAnalyses::all();
261}
262
263PreservedAnalyses
264AMDGPUPromoteAllocaToVectorPass::run(Function &F, FunctionAnalysisManager &AM) {
265 auto &LI = AM.getResult<LoopAnalysis>(IR&: F);
266 bool Changed = AMDGPUPromoteAllocaImpl(TM, *F.getParent(), LI)
267 .run(F, /*PromoteToLDS=*/false);
268 if (Changed) {
269 PreservedAnalyses PA;
270 PA.preserveSet<CFGAnalyses>();
271 return PA;
272 }
273 return PreservedAnalyses::all();
274}
275
276FunctionPass *llvm::createAMDGPUPromoteAlloca() {
277 return new AMDGPUPromoteAlloca();
278}
279
280bool AMDGPUPromoteAllocaImpl::collectAllocaUses(AllocaAnalysis &AA) const {
281 const auto RejectUser = [&](Instruction *Inst, Twine Msg) {
282 LLVM_DEBUG(dbgs() << " Cannot promote alloca: " << Msg << "\n"
283 << " " << *Inst << "\n");
284 return false;
285 };
286
287 SmallVector<Instruction *, 4> WorkList({AA.Alloca});
288 while (!WorkList.empty()) {
289 auto *Cur = WorkList.pop_back_val();
290 if (find(Range&: AA.Pointers, Val: Cur) != AA.Pointers.end())
291 continue;
292 AA.Pointers.insert(V: Cur);
293 for (auto &U : Cur->uses()) {
294 auto *Inst = cast<Instruction>(Val: U.getUser());
295 if (isa<StoreInst>(Val: Inst)) {
296 if (U.getOperandNo() != StoreInst::getPointerOperandIndex()) {
297 return RejectUser(Inst, "pointer escapes via store");
298 }
299 }
300 AA.Uses.push_back(Elt: &U);
301
302 if (isa<GetElementPtrInst>(Val: U.getUser())) {
303 WorkList.push_back(Elt: Inst);
304 } else if (auto *SI = dyn_cast<SelectInst>(Val: Inst)) {
305 // Only promote a select if we know that the other select operand is
306 // from another pointer that will also be promoted.
307 if (!binaryOpIsDerivedFromSameAlloca(Alloca: AA.Alloca, Val: Cur, UseInst: SI, OpIdx0: 1, OpIdx1: 2))
308 return RejectUser(Inst, "select from mixed objects");
309 WorkList.push_back(Elt: Inst);
310 AA.HaveSelectOrPHI = true;
311 } else if (auto *Phi = dyn_cast<PHINode>(Val: Inst)) {
312 // Repeat for phis.
313
314 // TODO: Handle more complex cases. We should be able to replace loops
315 // over arrays.
316 switch (Phi->getNumIncomingValues()) {
317 case 1:
318 break;
319 case 2:
320 if (!binaryOpIsDerivedFromSameAlloca(Alloca: AA.Alloca, Val: Cur, UseInst: Phi, OpIdx0: 0, OpIdx1: 1))
321 return RejectUser(Inst, "phi from mixed objects");
322 break;
323 default:
324 return RejectUser(Inst, "phi with too many operands");
325 }
326
327 WorkList.push_back(Elt: Inst);
328 AA.HaveSelectOrPHI = true;
329 }
330 }
331 }
332 return true;
333}
334
335void AMDGPUPromoteAllocaImpl::scoreAlloca(AllocaAnalysis &AA) const {
336 LLVM_DEBUG(dbgs() << "Scoring: " << *AA.Alloca << "\n");
337 unsigned Score = 0;
338 // Increment score by one for each user + a bonus for users within loops.
339 for (auto *U : AA.Uses) {
340 Instruction *Inst = cast<Instruction>(Val: U->getUser());
341 if (isa<GetElementPtrInst>(Val: Inst) || isa<SelectInst>(Val: Inst) ||
342 isa<PHINode>(Val: Inst))
343 continue;
344 unsigned UserScore =
345 1 + (LoopUserWeight * LI.getLoopDepth(BB: Inst->getParent()));
346 LLVM_DEBUG(dbgs() << " [+" << UserScore << "]:\t" << *Inst << "\n");
347 Score += UserScore;
348 }
349 LLVM_DEBUG(dbgs() << " => Final Score:" << Score << "\n");
350 AA.Score = Score;
351}
352
353void AMDGPUPromoteAllocaImpl::setFunctionLimits(const Function &F) {
354 // Load per function limits, overriding with global options where appropriate.
355 // R600 register tuples/aliasing are fragile with large vector promotions so
356 // apply architecture specific limit here.
357 const int R600MaxVectorRegs = 16;
358 MaxVectorRegs = F.getFnAttributeAsParsedInteger(
359 Kind: "amdgpu-promote-alloca-to-vector-max-regs",
360 Default: IsAMDGCN ? PromoteAllocaToVectorMaxRegs : R600MaxVectorRegs);
361 if (PromoteAllocaToVectorMaxRegs.getNumOccurrences())
362 MaxVectorRegs = PromoteAllocaToVectorMaxRegs;
363 VGPRBudgetRatio = F.getFnAttributeAsParsedInteger(
364 Kind: "amdgpu-promote-alloca-to-vector-vgpr-ratio",
365 Default: PromoteAllocaToVectorVGPRRatio);
366 if (PromoteAllocaToVectorVGPRRatio.getNumOccurrences())
367 VGPRBudgetRatio = PromoteAllocaToVectorVGPRRatio;
368}
369
370bool AMDGPUPromoteAllocaImpl::run(Function &F, bool PromoteToLDS) {
371 if (DisablePromoteAllocaToLDS && DisablePromoteAllocaToVector)
372 return false;
373
374 bool SufficientLDS = PromoteToLDS && hasSufficientLocalMem(F);
375 MaxVGPRs = IsAMDGCN ? getMaxVGPRs(LDSBytes: CurrentLocalMemUsage, TM, F) : 128;
376 setFunctionLimits(F);
377
378 unsigned VectorizationBudget =
379 (PromoteAllocaToVectorLimit ? PromoteAllocaToVectorLimit * 8
380 : (MaxVGPRs * 32)) /
381 VGPRBudgetRatio;
382
383 std::vector<AllocaAnalysis> Allocas;
384 for (Instruction &I : F.getEntryBlock()) {
385 if (AllocaInst *AI = dyn_cast<AllocaInst>(Val: &I)) {
386 // Array allocations are probably not worth handling, since an allocation
387 // of the array type is the canonical form.
388 if (!AI->isStaticAlloca() || AI->isArrayAllocation())
389 continue;
390
391 LLVM_DEBUG(dbgs() << "Analyzing: " << *AI << '\n');
392
393 AllocaAnalysis AA{AI};
394 if (collectAllocaUses(AA)) {
395 analyzePromoteToVector(AA);
396 if (PromoteToLDS)
397 analyzePromoteToLDS(AA);
398 if (AA.Vector.Ty || AA.LDS.Enable) {
399 scoreAlloca(AA);
400 Allocas.push_back(x: std::move(AA));
401 }
402 }
403 }
404 }
405
406 stable_sort(Range&: Allocas,
407 C: [](const auto &A, const auto &B) { return A.Score > B.Score; });
408
409 // clang-format off
410 LLVM_DEBUG(
411 dbgs() << "Sorted Worklist:\n";
412 for (const auto &AA : Allocas)
413 dbgs() << " " << *AA.Alloca << "\n";
414 );
415 // clang-format on
416
417 bool Changed = false;
418 SetVector<IntrinsicInst *> DeferredIntrs;
419 for (AllocaAnalysis &AA : Allocas) {
420 if (AA.Vector.Ty) {
421 std::optional<TypeSize> Size = AA.Alloca->getAllocationSize(DL);
422 assert(Size); // Expected to succeed on non-array alloca.
423 const unsigned AllocaCost = Size->getFixedValue() * 8;
424 // First, check if we have enough budget to vectorize this alloca.
425 if (AllocaCost <= VectorizationBudget) {
426 promoteAllocaToVector(AA);
427 Changed = true;
428 assert((VectorizationBudget - AllocaCost) < VectorizationBudget &&
429 "Underflow!");
430 VectorizationBudget -= AllocaCost;
431 LLVM_DEBUG(dbgs() << " Remaining vectorization budget:"
432 << VectorizationBudget << "\n");
433 continue;
434 } else {
435 LLVM_DEBUG(dbgs() << "Alloca too big for vectorization (size:"
436 << AllocaCost << ", budget:" << VectorizationBudget
437 << "): " << *AA.Alloca << "\n");
438 }
439 }
440
441 if (AA.LDS.Enable &&
442 tryPromoteAllocaToLDS(AA, SufficientLDS, DeferredIntrs))
443 Changed = true;
444 }
445 finishDeferredAllocaToLDSPromotion(DeferredIntrs);
446
447 // NOTE: tryPromoteAllocaToVector removes the alloca, so Allocas contains
448 // dangling pointers. If we want to reuse it past this point, the loop above
449 // would need to be updated to remove successfully promoted allocas.
450
451 return Changed;
452}
453
454// Checks if the instruction I is a memset user of the alloca AI that we can
455// deal with. Currently, only non-volatile memsets that affect the whole alloca
456// are handled.
457static bool isSupportedMemset(MemSetInst *I, AllocaInst *AI,
458 const DataLayout &DL) {
459 using namespace PatternMatch;
460 // For now we only care about non-volatile memsets that affect the whole type
461 // (start at index 0 and fill the whole alloca).
462 //
463 // TODO: Now that we moved to PromoteAlloca we could handle any memsets
464 // (except maybe volatile ones?) - we just need to use shufflevector if it
465 // only affects a subset of the vector.
466 const unsigned Size = DL.getTypeStoreSize(Ty: AI->getAllocatedType());
467 return I->getOperand(i_nocapture: 0) == AI &&
468 match(V: I->getOperand(i_nocapture: 2), P: m_SpecificInt(V: Size)) && !I->isVolatile();
469}
470
471static Value *calculateVectorIndex(Value *Ptr, AllocaAnalysis &AA) {
472 IRBuilder<> B(*AA.Alloca->getModule());
473
474 Ptr = Ptr->stripPointerCasts();
475 if (Ptr == AA.Alloca)
476 return B.getInt32(C: 0);
477
478 auto *GEP = cast<GetElementPtrInst>(Val: Ptr);
479 auto I = AA.Vector.GEPVectorIdx.find(Key: GEP);
480 assert(I != AA.Vector.GEPVectorIdx.end() && "Must have entry for GEP!");
481
482 if (!I->second.Full) {
483 Value *Result = nullptr;
484 B.SetInsertPoint(GEP);
485
486 if (I->second.VarIndex) {
487 Result = I->second.VarIndex;
488 Result = B.CreateSExtOrTrunc(V: Result, DestTy: B.getInt32Ty());
489
490 if (I->second.VarMul)
491 Result = B.CreateMul(LHS: Result, RHS: I->second.VarMul);
492
493 if (I->second.VarShift)
494 Result = B.CreateAShr(LHS: Result, RHS: I->second.VarShift, Name: "", /*isExact*/ true);
495 }
496
497 if (I->second.ConstIndex) {
498 if (Result)
499 Result = B.CreateAdd(LHS: Result, RHS: I->second.ConstIndex);
500 else
501 Result = I->second.ConstIndex;
502 }
503
504 if (!Result)
505 Result = B.getInt32(C: 0);
506
507 I->second.Full = Result;
508 }
509
510 return I->second.Full;
511}
512
513static std::optional<GEPToVectorIndex>
514computeGEPToVectorIndex(GetElementPtrInst *GEP, AllocaInst *Alloca,
515 Type *VecElemTy, const DataLayout &DL) {
516 // TODO: Extracting a "multiple of X" from a GEP might be a useful generic
517 // helper.
518 LLVMContext &Ctx = GEP->getContext();
519 unsigned BW = DL.getIndexTypeSizeInBits(Ty: GEP->getType());
520 SmallMapVector<Value *, APInt, 4> VarOffsets;
521 APInt ConstOffset(BW, 0);
522
523 // Walk backwards through nested GEPs to collect both constant and variable
524 // offsets, so that nested vector GEP chains can be lowered in one step.
525 //
526 // Given this IR fragment as input:
527 //
528 // %0 = alloca [10 x <2 x i32>], align 8, addrspace(5)
529 // %1 = getelementptr [10 x <2 x i32>], ptr addrspace(5) %0, i32 0, i32 %j
530 // %2 = getelementptr i8, ptr addrspace(5) %1, i32 4
531 // %3 = load i32, ptr addrspace(5) %2, align 4
532 //
533 // Combine both GEP operations in a single pass, producing:
534 // BasePtr = %0
535 // ConstOffset = 4
536 // VarOffsets = { %j -> element_size(<2 x i32>) }
537 //
538 // That lets us emit a single buffer_load directly into a VGPR, without ever
539 // allocating scratch memory for the intermediate pointer.
540 Value *CurPtr = GEP;
541 while (auto *CurGEP = dyn_cast<GetElementPtrInst>(Val: CurPtr)) {
542 if (!CurGEP->collectOffset(DL, BitWidth: BW, VariableOffsets&: VarOffsets, ConstantOffset&: ConstOffset))
543 return {};
544
545 // Move to the next outer pointer.
546 CurPtr = CurGEP->getPointerOperand();
547 }
548
549 assert(CurPtr == Alloca && "GEP not based on alloca");
550
551 int64_t VecElemSize = DL.getTypeAllocSize(Ty: VecElemTy);
552 if (VarOffsets.size() > 1)
553 return {};
554
555 // We support vector indices of the form ((VarIndex * stride) >> shift) + B.
556 // IndexQuot represents B. Check that the constant offset is a multiple
557 // of the vector element size.
558 if (ConstOffset.srem(RHS: VecElemSize) != 0)
559 return {};
560 APInt IndexQuot = ConstOffset.sdiv(RHS: VecElemSize);
561
562 GEPToVectorIndex Result;
563
564 if (!ConstOffset.isZero())
565 Result.ConstIndex = ConstantInt::get(Context&: Ctx, V: IndexQuot.sextOrTrunc(width: BW));
566
567 // If there are no variable offsets, only a constant offset, then we're done.
568 if (VarOffsets.empty())
569 return Result;
570
571 // Scale is the stride in the (A * stride) part. Check that there is only one
572 // variable offset and extract the scale factor.
573 const auto &VarOffset = VarOffsets.front();
574 auto ScaleOpt = VarOffset.second.tryZExtValue();
575 if (!ScaleOpt || *ScaleOpt == 0)
576 return {};
577
578 uint64_t Scale = *ScaleOpt;
579 Result.VarIndex = VarOffset.first;
580 auto *OffsetType = dyn_cast<IntegerType>(Val: Result.VarIndex->getType());
581 if (!OffsetType)
582 return {};
583
584 // The vector index for the variable part is: VarIndex * Scale / VecElemSize.
585 if (Scale >= (uint64_t)VecElemSize) {
586 if (Scale % VecElemSize != 0)
587 return {};
588
589 // Scale is a multiple of VecElemSize, so the index is just: VarIndex *
590 // (Scale / VecElemSize).
591 uint64_t VarMul = Scale / VecElemSize;
592 // Only the multiplier is needed.
593 if (VarMul != 1)
594 Result.VarMul = ConstantInt::get(Context&: Ctx, V: APInt(BW, VarMul));
595 } else {
596 if ((uint64_t)VecElemSize % Scale != 0)
597 return {};
598
599 // VecElemSize is a multiple of Scale, so the index is just: VarIndex /
600 // (VecElemSize / Scale).
601 uint64_t Divisor = VecElemSize / Scale;
602 // The divisor must be a power of 2 so we can use a right shift.
603 if (!isPowerOf2_64(Value: Divisor))
604 return {};
605
606 // VarIndex must be known to be divisible by that divisor.
607 KnownBits KB = computeKnownBits(V: VarOffset.first, DL);
608 if (KB.countMinTrailingZeros() < Log2_64(Value: Divisor))
609 return {};
610
611 Result.VarShift = ConstantInt::get(Context&: Ctx, V: APInt(BW, Log2_64(Value: Divisor)));
612 }
613
614 return Result;
615}
616
617/// Promotes a single user of the alloca to a vector form.
618///
619/// \param Inst Instruction to be promoted.
620/// \param DL Module Data Layout.
621/// \param AA Alloca Analysis.
622/// \param VecStoreSize Size of \p VectorTy in bytes.
623/// \param ElementSize Size of \p VectorTy element type in bytes.
624/// \param CurVal Current value of the vector (e.g. last stored value)
625/// \param[out] DeferredLoads \p Inst is added to this vector if it can't
626/// be promoted now. This happens when promoting requires \p
627/// CurVal, but \p CurVal is nullptr.
628/// \return the stored value if \p Inst would have written to the alloca, or
629/// nullptr otherwise.
630static Value *promoteAllocaUserToVector(Instruction *Inst, const DataLayout &DL,
631 AllocaAnalysis &AA,
632 unsigned VecStoreSize,
633 unsigned ElementSize,
634 function_ref<Value *()> GetCurVal) {
635 // Note: we use InstSimplifyFolder because it can leverage the DataLayout
636 // to do more folding, especially in the case of vector splats.
637 IRBuilder<InstSimplifyFolder> Builder(Inst->getIterator(),
638 InstSimplifyFolder(DL));
639
640 Type *VecEltTy = AA.Vector.Ty->getElementType();
641
642 switch (Inst->getOpcode()) {
643 case Instruction::Load: {
644 Value *CurVal = GetCurVal();
645 Value *Index =
646 calculateVectorIndex(Ptr: cast<LoadInst>(Val: Inst)->getPointerOperand(), AA);
647
648 // We're loading the full vector.
649 Type *AccessTy = Inst->getType();
650 TypeSize AccessSize = DL.getTypeStoreSize(Ty: AccessTy);
651 if (Constant *CI = dyn_cast<Constant>(Val: Index)) {
652 if (CI->isNullValue() && AccessSize == VecStoreSize) {
653 Inst->replaceAllUsesWith(
654 V: Builder.CreateBitPreservingCastChain(DL, V: CurVal, NewTy: AccessTy));
655 return nullptr;
656 }
657 }
658
659 // Loading a subvector, or a scalar that spans several elements.
660 TypeSize EltSize = DL.getTypeStoreSize(Ty: VecEltTy);
661 assert(AccessSize.isKnownMultipleOf(EltSize) &&
662 "promotable access must cover a whole number of elements");
663 const unsigned NumLoadedElts = AccessSize / EltSize;
664 if (NumLoadedElts > 1) {
665 auto *SubVecTy = FixedVectorType::get(ElementType: VecEltTy, NumElts: NumLoadedElts);
666 assert(DL.getTypeStoreSize(SubVecTy) == DL.getTypeStoreSize(AccessTy));
667
668 // If idx is dynamic, then sandwich load with bitcasts.
669 // ie. VectorTy SubVecTy AccessTy
670 // <64 x i8> -> <16 x i8> <8 x i16>
671 // <64 x i8> -> <4 x i128> -> i128 -> <8 x i16>
672 // Extracting subvector with dynamic index has very large expansion in
673 // the amdgpu backend. Limit to pow2.
674 FixedVectorType *VectorTy = AA.Vector.Ty;
675 TypeSize NumBits = DL.getTypeStoreSize(Ty: SubVecTy) * 8u;
676 uint64_t LoadAlign = cast<LoadInst>(Val: Inst)->getAlign().value();
677 bool IsAlignedLoad = NumBits <= (LoadAlign * 8u);
678 unsigned TotalNumElts = VectorTy->getNumElements();
679 bool IsProperlyDivisible = TotalNumElts % NumLoadedElts == 0;
680 if (!isa<ConstantInt>(Val: Index) &&
681 llvm::isPowerOf2_32(Value: SubVecTy->getNumElements()) &&
682 IsProperlyDivisible && IsAlignedLoad) {
683 IntegerType *NewElemTy = Builder.getIntNTy(N: NumBits);
684 const unsigned NewNumElts =
685 DL.getTypeStoreSize(Ty: VectorTy) * 8u / NumBits;
686 const unsigned LShrAmt = llvm::Log2_32(Value: SubVecTy->getNumElements());
687 FixedVectorType *BitCastTy =
688 FixedVectorType::get(ElementType: NewElemTy, NumElts: NewNumElts);
689 Value *BCVal =
690 Builder.CreateBitPreservingCastChain(DL, V: CurVal, NewTy: BitCastTy);
691 Value *NewIdx = Builder.CreateLShr(
692 LHS: Index, RHS: ConstantInt::get(Ty: Index->getType(), V: LShrAmt));
693 Value *ExtVal = Builder.CreateExtractElement(Vec: BCVal, Idx: NewIdx);
694 Value *BCOut =
695 Builder.CreateBitPreservingCastChain(DL, V: ExtVal, NewTy: AccessTy);
696 Inst->replaceAllUsesWith(V: BCOut);
697 return nullptr;
698 }
699
700 Value *SubVec = PoisonValue::get(T: SubVecTy);
701 for (unsigned K = 0; K < NumLoadedElts; ++K) {
702 Value *CurIdx =
703 Builder.CreateAdd(LHS: Index, RHS: ConstantInt::get(Ty: Index->getType(), V: K));
704 SubVec = Builder.CreateInsertElement(
705 Vec: SubVec, NewElt: Builder.CreateExtractElement(Vec: CurVal, Idx: CurIdx), Idx: K);
706 }
707
708 Inst->replaceAllUsesWith(
709 V: Builder.CreateBitPreservingCastChain(DL, V: SubVec, NewTy: AccessTy));
710 return nullptr;
711 }
712
713 // We're loading one element.
714 Value *ExtractElement = Builder.CreateExtractElement(Vec: CurVal, Idx: Index);
715 if (AccessTy != VecEltTy)
716 ExtractElement = Builder.CreateBitOrPointerCast(V: ExtractElement, DestTy: AccessTy);
717
718 Inst->replaceAllUsesWith(V: ExtractElement);
719 return nullptr;
720 }
721 case Instruction::Store: {
722 // For stores, it's a bit trickier and it depends on whether we're storing
723 // the full vector or not. If we're storing the full vector, we don't need
724 // to know the current value. If this is a store of a single element, we
725 // need to know the value.
726 StoreInst *SI = cast<StoreInst>(Val: Inst);
727 Value *Index = calculateVectorIndex(Ptr: SI->getPointerOperand(), AA);
728 Value *Val = SI->getValueOperand();
729
730 // We're storing the full vector, we can handle this without knowing CurVal.
731 Type *AccessTy = Val->getType();
732 TypeSize AccessSize = DL.getTypeStoreSize(Ty: AccessTy);
733 if (Constant *CI = dyn_cast<Constant>(Val: Index)) {
734 if (CI->isNullValue() && AccessSize == VecStoreSize) {
735 Value *Result =
736 Builder.CreateBitPreservingCastChain(DL, V: Val, NewTy: AA.Vector.Ty);
737 // If Result is a load from this alloca, it will later be RAUW'd and
738 // deleted. The SSAUpdater holds a raw Value* that RAUW doesn't update,
739 // leaving a dangling pointer. Wrap in a freeze to create a fresh value
740 // the SSAUpdater can safely hold; the freeze's operand is a proper IR
741 // use that RAUW does update.
742 if (isa<LoadInst>(Val: Result))
743 Result = Builder.CreateFreeze(V: Result);
744 return Result;
745 }
746 }
747
748 // Storing a subvector, or a scalar that spans several elements.
749 TypeSize EltSize = DL.getTypeStoreSize(Ty: VecEltTy);
750 assert(AccessSize.isKnownMultipleOf(EltSize) &&
751 "promotable access must cover a whole number of elements");
752 const unsigned NumWrittenElts = AccessSize / EltSize;
753 if (NumWrittenElts > 1) {
754 const unsigned NumVecElts = AA.Vector.Ty->getNumElements();
755 auto *SubVecTy = FixedVectorType::get(ElementType: VecEltTy, NumElts: NumWrittenElts);
756 assert(DL.getTypeStoreSize(SubVecTy) == DL.getTypeStoreSize(AccessTy));
757
758 Val = Builder.CreateBitPreservingCastChain(DL, V: Val, NewTy: SubVecTy);
759 Value *CurVec = GetCurVal();
760 for (unsigned K = 0, NumElts = std::min(a: NumWrittenElts, b: NumVecElts);
761 K < NumElts; ++K) {
762 Value *CurIdx =
763 Builder.CreateAdd(LHS: Index, RHS: ConstantInt::get(Ty: Index->getType(), V: K));
764 CurVec = Builder.CreateInsertElement(
765 Vec: CurVec, NewElt: Builder.CreateExtractElement(Vec: Val, Idx: K), Idx: CurIdx);
766 }
767 return CurVec;
768 }
769
770 if (Val->getType() != VecEltTy)
771 Val = Builder.CreateBitOrPointerCast(V: Val, DestTy: VecEltTy);
772 return Builder.CreateInsertElement(Vec: GetCurVal(), NewElt: Val, Idx: Index);
773 }
774 case Instruction::Call: {
775 if (auto *MTI = dyn_cast<MemTransferInst>(Val: Inst)) {
776 // For memcpy, we need to know curval.
777 ConstantInt *Length = cast<ConstantInt>(Val: MTI->getLength());
778 unsigned NumCopied = Length->getZExtValue() / ElementSize;
779 MemTransferInfo *TI = &AA.Vector.TransferInfo[MTI];
780 unsigned SrcBegin = TI->SrcIndex->getZExtValue();
781 unsigned DestBegin = TI->DestIndex->getZExtValue();
782
783 SmallVector<int> Mask;
784 for (unsigned Idx = 0; Idx < AA.Vector.Ty->getNumElements(); ++Idx) {
785 if (Idx >= DestBegin && Idx < DestBegin + NumCopied) {
786 Mask.push_back(Elt: SrcBegin < AA.Vector.Ty->getNumElements()
787 ? SrcBegin++
788 : PoisonMaskElem);
789 } else {
790 Mask.push_back(Elt: Idx);
791 }
792 }
793
794 return Builder.CreateShuffleVector(V: GetCurVal(), Mask);
795 }
796
797 if (auto *MSI = dyn_cast<MemSetInst>(Val: Inst)) {
798 // For memset, we don't need to know the previous value because we
799 // currently only allow memsets that cover the whole alloca.
800 Value *Elt = MSI->getOperand(i_nocapture: 1);
801 const unsigned BytesPerElt = DL.getTypeStoreSize(Ty: VecEltTy);
802 if (BytesPerElt > 1) {
803 Value *EltBytes = Builder.CreateVectorSplat(NumElts: BytesPerElt, V: Elt);
804
805 // If the element type of the vector is a pointer, we need to first cast
806 // to an integer, then use a PtrCast.
807 if (VecEltTy->isPointerTy()) {
808 Type *PtrInt = Builder.getIntNTy(N: BytesPerElt * 8);
809 Elt = Builder.CreateBitCast(V: EltBytes, DestTy: PtrInt);
810 Elt = Builder.CreateIntToPtr(V: Elt, DestTy: VecEltTy);
811 } else
812 Elt = Builder.CreateBitCast(V: EltBytes, DestTy: VecEltTy);
813 }
814
815 return Builder.CreateVectorSplat(EC: AA.Vector.Ty->getElementCount(), V: Elt);
816 }
817
818 if (auto *Intr = dyn_cast<IntrinsicInst>(Val: Inst)) {
819 if (Intr->getIntrinsicID() == Intrinsic::objectsize) {
820 Intr->replaceAllUsesWith(
821 V: Builder.getIntN(N: Intr->getType()->getIntegerBitWidth(),
822 C: DL.getTypeAllocSize(Ty: AA.Vector.Ty)));
823 return nullptr;
824 }
825 }
826
827 llvm_unreachable("Unsupported call when promoting alloca to vector");
828 }
829
830 default:
831 llvm_unreachable("Inconsistency in instructions promotable to vector");
832 }
833
834 llvm_unreachable("Did not return after promoting instruction!");
835}
836
837static bool isSupportedAccessType(FixedVectorType *VecTy, Type *AccessTy,
838 const DataLayout &DL) {
839 // An access that covers several elements can work if its size is a multiple
840 // of the size of the alloca's vector element type, since it can be split
841 // across consecutive elements. This covers accesses by a vector type, as well
842 // as scalar accesses that are wider than one element, which happens when an
843 // object is written one element at a time but read back in wider pieces.
844 //
845 // Examples:
846 // - VecTy = <8 x float>, AccessTy = <4 x float> -> OK
847 // - VecTy = <4 x double>, AccessTy = <2 x float> -> OK
848 // - VecTy = <4 x double>, AccessTy = <3 x float> -> NOT OK
849 // - 3*32 is not a multiple of 64
850 // - VecTy = <8 x i32>, AccessTy = i64 -> OK
851 //
852 // We could handle more complicated cases, but it'd make things a lot more
853 // complicated.
854 if (isa<FixedVectorType>(Val: AccessTy) || AccessTy->isIntegerTy() ||
855 AccessTy->isFloatingPointTy()) {
856 TypeSize AccTS = DL.getTypeStoreSize(Ty: AccessTy);
857 TypeSize VecTS = DL.getTypeStoreSize(Ty: VecTy->getElementType());
858 // If the type size and the store size don't match, we would need to do more
859 // than just bitcast to translate between an extracted/insertable subvectors
860 // and the accessed value.
861 if (AccTS * 8 == DL.getTypeSizeInBits(Ty: AccessTy) && AccTS > VecTS &&
862 AccTS.isKnownMultipleOf(RHS: VecTS))
863 return true;
864 }
865
866 // An access that covers exactly one element only needs a cast.
867 return CastInst::isBitOrNoopPointerCastable(SrcTy: VecTy->getElementType(), DestTy: AccessTy,
868 DL);
869}
870
871/// Iterates over an instruction worklist that may contain multiple instructions
872/// from the same basic block, but in a different order.
873template <typename InstContainer>
874static void forEachWorkListItem(const InstContainer &WorkList,
875 std::function<void(Instruction *)> Fn) {
876 // Bucket up uses of the alloca by the block they occur in.
877 // This is important because we have to handle multiple defs/uses in a block
878 // ourselves: SSAUpdater is purely for cross-block references.
879 DenseMap<BasicBlock *, SmallDenseSet<Instruction *>> UsesByBlock;
880 for (Instruction *User : WorkList)
881 UsesByBlock[User->getParent()].insert(V: User);
882
883 for (Instruction *User : WorkList) {
884 BasicBlock *BB = User->getParent();
885 auto &BlockUses = UsesByBlock[BB];
886
887 // Already processed, skip.
888 if (BlockUses.empty())
889 continue;
890
891 // Only user in the block, directly process it.
892 if (BlockUses.size() == 1) {
893 Fn(User);
894 continue;
895 }
896
897 // Multiple users in the block, do a linear scan to see users in order.
898 for (Instruction &Inst : *BB) {
899 if (!BlockUses.contains(V: &Inst))
900 continue;
901
902 Fn(&Inst);
903 }
904
905 // Clear the block so we know it's been processed.
906 BlockUses.clear();
907 }
908}
909
910/// Find an insert point after an alloca, after all other allocas clustered at
911/// the start of the block.
912static BasicBlock::iterator skipToNonAllocaInsertPt(BasicBlock &BB,
913 BasicBlock::iterator I) {
914 for (BasicBlock::iterator E = BB.end(); I != E && isa<AllocaInst>(Val: *I); ++I)
915 ;
916 return I;
917}
918
919/// Peel nested aggregates down to a single uniform element type, multiplying
920/// NumElems by the element count of each layer peeled.
921static Type *peelAggregateToElementType(Type *Ty, uint64_t &NumElems) {
922 while (true) {
923 if (auto *ArrayTy = dyn_cast<ArrayType>(Val: Ty)) {
924 NumElems *= ArrayTy->getNumElements();
925 Ty = ArrayTy->getElementType();
926 continue;
927 }
928
929 auto *StructTy = dyn_cast<StructType>(Val: Ty);
930 if (!StructTy || !StructTy->containsHomogeneousTypes())
931 break;
932
933 NumElems *= StructTy->getNumElements();
934 Ty = StructTy->getElementType(N: 0);
935 }
936
937 return Ty;
938}
939
940FixedVectorType *
941AMDGPUPromoteAllocaImpl::getVectorTypeForAlloca(Type *AllocaTy) const {
942 if (DisablePromoteAllocaToVector) {
943 LLVM_DEBUG(dbgs() << " Promote alloca to vectors is disabled\n");
944 return nullptr;
945 }
946
947 auto *VectorTy = dyn_cast<FixedVectorType>(Val: AllocaTy);
948 if (AllocaTy->isAggregateType()) {
949 uint64_t NumElems = 1;
950 Type *ElemTy = peelAggregateToElementType(Ty: AllocaTy, NumElems);
951
952 // Check for array of vectors
953 auto *InnerVectorTy = dyn_cast<FixedVectorType>(Val: ElemTy);
954 if (InnerVectorTy) {
955 NumElems *= InnerVectorTy->getNumElements();
956 ElemTy = InnerVectorTy->getElementType();
957 }
958
959 if (VectorType::isValidElementType(ElemTy) && NumElems > 0) {
960 unsigned ElementSize = DL.getTypeSizeInBits(Ty: ElemTy) / 8;
961 if (ElementSize > 0) {
962 unsigned AllocaSize = DL.getTypeStoreSize(Ty: AllocaTy);
963 // Expand vector if required to match padding of inner type,
964 // i.e. odd size subvectors.
965 // Storage size of new vector must match that of alloca for correct
966 // behaviour of byte offsets and GEP computation.
967 if (NumElems * ElementSize != AllocaSize)
968 NumElems = AllocaSize / ElementSize;
969 if (NumElems > 0 && (AllocaSize % ElementSize) == 0)
970 VectorTy = FixedVectorType::get(ElementType: ElemTy, NumElts: NumElems);
971 }
972 }
973 }
974 if (!VectorTy) {
975 LLVM_DEBUG(dbgs() << " Cannot convert type to vector\n");
976 return nullptr;
977 }
978
979 const unsigned MaxElements =
980 (MaxVectorRegs * 32) / DL.getTypeSizeInBits(Ty: VectorTy->getElementType());
981
982 if (VectorTy->getNumElements() > MaxElements ||
983 VectorTy->getNumElements() < 2) {
984 LLVM_DEBUG(dbgs() << " " << *VectorTy
985 << " has an unsupported number of elements\n");
986 return nullptr;
987 }
988
989 Type *VecEltTy = VectorTy->getElementType();
990 unsigned ElementSizeInBits = DL.getTypeSizeInBits(Ty: VecEltTy);
991 if (ElementSizeInBits != DL.getTypeAllocSizeInBits(Ty: VecEltTy)) {
992 LLVM_DEBUG(dbgs() << " Cannot convert to vector if the allocation size "
993 "does not match the type's size\n");
994 return nullptr;
995 }
996
997 return VectorTy;
998}
999
1000void AMDGPUPromoteAllocaImpl::analyzePromoteToVector(AllocaAnalysis &AA) const {
1001 if (AA.HaveSelectOrPHI) {
1002 LLVM_DEBUG(dbgs() << " Cannot convert to vector due to select or phi\n");
1003 return;
1004 }
1005
1006 Type *AllocaTy = AA.Alloca->getAllocatedType();
1007 AA.Vector.Ty = getVectorTypeForAlloca(AllocaTy);
1008 if (!AA.Vector.Ty)
1009 return;
1010
1011 const auto RejectUser = [&](Instruction *Inst, Twine Msg) {
1012 LLVM_DEBUG(dbgs() << " Cannot promote alloca to vector: " << Msg << "\n"
1013 << " " << *Inst << "\n");
1014 AA.Vector.Ty = nullptr;
1015 };
1016
1017 Type *VecEltTy = AA.Vector.Ty->getElementType();
1018 unsigned ElementSize = DL.getTypeSizeInBits(Ty: VecEltTy) / 8;
1019 assert(ElementSize > 0);
1020 for (auto *U : AA.Uses) {
1021 Instruction *Inst = cast<Instruction>(Val: U->getUser());
1022
1023 if (Value *Ptr = getLoadStorePointerOperand(V: Inst)) {
1024 assert(!isa<StoreInst>(Inst) ||
1025 U->getOperandNo() == StoreInst::getPointerOperandIndex());
1026
1027 Type *AccessTy = getLoadStoreType(I: Inst);
1028 if (AccessTy->isAggregateType())
1029 return RejectUser(Inst, "unsupported load/store as aggregate");
1030 assert(!AccessTy->isAggregateType() || AccessTy->isArrayTy());
1031
1032 // Check that this is a simple access of a vector element.
1033 bool IsSimple = isa<LoadInst>(Val: Inst) ? cast<LoadInst>(Val: Inst)->isSimple()
1034 : cast<StoreInst>(Val: Inst)->isSimple();
1035 if (!IsSimple)
1036 return RejectUser(Inst, "not a simple load or store");
1037
1038 Ptr = Ptr->stripPointerCasts();
1039
1040 // Alloca already accessed as vector.
1041 if (Ptr == AA.Alloca &&
1042 DL.getTypeStoreSize(Ty: AA.Alloca->getAllocatedType()) ==
1043 DL.getTypeStoreSize(Ty: AccessTy)) {
1044 AA.Vector.Worklist.push_back(Elt: Inst);
1045 continue;
1046 }
1047
1048 if (!isSupportedAccessType(VecTy: AA.Vector.Ty, AccessTy, DL))
1049 return RejectUser(Inst, "not a supported access type");
1050
1051 AA.Vector.Worklist.push_back(Elt: Inst);
1052 continue;
1053 }
1054
1055 if (auto *GEP = dyn_cast<GetElementPtrInst>(Val: Inst)) {
1056 // If we can't compute a vector index from this GEP, then we can't
1057 // promote this alloca to vector.
1058 auto Index = computeGEPToVectorIndex(GEP, Alloca: AA.Alloca, VecElemTy: VecEltTy, DL);
1059 if (!Index)
1060 return RejectUser(Inst, "cannot compute vector index for GEP");
1061
1062 AA.Vector.GEPVectorIdx[GEP] = std::move(Index.value());
1063 AA.Vector.UsersToRemove.push_back(Elt: Inst);
1064 continue;
1065 }
1066
1067 if (MemSetInst *MSI = dyn_cast<MemSetInst>(Val: Inst);
1068 MSI && isSupportedMemset(I: MSI, AI: AA.Alloca, DL)) {
1069 AA.Vector.Worklist.push_back(Elt: Inst);
1070 continue;
1071 }
1072
1073 if (MemTransferInst *TransferInst = dyn_cast<MemTransferInst>(Val: Inst)) {
1074 if (TransferInst->isVolatile())
1075 return RejectUser(Inst, "mem transfer inst is volatile");
1076
1077 ConstantInt *Len = dyn_cast<ConstantInt>(Val: TransferInst->getLength());
1078 if (!Len || (Len->getZExtValue() % ElementSize))
1079 return RejectUser(Inst, "mem transfer inst length is non-constant or "
1080 "not a multiple of the vector element size");
1081
1082 auto getConstIndexIntoAlloca = [&](Value *Ptr) -> ConstantInt * {
1083 if (Ptr == AA.Alloca)
1084 return ConstantInt::get(Context&: Ptr->getContext(), V: APInt(32, 0));
1085
1086 GetElementPtrInst *GEP = cast<GetElementPtrInst>(Val: Ptr);
1087 const auto &GEPI = AA.Vector.GEPVectorIdx.find(Key: GEP)->second;
1088 if (GEPI.VarIndex)
1089 return nullptr;
1090 if (GEPI.ConstIndex)
1091 return GEPI.ConstIndex;
1092 return ConstantInt::get(Context&: Ptr->getContext(), V: APInt(32, 0));
1093 };
1094
1095 MemTransferInfo *TI =
1096 &AA.Vector.TransferInfo.try_emplace(Key: TransferInst).first->second;
1097 unsigned OpNum = U->getOperandNo();
1098 if (OpNum == 0) {
1099 Value *Dest = TransferInst->getDest();
1100 ConstantInt *Index = getConstIndexIntoAlloca(Dest);
1101 if (!Index)
1102 return RejectUser(Inst, "could not calculate constant dest index");
1103 TI->DestIndex = Index;
1104 } else {
1105 assert(OpNum == 1);
1106 Value *Src = TransferInst->getSource();
1107 ConstantInt *Index = getConstIndexIntoAlloca(Src);
1108 if (!Index)
1109 return RejectUser(Inst, "could not calculate constant src index");
1110 TI->SrcIndex = Index;
1111 }
1112 continue;
1113 }
1114
1115 if (auto *Intr = dyn_cast<IntrinsicInst>(Val: Inst)) {
1116 if (Intr->getIntrinsicID() == Intrinsic::objectsize) {
1117 AA.Vector.Worklist.push_back(Elt: Inst);
1118 continue;
1119 }
1120 }
1121
1122 // Ignore assume-like intrinsics and comparisons used in assumes.
1123 if (isAssumeLikeIntrinsic(I: Inst)) {
1124 if (!Inst->use_empty())
1125 return RejectUser(Inst, "assume-like intrinsic cannot have any users");
1126 AA.Vector.UsersToRemove.push_back(Elt: Inst);
1127 continue;
1128 }
1129
1130 if (isa<ICmpInst>(Val: Inst) && all_of(Range: Inst->users(), P: [](User *U) {
1131 return isAssumeLikeIntrinsic(I: cast<Instruction>(Val: U));
1132 })) {
1133 AA.Vector.UsersToRemove.push_back(Elt: Inst);
1134 continue;
1135 }
1136
1137 return RejectUser(Inst, "unhandled alloca user");
1138 }
1139
1140 // Follow-up check to ensure we've seen both sides of all transfer insts.
1141 for (const auto &Entry : AA.Vector.TransferInfo) {
1142 const MemTransferInfo &TI = Entry.second;
1143 if (!TI.SrcIndex || !TI.DestIndex)
1144 return RejectUser(Entry.first,
1145 "mem transfer inst between different objects");
1146 AA.Vector.Worklist.push_back(Elt: Entry.first);
1147 }
1148}
1149
1150void AMDGPUPromoteAllocaImpl::promoteAllocaToVector(AllocaAnalysis &AA) {
1151 LLVM_DEBUG(dbgs() << "Promoting to vectors: " << *AA.Alloca << '\n');
1152 LLVM_DEBUG(dbgs() << " type conversion: " << *AA.Alloca->getAllocatedType()
1153 << " -> " << *AA.Vector.Ty << '\n');
1154 const unsigned VecStoreSize = DL.getTypeStoreSize(Ty: AA.Vector.Ty);
1155
1156 Type *VecEltTy = AA.Vector.Ty->getElementType();
1157 const unsigned ElementSize = DL.getTypeSizeInBits(Ty: VecEltTy) / 8;
1158
1159 // Alloca is uninitialized memory. Imitate that by making the first value
1160 // undef.
1161 SSAUpdater Updater;
1162 Updater.Initialize(Ty: AA.Vector.Ty, Name: "promotealloca");
1163
1164 BasicBlock *EntryBB = AA.Alloca->getParent();
1165 BasicBlock::iterator InitInsertPos =
1166 skipToNonAllocaInsertPt(BB&: *EntryBB, I: AA.Alloca->getIterator());
1167 IRBuilder<> Builder(&*InitInsertPos);
1168 Value *AllocaInitValue = Builder.CreateFreeze(V: PoisonValue::get(T: AA.Vector.Ty));
1169 AllocaInitValue->takeName(V: AA.Alloca);
1170
1171 Updater.AddAvailableValue(BB: AA.Alloca->getParent(), V: AllocaInitValue);
1172
1173 // First handle the initial worklist, in basic block order.
1174 //
1175 // Insert a placeholder whenever we need the vector value at the top of a
1176 // basic block.
1177 SmallSetVector<Instruction *, 8> Placeholders;
1178 forEachWorkListItem(WorkList: AA.Vector.Worklist, Fn: [&](Instruction *I) {
1179 BasicBlock *BB = I->getParent();
1180 auto GetCurVal = [&]() -> Value * {
1181 if (Value *CurVal = Updater.FindValueForBlock(BB))
1182 return CurVal;
1183
1184 if (!Placeholders.empty() && Placeholders.back()->getParent() == BB)
1185 return Placeholders.back();
1186
1187 // If the current value in the basic block is not yet known, insert a
1188 // placeholder that we will replace later.
1189 IRBuilder<> Builder(I);
1190 auto *Placeholder = cast<Instruction>(Val: Builder.CreateFreeze(
1191 V: PoisonValue::get(T: AA.Vector.Ty), Name: "promotealloca.placeholder"));
1192 Placeholders.insert(X: Placeholder);
1193 return Placeholders.back();
1194 };
1195
1196 Value *Result = promoteAllocaUserToVector(Inst: I, DL, AA, VecStoreSize,
1197 ElementSize, GetCurVal);
1198 // If the returned result is a placeholder, it means the instruction does
1199 // not really modify the alloca. So no need to make it being available value
1200 // to SSAUpdater.
1201 // This will stop placeholder being cached in SSAUpdater. The cached
1202 // placeholder may cause stale pointer being referenced when doing
1203 // placeholder replacement.
1204 if (Result && (!isa<Instruction>(Val: Result) ||
1205 !Placeholders.contains(key: cast<Instruction>(Val: Result))))
1206 Updater.AddAvailableValue(BB, V: Result);
1207 });
1208
1209 // Now fixup the placeholders.
1210 for (Instruction *Placeholder : Placeholders) {
1211 Placeholder->replaceAllUsesWith(
1212 V: Updater.GetValueInMiddleOfBlock(BB: Placeholder->getParent()));
1213 Placeholder->eraseFromParent();
1214 }
1215
1216 // Delete all instructions.
1217 for (Instruction *I : AA.Vector.Worklist) {
1218 assert(I->use_empty());
1219 I->eraseFromParent();
1220 }
1221
1222 // Delete all the users that are known to be removeable.
1223 for (Instruction *I : reverse(C&: AA.Vector.UsersToRemove)) {
1224 I->dropDroppableUses();
1225 assert(I->use_empty());
1226 I->eraseFromParent();
1227 }
1228
1229 // Alloca should now be dead too.
1230 assert(AA.Alloca->use_empty());
1231 AA.Alloca->eraseFromParent();
1232}
1233
1234std::pair<Value *, Value *>
1235AMDGPUPromoteAllocaImpl::getLocalSizeYZ(IRBuilder<> &Builder) {
1236 Function &F = *Builder.GetInsertBlock()->getParent();
1237 const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, F);
1238
1239 if (!IsAMDHSA) {
1240 CallInst *LocalSizeY = Builder.CreateIntrinsicWithoutFolding(
1241 ID: Intrinsic::r600_read_local_size_y, Args: {});
1242 CallInst *LocalSizeZ = Builder.CreateIntrinsicWithoutFolding(
1243 ID: Intrinsic::r600_read_local_size_z, Args: {});
1244
1245 ST.makeLIDRangeMetadata(I: LocalSizeY);
1246 ST.makeLIDRangeMetadata(I: LocalSizeZ);
1247
1248 return std::pair(LocalSizeY, LocalSizeZ);
1249 }
1250
1251 // We must read the size out of the dispatch pointer.
1252 assert(IsAMDGCN);
1253
1254 // We are indexing into this struct, and want to extract the workgroup_size_*
1255 // fields.
1256 //
1257 // typedef struct hsa_kernel_dispatch_packet_s {
1258 // uint16_t header;
1259 // uint16_t setup;
1260 // uint16_t workgroup_size_x ;
1261 // uint16_t workgroup_size_y;
1262 // uint16_t workgroup_size_z;
1263 // uint16_t reserved0;
1264 // uint32_t grid_size_x ;
1265 // uint32_t grid_size_y ;
1266 // uint32_t grid_size_z;
1267 //
1268 // uint32_t private_segment_size;
1269 // uint32_t group_segment_size;
1270 // uint64_t kernel_object;
1271 //
1272 // #ifdef HSA_LARGE_MODEL
1273 // void *kernarg_address;
1274 // #elif defined HSA_LITTLE_ENDIAN
1275 // void *kernarg_address;
1276 // uint32_t reserved1;
1277 // #else
1278 // uint32_t reserved1;
1279 // void *kernarg_address;
1280 // #endif
1281 // uint64_t reserved2;
1282 // hsa_signal_t completion_signal; // uint64_t wrapper
1283 // } hsa_kernel_dispatch_packet_t
1284 //
1285 CallInst *DispatchPtr =
1286 Builder.CreateIntrinsicWithoutFolding(ID: Intrinsic::amdgcn_dispatch_ptr, Args: {});
1287 DispatchPtr->addRetAttr(Kind: Attribute::NoAlias);
1288 DispatchPtr->addRetAttr(Kind: Attribute::NonNull);
1289 F.removeFnAttr(Kind: "amdgpu-no-dispatch-ptr");
1290
1291 // Size of the dispatch packet struct.
1292 DispatchPtr->addDereferenceableRetAttr(Bytes: 64);
1293
1294 Type *I32Ty = Type::getInt32Ty(C&: Mod.getContext());
1295
1296 // We could do a single 64-bit load here, but it's likely that the basic
1297 // 32-bit and extract sequence is already present, and it is probably easier
1298 // to CSE this. The loads should be mergeable later anyway.
1299 Value *GEPXY = Builder.CreateConstInBoundsGEP1_64(Ty: I32Ty, Ptr: DispatchPtr, Idx0: 1);
1300 LoadInst *LoadXY = Builder.CreateAlignedLoad(Ty: I32Ty, Ptr: GEPXY, Align: Align(4));
1301
1302 Value *GEPZU = Builder.CreateConstInBoundsGEP1_64(Ty: I32Ty, Ptr: DispatchPtr, Idx0: 2);
1303 LoadInst *LoadZU = Builder.CreateAlignedLoad(Ty: I32Ty, Ptr: GEPZU, Align: Align(4));
1304
1305 MDNode *MD = MDNode::get(Context&: Mod.getContext(), MDs: {});
1306 LoadXY->setMetadata(KindID: LLVMContext::MD_invariant_load, Node: MD);
1307 LoadZU->setMetadata(KindID: LLVMContext::MD_invariant_load, Node: MD);
1308 ST.makeLIDRangeMetadata(I: LoadZU);
1309
1310 // Extract y component. Upper half of LoadZU should be zero already.
1311 Value *Y = Builder.CreateLShr(LHS: LoadXY, RHS: 16);
1312
1313 return std::pair(Y, LoadZU);
1314}
1315
1316Value *AMDGPUPromoteAllocaImpl::getWorkitemID(IRBuilder<> &Builder,
1317 unsigned N) {
1318 Function *F = Builder.GetInsertBlock()->getParent();
1319 const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, F: *F);
1320 Intrinsic::ID IntrID = Intrinsic::not_intrinsic;
1321 StringRef AttrName;
1322
1323 switch (N) {
1324 case 0:
1325 IntrID = IsAMDGCN ? (Intrinsic::ID)Intrinsic::amdgcn_workitem_id_x
1326 : (Intrinsic::ID)Intrinsic::r600_read_tidig_x;
1327 AttrName = "amdgpu-no-workitem-id-x";
1328 break;
1329 case 1:
1330 IntrID = IsAMDGCN ? (Intrinsic::ID)Intrinsic::amdgcn_workitem_id_y
1331 : (Intrinsic::ID)Intrinsic::r600_read_tidig_y;
1332 AttrName = "amdgpu-no-workitem-id-y";
1333 break;
1334
1335 case 2:
1336 IntrID = IsAMDGCN ? (Intrinsic::ID)Intrinsic::amdgcn_workitem_id_z
1337 : (Intrinsic::ID)Intrinsic::r600_read_tidig_z;
1338 AttrName = "amdgpu-no-workitem-id-z";
1339 break;
1340 default:
1341 llvm_unreachable("invalid dimension");
1342 }
1343
1344 Function *WorkitemIdFn = Intrinsic::getOrInsertDeclaration(M: &Mod, id: IntrID);
1345 CallInst *CI = Builder.CreateCall(Callee: WorkitemIdFn);
1346 ST.makeLIDRangeMetadata(I: CI);
1347 F->removeFnAttr(Kind: AttrName);
1348
1349 return CI;
1350}
1351
1352static bool isCallPromotable(CallInst *CI) {
1353 IntrinsicInst *II = dyn_cast<IntrinsicInst>(Val: CI);
1354 if (!II)
1355 return false;
1356
1357 switch (II->getIntrinsicID()) {
1358 case Intrinsic::memcpy:
1359 case Intrinsic::memmove:
1360 case Intrinsic::memset:
1361 case Intrinsic::lifetime_start:
1362 case Intrinsic::lifetime_end:
1363 case Intrinsic::invariant_start:
1364 case Intrinsic::invariant_end:
1365 case Intrinsic::launder_invariant_group:
1366 case Intrinsic::objectsize:
1367 return true;
1368 default:
1369 return false;
1370 }
1371}
1372
1373bool AMDGPUPromoteAllocaImpl::binaryOpIsDerivedFromSameAlloca(
1374 Value *BaseAlloca, Value *Val, Instruction *Inst, int OpIdx0,
1375 int OpIdx1) const {
1376 // Figure out which operand is the one we might not be promoting.
1377 Value *OtherOp = Inst->getOperand(i: OpIdx0);
1378 if (Val == OtherOp)
1379 OtherOp = Inst->getOperand(i: OpIdx1);
1380
1381 if (isa<ConstantPointerNull, ConstantAggregateZero>(Val: OtherOp))
1382 return true;
1383
1384 // TODO: getUnderlyingObject will not work on a vector getelementptr
1385 Value *OtherObj = getUnderlyingObject(V: OtherOp);
1386 if (!isa<AllocaInst>(Val: OtherObj))
1387 return false;
1388
1389 // TODO: We should be able to replace undefs with the right pointer type.
1390
1391 // TODO: If we know the other base object is another promotable
1392 // alloca, not necessarily this alloca, we can do this. The
1393 // important part is both must have the same address space at
1394 // the end.
1395 if (OtherObj != BaseAlloca) {
1396 LLVM_DEBUG(
1397 dbgs() << "Found a binary instruction with another alloca object\n");
1398 return false;
1399 }
1400
1401 return true;
1402}
1403
1404void AMDGPUPromoteAllocaImpl::analyzePromoteToLDS(AllocaAnalysis &AA) const {
1405 if (DisablePromoteAllocaToLDS) {
1406 LLVM_DEBUG(dbgs() << " Promote alloca to LDS is disabled\n");
1407 return;
1408 }
1409
1410 // Don't promote the alloca to LDS for shader calling conventions as the work
1411 // item ID intrinsics are not supported for these calling conventions.
1412 // Furthermore not all LDS is available for some of the stages.
1413 const Function &ContainingFunction = *AA.Alloca->getFunction();
1414 CallingConv::ID CC = ContainingFunction.getCallingConv();
1415
1416 switch (CC) {
1417 case CallingConv::AMDGPU_KERNEL:
1418 case CallingConv::SPIR_KERNEL:
1419 break;
1420 default:
1421 LLVM_DEBUG(
1422 dbgs()
1423 << " promote alloca to LDS not supported with calling convention.\n");
1424 return;
1425 }
1426
1427 for (Use *Use : AA.Uses) {
1428 auto *User = Use->getUser();
1429
1430 if (CallInst *CI = dyn_cast<CallInst>(Val: User)) {
1431 if (!isCallPromotable(CI))
1432 return;
1433
1434 if (find(Range&: AA.LDS.Worklist, Val: User) == AA.LDS.Worklist.end())
1435 AA.LDS.Worklist.push_back(Elt: User);
1436 continue;
1437 }
1438
1439 Instruction *UseInst = cast<Instruction>(Val: User);
1440 if (UseInst->getOpcode() == Instruction::PtrToInt)
1441 return;
1442
1443 if (LoadInst *LI = dyn_cast<LoadInst>(Val: UseInst)) {
1444 if (LI->isVolatile())
1445 return;
1446 continue;
1447 }
1448
1449 if (StoreInst *SI = dyn_cast<StoreInst>(Val: UseInst)) {
1450 if (SI->isVolatile())
1451 return;
1452 continue;
1453 }
1454
1455 if (AtomicRMWInst *RMW = dyn_cast<AtomicRMWInst>(Val: UseInst)) {
1456 if (RMW->isVolatile())
1457 return;
1458 continue;
1459 }
1460
1461 if (AtomicCmpXchgInst *CAS = dyn_cast<AtomicCmpXchgInst>(Val: UseInst)) {
1462 if (CAS->isVolatile())
1463 return;
1464 continue;
1465 }
1466
1467 // Only promote a select if we know that the other select operand
1468 // is from another pointer that will also be promoted.
1469 if (ICmpInst *ICmp = dyn_cast<ICmpInst>(Val: UseInst)) {
1470 if (!binaryOpIsDerivedFromSameAlloca(BaseAlloca: AA.Alloca, Val: Use->get(), Inst: ICmp, OpIdx0: 0, OpIdx1: 1))
1471 return;
1472
1473 // May need to rewrite constant operands.
1474 if (find(Range&: AA.LDS.Worklist, Val: User) == AA.LDS.Worklist.end())
1475 AA.LDS.Worklist.push_back(Elt: ICmp);
1476 continue;
1477 }
1478
1479 if (GetElementPtrInst *GEP = dyn_cast<GetElementPtrInst>(Val: UseInst)) {
1480 // Be conservative if an address could be computed outside the bounds of
1481 // the alloca.
1482 if (!GEP->isInBounds())
1483 return;
1484 } else if (!isa<ExtractElementInst, SelectInst, PHINode>(Val: User)) {
1485 // Do not promote vector/aggregate type instructions. It is hard to track
1486 // their users.
1487
1488 // Do not promote addrspacecast.
1489 //
1490 // TODO: If we know the address is only observed through flat pointers, we
1491 // could still promote.
1492 return;
1493 }
1494
1495 if (find(Range&: AA.LDS.Worklist, Val: User) == AA.LDS.Worklist.end())
1496 AA.LDS.Worklist.push_back(Elt: User);
1497 }
1498
1499 AA.LDS.Enable = true;
1500}
1501
1502bool AMDGPUPromoteAllocaImpl::hasSufficientLocalMem(const Function &F) {
1503
1504 FunctionType *FTy = F.getFunctionType();
1505 const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, F);
1506
1507 // If the function has any arguments in the local address space, then it's
1508 // possible these arguments require the entire local memory space, so
1509 // we cannot use local memory in the pass.
1510 for (Type *ParamTy : FTy->params()) {
1511 PointerType *PtrTy = dyn_cast<PointerType>(Val: ParamTy);
1512 if (PtrTy && PtrTy->getAddressSpace() == AMDGPUAS::LOCAL_ADDRESS) {
1513 LocalMemLimit = 0;
1514 LLVM_DEBUG(dbgs() << "Function has local memory argument. Promoting to "
1515 "local memory disabled.\n");
1516 return false;
1517 }
1518 }
1519
1520 LocalMemLimit = ST.getAddressableLocalMemorySize();
1521 if (LocalMemLimit == 0)
1522 return false;
1523
1524 SmallVector<const Constant *, 16> Stack;
1525 SmallPtrSet<const Constant *, 8> VisitedConstants;
1526 SmallPtrSet<const GlobalVariable *, 8> UsedLDS;
1527
1528 auto visitUsers = [&](const GlobalVariable *GV, const Constant *Val) -> bool {
1529 for (const User *U : Val->users()) {
1530 if (const Instruction *Use = dyn_cast<Instruction>(Val: U)) {
1531 if (Use->getFunction() == &F)
1532 return true;
1533 } else {
1534 const Constant *C = cast<Constant>(Val: U);
1535 if (VisitedConstants.insert(Ptr: C).second)
1536 Stack.push_back(Elt: C);
1537 }
1538 }
1539
1540 return false;
1541 };
1542
1543 for (GlobalVariable &GV : Mod.globals()) {
1544 if (GV.getAddressSpace() != AMDGPUAS::LOCAL_ADDRESS)
1545 continue;
1546
1547 if (visitUsers(&GV, &GV)) {
1548 UsedLDS.insert(Ptr: &GV);
1549 Stack.clear();
1550 continue;
1551 }
1552
1553 // For any ConstantExpr uses, we need to recursively search the users until
1554 // we see a function.
1555 while (!Stack.empty()) {
1556 const Constant *C = Stack.pop_back_val();
1557 if (visitUsers(&GV, C)) {
1558 UsedLDS.insert(Ptr: &GV);
1559 Stack.clear();
1560 break;
1561 }
1562 }
1563 }
1564
1565 SmallVector<std::pair<uint64_t, Align>, 16> AllocatedSizes;
1566 AllocatedSizes.reserve(N: UsedLDS.size());
1567
1568 for (const GlobalVariable *GV : UsedLDS) {
1569 Align Alignment =
1570 DL.getValueOrABITypeAlignment(Alignment: GV->getAlign(), Ty: GV->getValueType());
1571 uint64_t AllocSize = GV->getGlobalSize(DL);
1572
1573 // HIP uses an extern unsized array in local address space for dynamically
1574 // allocated shared memory. In that case, we have to disable the promotion.
1575 if (GV->hasExternalLinkage() && AllocSize == 0) {
1576 LocalMemLimit = 0;
1577 LLVM_DEBUG(dbgs() << "Function has a reference to externally allocated "
1578 "local memory. Promoting to local memory "
1579 "disabled.\n");
1580 return false;
1581 }
1582
1583 AllocatedSizes.emplace_back(Args&: AllocSize, Args&: Alignment);
1584 }
1585
1586 // Sort to try to estimate the worst case alignment padding
1587 //
1588 // FIXME: We should really do something to fix the addresses to a more optimal
1589 // value instead
1590 llvm::sort(C&: AllocatedSizes, Comp: llvm::less_second());
1591
1592 // Check how much local memory is being used by global objects
1593 CurrentLocalMemUsage = 0;
1594
1595 // FIXME: Try to account for padding here. The real padding and address is
1596 // currently determined from the inverse order of uses in the function when
1597 // legalizing, which could also potentially change. We try to estimate the
1598 // worst case here, but we probably should fix the addresses earlier.
1599 for (auto Alloc : AllocatedSizes) {
1600 CurrentLocalMemUsage = alignTo(Size: CurrentLocalMemUsage, A: Alloc.second);
1601 CurrentLocalMemUsage += Alloc.first;
1602 }
1603
1604 unsigned MaxOccupancy =
1605 ST.getWavesPerEU(FlatWorkGroupSizes: ST.getFlatWorkGroupSizes(F), LDSBytes: CurrentLocalMemUsage, F)
1606 .second;
1607
1608 // Round up to the next tier of usage.
1609 unsigned MaxSizeWithWaveCount =
1610 ST.getMaxLocalMemSizeWithWaveCount(WaveCount: MaxOccupancy, F);
1611
1612 // Program may already use more LDS than is usable at maximum occupancy.
1613 if (CurrentLocalMemUsage > MaxSizeWithWaveCount)
1614 return false;
1615
1616 LocalMemLimit = MaxSizeWithWaveCount;
1617
1618 LLVM_DEBUG(dbgs() << F.getName() << " uses " << CurrentLocalMemUsage
1619 << " bytes of LDS\n"
1620 << " Rounding size to " << MaxSizeWithWaveCount
1621 << " with a maximum occupancy of " << MaxOccupancy << '\n'
1622 << " and " << (LocalMemLimit - CurrentLocalMemUsage)
1623 << " available for promotion\n");
1624
1625 return true;
1626}
1627
1628// FIXME: Should try to pick the most likely to be profitable allocas first.
1629bool AMDGPUPromoteAllocaImpl::tryPromoteAllocaToLDS(
1630 AllocaAnalysis &AA, bool SufficientLDS,
1631 SetVector<IntrinsicInst *> &DeferredIntrs) {
1632 LLVM_DEBUG(dbgs() << "Trying to promote to LDS: " << *AA.Alloca << '\n');
1633
1634 // Not likely to have sufficient local memory for promotion.
1635 if (!SufficientLDS)
1636 return false;
1637
1638 IRBuilder<> Builder(AA.Alloca);
1639
1640 const Function &ContainingFunction = *AA.Alloca->getParent()->getParent();
1641 const AMDGPUSubtarget &ST = AMDGPUSubtarget::get(TM, F: ContainingFunction);
1642 unsigned WorkGroupSize = ST.getFlatWorkGroupSizes(F: ContainingFunction).second;
1643
1644 Align Alignment = AA.Alloca->getAlign();
1645
1646 // FIXME: This computed padding is likely wrong since it depends on inverse
1647 // usage order.
1648 //
1649 // FIXME: It is also possible that if we're allowed to use all of the memory
1650 // could end up using more than the maximum due to alignment padding.
1651
1652 uint32_t NewSize = alignTo(Size: CurrentLocalMemUsage, A: Alignment);
1653 std::optional<TypeSize> ElemSize = AA.Alloca->getAllocationSize(DL);
1654 if (!ElemSize || ElemSize->isScalable())
1655 return false;
1656 TypeSize AllocSize = WorkGroupSize * *ElemSize;
1657 NewSize += AllocSize.getFixedValue();
1658
1659 if (NewSize > LocalMemLimit) {
1660 LLVM_DEBUG(dbgs() << " " << AllocSize
1661 << " bytes of local memory not available to promote\n");
1662 return false;
1663 }
1664
1665 CurrentLocalMemUsage = NewSize;
1666
1667 LLVM_DEBUG(dbgs() << "Promoting alloca to local memory\n");
1668
1669 Function *F = AA.Alloca->getFunction();
1670
1671 Type *GVTy = ArrayType::get(ElementType: AA.Alloca->getAllocatedType(), NumElements: WorkGroupSize);
1672 GlobalVariable *GV = new GlobalVariable(
1673 Mod, GVTy, false, GlobalValue::InternalLinkage, PoisonValue::get(T: GVTy),
1674 Twine(F->getName()) + Twine('.') + AA.Alloca->getName(), nullptr,
1675 GlobalVariable::NotThreadLocal, AMDGPUAS::LOCAL_ADDRESS);
1676 GV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
1677 GV->setAlignment(AA.Alloca->getAlign());
1678
1679 Value *TCntY, *TCntZ;
1680
1681 std::tie(args&: TCntY, args&: TCntZ) = getLocalSizeYZ(Builder);
1682 Value *TIdX = getWorkitemID(Builder, N: 0);
1683 Value *TIdY = getWorkitemID(Builder, N: 1);
1684 Value *TIdZ = getWorkitemID(Builder, N: 2);
1685
1686 Value *Tmp0 = Builder.CreateMul(LHS: TCntY, RHS: TCntZ, Name: "", HasNUW: true, HasNSW: true);
1687 Tmp0 = Builder.CreateMul(LHS: Tmp0, RHS: TIdX);
1688 Value *Tmp1 = Builder.CreateMul(LHS: TIdY, RHS: TCntZ, Name: "", HasNUW: true, HasNSW: true);
1689 Value *TID = Builder.CreateAdd(LHS: Tmp0, RHS: Tmp1);
1690 TID = Builder.CreateAdd(LHS: TID, RHS: TIdZ);
1691
1692 LLVMContext &Context = Mod.getContext();
1693 Value *Indices[] = {Constant::getNullValue(Ty: Type::getInt32Ty(C&: Context)), TID};
1694
1695 Value *Offset = Builder.CreateInBoundsGEP(Ty: GVTy, Ptr: GV, IdxList: Indices);
1696 AA.Alloca->mutateType(Ty: Offset->getType());
1697 AA.Alloca->replaceAllUsesWith(V: Offset);
1698 AA.Alloca->eraseFromParent();
1699
1700 PointerType *NewPtrTy = PointerType::get(C&: Context, AddressSpace: AMDGPUAS::LOCAL_ADDRESS);
1701
1702 for (Value *V : AA.LDS.Worklist) {
1703 CallInst *Call = dyn_cast<CallInst>(Val: V);
1704 if (!Call) {
1705 if (ICmpInst *CI = dyn_cast<ICmpInst>(Val: V)) {
1706 Value *LHS = CI->getOperand(i_nocapture: 0);
1707 Value *RHS = CI->getOperand(i_nocapture: 1);
1708
1709 Type *NewTy = LHS->getType()->getWithNewType(EltTy: NewPtrTy);
1710 if (isa<ConstantPointerNull, ConstantAggregateZero>(Val: LHS))
1711 CI->setOperand(i_nocapture: 0, Val_nocapture: Constant::getNullValue(Ty: NewTy));
1712
1713 if (isa<ConstantPointerNull, ConstantAggregateZero>(Val: RHS))
1714 CI->setOperand(i_nocapture: 1, Val_nocapture: Constant::getNullValue(Ty: NewTy));
1715
1716 continue;
1717 }
1718
1719 // The operand's value should be corrected on its own and we don't want to
1720 // touch the users.
1721 if (isa<AddrSpaceCastInst>(Val: V))
1722 continue;
1723
1724 assert(V->getType()->isPtrOrPtrVectorTy());
1725
1726 Type *NewTy = V->getType()->getWithNewType(EltTy: NewPtrTy);
1727 V->mutateType(Ty: NewTy);
1728
1729 // Adjust the types of any constant operands.
1730 if (SelectInst *SI = dyn_cast<SelectInst>(Val: V)) {
1731 if (isa<ConstantPointerNull, ConstantAggregateZero>(Val: SI->getOperand(i_nocapture: 1)))
1732 SI->setOperand(i_nocapture: 1, Val_nocapture: Constant::getNullValue(Ty: NewTy));
1733
1734 if (isa<ConstantPointerNull, ConstantAggregateZero>(Val: SI->getOperand(i_nocapture: 2)))
1735 SI->setOperand(i_nocapture: 2, Val_nocapture: Constant::getNullValue(Ty: NewTy));
1736 } else if (PHINode *Phi = dyn_cast<PHINode>(Val: V)) {
1737 for (unsigned I = 0, E = Phi->getNumIncomingValues(); I != E; ++I) {
1738 if (isa<ConstantPointerNull, ConstantAggregateZero>(
1739 Val: Phi->getIncomingValue(i: I)))
1740 Phi->setIncomingValue(i: I, V: Constant::getNullValue(Ty: NewTy));
1741 }
1742 }
1743
1744 continue;
1745 }
1746
1747 IntrinsicInst *Intr = cast<IntrinsicInst>(Val: Call);
1748 Builder.SetInsertPoint(Intr);
1749 switch (Intr->getIntrinsicID()) {
1750 case Intrinsic::lifetime_start:
1751 case Intrinsic::lifetime_end:
1752 // These intrinsics are for address space 0 only
1753 Intr->eraseFromParent();
1754 continue;
1755 case Intrinsic::memcpy:
1756 case Intrinsic::memmove:
1757 // These have 2 pointer operands. In case if second pointer also needs
1758 // to be replaced we defer processing of these intrinsics until all
1759 // other values are processed.
1760 DeferredIntrs.insert(X: Intr);
1761 continue;
1762 case Intrinsic::memset: {
1763 MemSetInst *MemSet = cast<MemSetInst>(Val: Intr);
1764 Builder.CreateMemSet(Ptr: MemSet->getRawDest(), Val: MemSet->getValue(),
1765 Size: MemSet->getLength(), Align: MemSet->getDestAlign(),
1766 isVolatile: MemSet->isVolatile());
1767 Intr->eraseFromParent();
1768 continue;
1769 }
1770 case Intrinsic::invariant_start:
1771 case Intrinsic::invariant_end:
1772 case Intrinsic::launder_invariant_group: {
1773 assert(Intr->getArgOperand(Intr->arg_size() - 1)->getType() == NewPtrTy &&
1774 "pointer operand should already have been promoted");
1775 Function *NewF = Intrinsic::getOrInsertDeclaration(
1776 M: Intr->getModule(), id: Intr->getIntrinsicID(), OverloadTys: NewPtrTy);
1777 Intr->mutateType(Ty: NewF->getReturnType());
1778 Intr->setCalledFunction(NewF);
1779 continue;
1780 }
1781 case Intrinsic::objectsize: {
1782 Value *Src = Intr->getOperand(i_nocapture: 0);
1783
1784 Value *NewCall = Builder.CreateIntrinsic(
1785 ID: Intrinsic::objectsize,
1786 OverloadTypes: {Intr->getType(), PointerType::get(C&: Context, AddressSpace: AMDGPUAS::LOCAL_ADDRESS)},
1787 Args: {Src, Intr->getOperand(i_nocapture: 1), Intr->getOperand(i_nocapture: 2), Intr->getOperand(i_nocapture: 3)});
1788 Intr->replaceAllUsesWith(V: NewCall);
1789 Intr->eraseFromParent();
1790 continue;
1791 }
1792 default:
1793 Intr->print(O&: errs());
1794 llvm_unreachable("Don't know how to promote alloca intrinsic use.");
1795 }
1796 }
1797
1798 return true;
1799}
1800
1801void AMDGPUPromoteAllocaImpl::finishDeferredAllocaToLDSPromotion(
1802 SetVector<IntrinsicInst *> &DeferredIntrs) {
1803
1804 for (IntrinsicInst *Intr : DeferredIntrs) {
1805 IRBuilder<> Builder(Intr);
1806 Builder.SetInsertPoint(Intr);
1807 Intrinsic::ID ID = Intr->getIntrinsicID();
1808 assert(ID == Intrinsic::memcpy || ID == Intrinsic::memmove);
1809
1810 MemTransferInst *MI = cast<MemTransferInst>(Val: Intr);
1811 auto *B = Builder.CreateMemTransferInst(
1812 IntrID: ID, Dst: MI->getRawDest(), DstAlign: MI->getDestAlign(), Src: MI->getRawSource(),
1813 SrcAlign: MI->getSourceAlign(), Size: MI->getLength(), isVolatile: MI->isVolatile());
1814
1815 for (unsigned I = 0; I != 2; ++I) {
1816 if (uint64_t Bytes = Intr->getParamDereferenceableBytes(i: I)) {
1817 B->addDereferenceableParamAttr(i: I, Bytes);
1818 }
1819 }
1820
1821 Intr->eraseFromParent();
1822 }
1823}
1824